clang/test/CodeGenCUDA/grid-constant.cu

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 5
   2 // RUN: %clang_cc1 "-triple" "nvptx64-nvidia-cuda" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
   3
   4 #include "Inputs/cuda.h"
   5
   6 struct S {};
   7
   8 __global__ void kernel(__grid_constant__ const S gc_arg1, int arg2, __grid_constant__ const int gc_arg3) {}
   9
  10 // dependent arguments get diagnosed after instantiation.
  11 template <typename T>
  12 __global__ void tkernel_const(__grid_constant__ const T arg) {}
  13
  14 template <typename T>
  15 __global__ void tkernel(int dummy, __grid_constant__ T arg) {}
  16
  17 void foo() {
  18   tkernel_const<const S><<<1,1>>>({});
  19   tkernel_const<S><<<1,1>>>({});
  20   tkernel<const S><<<1,1>>>(1, {});
  21 }
  22 //.
  23 //.
  24 // CHECK: [[META0:![0-9]+]] = !{ptr @_Z6kernel1Sii, !"kernel", i32 1, !"grid_constant", [[META1:![0-9]+]]}
  25 // CHECK: [[META1]] = !{i32 1, i32 3}
  26 // CHECK: [[META2:![0-9]+]] = !{ptr @_Z13tkernel_constIK1SEvT_, !"kernel", i32 1, !"grid_constant", [[META3:![0-9]+]]}
  27 // CHECK: [[META3]] = !{i32 1}
  28 // CHECK: [[META4:![0-9]+]] = !{ptr @_Z13tkernel_constI1SEvT_, !"kernel", i32 1, !"grid_constant", [[META3]]}
  29 // CHECK: [[META5:![0-9]+]] = !{ptr @_Z7tkernelIK1SEviT_, !"kernel", i32 1, !"grid_constant", [[META6:![0-9]+]]}
  30 // CHECK: [[META6]] = !{i32 2}
  31 //.