1 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \ 2 // RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \ 3 // RUN: | FileCheck -check-prefixes=CHECK,DEFAULT %s 4 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa --gpu-max-threads-per-block=1024 \ 5 // RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \ 6 // RUN: | FileCheck -check-prefixes=CHECK,MAX1024 %s 7 // RUN: %clang_cc1 -triple nvptx \ 8 // RUN: -fcuda-is-device -emit-llvm -o - %s | FileCheck %s \ 9 // RUN: -check-prefix=NAMD 10 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \ 11 // RUN: -verify -o - -x hip %s | FileCheck -check-prefix=NAMD %s 12 13 #include "Inputs/cuda.h" 14 15 __global__ void flat_work_group_size_default() { 16 // CHECK: define{{.*}} amdgpu_kernel void @_Z28flat_work_group_size_defaultv() [[FLAT_WORK_GROUP_SIZE_DEFAULT:#[0-9]+]] 17 } 18 19 __attribute__((amdgpu_flat_work_group_size(32, 64))) // expected-no-diagnostics 20 __global__ void flat_work_group_size_32_64() { 21 // CHECK: define{{.*}} amdgpu_kernel void @_Z26flat_work_group_size_32_64v() [[FLAT_WORK_GROUP_SIZE_32_64:#[0-9]+]] 22 } 23 __attribute__((amdgpu_waves_per_eu(2))) // expected-no-diagnostics 24 __global__ void waves_per_eu_2() { 25 // CHECK: define{{.*}} amdgpu_kernel void @_Z14waves_per_eu_2v() [[WAVES_PER_EU_2:#[0-9]+]] 26 } 27 __attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics 28 __global__ void num_sgpr_32() { 29 // CHECK: define{{.*}} amdgpu_kernel void @_Z11num_sgpr_32v() [[NUM_SGPR_32:#[0-9]+]] 30 } 31 __attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics 32 __global__ void num_vgpr_64() { 33 // CHECK: define{{.*}} amdgpu_kernel void @_Z11num_vgpr_64v() [[NUM_VGPR_64:#[0-9]+]] 34 } 35 36 // Make sure this is silently accepted on other targets. 37 // NAMD-NOT: "amdgpu-flat-work-group-size" 38 // NAMD-NOT: "amdgpu-waves-per-eu" 39 // NAMD-NOT: "amdgpu-num-vgpr" 40 // NAMD-NOT: "amdgpu-num-sgpr" 41 42 // DEFAULT-DAG: attributes [[FLAT_WORK_GROUP_SIZE_DEFAULT]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024"{{.*}}"uniform-work-group-size"="true" 43 // MAX1024-DAG: attributes [[FLAT_WORK_GROUP_SIZE_DEFAULT]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024" 44 // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64]] = {{.*}}"amdgpu-flat-work-group-size"="32,64" 45 // CHECK-DAG: attributes [[WAVES_PER_EU_2]] = {{.*}}"amdgpu-waves-per-eu"="2" 46 // CHECK-DAG: attributes [[NUM_SGPR_32]] = {{.*}}"amdgpu-num-sgpr"="32" 47 // CHECK-DAG: attributes [[NUM_VGPR_64]] = {{.*}}"amdgpu-num-vgpr"="64" 48