1 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
2 // RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
3 // RUN: | FileCheck -check-prefix=PRECOV5 %s
6 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
7 // RUN: -fcuda-is-device -mcode-object-version=5 -emit-llvm -o - -x hip %s \
8 // RUN: | FileCheck -check-prefix=COV5 %s
10 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
11 // RUN: -fcuda-is-device -mcode-object-version=none -emit-llvm -o - -x hip %s \
12 // RUN: | FileCheck -check-prefix=COVNONE %s
14 #include "Inputs/cuda.h"
16 // PRECOV5-LABEL: test_get_workgroup_size
17 // PRECOV5: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
18 // PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4
19 // PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
20 // PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6
21 // PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
22 // PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8
23 // PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
25 // COV5-LABEL: test_get_workgroup_size
26 // COV5: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
27 // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12
28 // COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
29 // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14
30 // COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
31 // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16
32 // COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
35 // COVNONE-LABEL: test_get_workgroup_size
36 // COVNONE: load i32, ptr addrspace(4) @llvm.amdgcn.abi.version
37 // COVNONE: [[ABI5_X:%.*]] = icmp sge i32 %{{.*}}, 500
38 // COVNONE: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
39 // COVNONE: [[GEP_5_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12
40 // COVNONE: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
41 // COVNONE: [[GEP_4_X:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4
42 // COVNONE: select i1 [[ABI5_X]], ptr addrspace(4) [[GEP_5_X]], ptr addrspace(4) [[GEP_4_X]]
43 // COVNONE: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
45 // COVNONE: load i32, ptr addrspace(4) @llvm.amdgcn.abi.version
46 // COVNONE: [[ABI5_Y:%.*]] = icmp sge i32 %{{.*}}, 500
47 // COVNONE: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
48 // COVNONE: [[GEP_5_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14
49 // COVNONE: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
50 // COVNONE: [[GEP_4_Y:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6
51 // COVNONE: select i1 [[ABI5_Y]], ptr addrspace(4) [[GEP_5_Y]], ptr addrspace(4) [[GEP_4_Y]]
52 // COVNONE: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
54 // COVNONE: load i32, ptr addrspace(4) @llvm.amdgcn.abi.version
55 // COVNONE: [[ABI5_Z:%.*]] = icmp sge i32 %{{.*}}, 500
56 // COVNONE: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
57 // COVNONE: [[GEP_5_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16
58 // COVNONE: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
59 // COVNONE: [[GEP_4_Z:%.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8
60 // COVNONE: select i1 [[ABI5_Z]], ptr addrspace(4) [[GEP_5_Z]], ptr addrspace(4) [[GEP_4_Z]]
61 // COVNONE: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
63 __device__ void test_get_workgroup_size(int d, int *out)
66 case 0: *out = __builtin_amdgcn_workgroup_size_x(); break;
67 case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;
68 case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;
73 // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}