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

   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
   4
   5
   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
   9
  10 #include "Inputs/cuda.h"
  11
  12 // PRECOV5-LABEL: test_get_workgroup_size
  13 // PRECOV5: call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
  14 // PRECOV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 4
  15 // PRECOV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
  16 // PRECOV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 6
  17 // PRECOV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
  18 // PRECOV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 8
  19 // PRECOV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
  20
  21 // COV5-LABEL: test_get_workgroup_size
  22 // COV5: call align 8 dereferenceable(256) i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
  23 // COV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 12
  24 // COV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
  25 // COV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 14
  26 // COV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
  27 // COV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 16
  28 // COV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
  29 __device__ void test_get_workgroup_size(int d, int *out)
  30 {
  31   switch (d) {
  32   case 0: *out = __builtin_amdgcn_workgroup_size_x(); break;
  33   case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;
  34   case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;
  35   default: *out = 0;
  36   }
  37 }
  38
  39 // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}