clang/test/CodeGenCUDA/anon-ns.cu

   1 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
   2 // RUN:   -aux-triple x86_64-unknown-linux-gnu -std=c++17 -fgpu-rdc \
   3 // RUN:   -emit-llvm -o - -x hip %s > %t.dev
   4
   5 // RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
   6 // RUN:   -aux-triple amdgcn-amd-amdhsa -std=c++17 -fgpu-rdc \
   7 // RUN:   -emit-llvm -o - -x hip %s > %t.host
   8
   9 // RUN: cat %t.dev %t.host | FileCheck -check-prefixes=HIP,COMMON %s
  10 // RUN: cat %t.dev %t.host | FileCheck -check-prefixes=COMNEG %s
  11
  12 // RUN: echo "GPU binary" > %t.fatbin
  13
  14 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device -cuid=abc \
  15 // RUN:   -aux-triple x86_64-unknown-linux-gnu -std=c++17 -fgpu-rdc \
  16 // RUN:   -emit-llvm -o - %s > %t.dev
  17
  18 // RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
  19 // RUN:   -aux-triple nvptx -std=c++17 -fgpu-rdc -fcuda-include-gpubinary %t.fatbin \
  20 // RUN:   -emit-llvm -o - %s > %t.host
  21
  22 // RUN: cat %t.dev %t.host | FileCheck -check-prefixes=CUDA,COMMON %s
  23 // RUN: cat %t.dev %t.host | FileCheck -check-prefixes=COMNEG %s
  24
  25 #include "Inputs/cuda.h"
  26
  27 // HIP-DAG: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv\.intern\.b04fd23c98500190]](
  28 // HIP-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT_\.intern\.b04fd23c98500190]](
  29 // HIP-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_1UlvE_EEvT_\.intern\.b04fd23c98500190]](
  30 // HIP-DAG: @[[VM:_ZN12_GLOBAL__N_12vmE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global
  31 // HIP-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE\.static\.b04fd23c98500190]] = addrspace(4) externally_initialized global
  32 // HIP-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global
  33
  34 // CUDA-DAG: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv__intern__b04fd23c98500190]](
  35 // CUDA-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT___intern__b04fd23c98500190]](
  36 // CUDA-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_1UlvE_EEvT___intern__b04fd23c98500190]](
  37 // CUDA-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE__static__b04fd23c98500190]] = addrspace(4) externally_initialized global
  38 // CUDA-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE__static__b04fd23c98500190]] = addrspace(1) externally_initialized global
  39
  40 // COMMON-DAG: @_ZN12_GLOBAL__N_12vdE = internal addrspace(1) global
  41 // COMNEG-NOT: @{{.*}} = {{.*}} c"_ZN12_GLOBAL__N_12vdE{{.*}}\00"
  42
  43 // HIP-DAG: @llvm.compiler.used = {{.*}}@[[VM]]{{.*}}@[[VT]]{{.*}}@[[VC]]
  44 // CUDA-DAG: @llvm.compiler.used = {{.*}}@[[VT]]{{.*}}@[[VC]]
  45
  46 // COMMON-DAG: @[[KERNSTR:.*]] = {{.*}} c"[[KERN]]\00"
  47 // COMMON-DAG: @[[KTXSTR:.*]] = {{.*}} c"[[KTX]]\00"
  48 // COMMON-DAG: @[[KTLSTR:.*]] = {{.*}} c"[[KTL]]\00"
  49 // HIP-DAG: @[[VMSTR:.*]] = {{.*}} c"[[VM]]\00"
  50 // COMMON-DAG: @[[VCSTR:.*]] = {{.*}} c"[[VC]]\00"
  51 // COMMON-DAG: @[[VTSTR:.*]] = {{.*}} c"[[VT]]\00"
  52
  53 // COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KERNSTR]]
  54 // COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KTXSTR]]
  55 // COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KTLSTR]]
  56 // HIP-DAG: call void @__{{.*}}RegisterManagedVar({{.*}}@[[VMSTR]]
  57 // COMMON-DAG: call void @__{{.*}}RegisterVar({{.*}}@[[VCSTR]]
  58 // COMMON-DAG: call void @__{{.*}}RegisterVar({{.*}}@[[VTSTR]]
  59
  60 template <typename T>
  61 __global__ void kt(T x) {}
  62
  63 template <typename T>
  64 __device__ T vt;
  65
  66 namespace {
  67   struct X {};
  68   X x;
  69   auto lambda = [](){};
  70 #if __HIP__
  71   __managed__ int vm = 1;
  72 #endif
  73   __constant__ int vc = 2;
  74
  75   // C should not be externalized since it is used by device code only.
  76   __device__ int vd = 3;
  77   __global__ void kernel() { vd = 4; }
  78 }
  79
  80 template<typename T>
  81 void getSymbol(T *x) {}
  82
  83 void test() {
  84   kernel<<<1, 1>>>();
  85
  86   kt<<<1, 1>>>(x);
  87
  88   kt<<<1, 1>>>(lambda);
  89
  90   // A, B, and tempVar<X> should be externalized since they are
  91   // used by host code.
  92 #if __HIP__
  93   getSymbol(&vm);
  94 #endif
  95   getSymbol(&vc);
  96   getSymbol(&vt<X>);
  97 }