clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu

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