clang/test/CodeGenCUDA/host-used-extern.cu

   1 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
   2 // RUN:   -fgpu-rdc -std=c++11 -emit-llvm -o - -target-cpu gfx906 | FileCheck %s
   3
   4 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
   5 // RUN:   -fgpu-rdc -std=c++11 -emit-llvm -o - -target-cpu gfx906 \
   6 // RUN:   | FileCheck -check-prefix=NEG %s
   7
   8 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
   9 // RUN:   -std=c++11 -emit-llvm -o - -target-cpu gfx906 \
  10 // RUN:   | FileCheck -check-prefixes=NEG,NORDC %s
  11
  12 // RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -x hip %s \
  13 // RUN:   -fgpu-rdc -std=c++11 -emit-llvm -o - \
  14 // RUN:   | FileCheck -check-prefix=HOST-NEG %s
  15
  16
  17 #include "Inputs/cuda.h"
  18
  19 // CHECK-LABEL: @__clang_gpu_used_external = internal {{.*}}global
  20 // CHECK-DAG: @_Z7kernel1v
  21 // CHECK-DAG: @_Z7kernel4v
  22 // CHECK-DAG: @var1
  23 // CHECK-LABEL: @llvm.compiler.used = {{.*}} @__clang_gpu_used_external
  24
  25 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel2v
  26 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel3v
  27 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel5v
  28 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @var2
  29 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @var3
  30 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @ext_shvar
  31 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @shvar
  32 // NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel1v
  33 // NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel4v
  34 // NORDC-NOT: @__clang_gpu_used_external = {{.*}} @var1
  35 // HOST-NEG-NOT: call void @__hipRegisterVar({{.*}}, ptr @ext_shvar
  36 // HOST-NEG-NOT: call void @__hipRegisterVar({{.*}}, ptr @shvar
  37 __global__ void kernel1();
  38
  39 // kernel2 is not marked as used since it is a definition.
  40 __global__ void kernel2() {}
  41
  42 // kernel3 is not marked as used since it is not called by host function.
  43 __global__ void kernel3();
  44
  45 // kernel4 is marked as used even though it is not called.
  46 __global__ void kernel4();
  47
  48 // kernel5 is not marked as used since it is called by host function
  49 // with weak_odr linkage, which may be dropped by linker.
  50 __global__ void kernel5();
  51
  52 extern __device__ int var1;
  53
  54 __device__ int var2;
  55
  56 extern __device__ int var3;
  57
  58 void use(int *p);
  59
  60 void test() {
  61   kernel1<<<1, 1>>>();
  62   void *p = (void*)kernel4;
  63   use(&var1);
  64 }
  65
  66 __global__ void test_lambda_using_extern_shared() {
  67   extern __shared__ int ext_shvar[];
  68   __shared__ int shvar[10];
  69   auto lambda = [&]() {
  70     ext_shvar[0] = 1;
  71     shvar[0] = 2;
  72   };
  73   lambda();
  74 }
  75
  76 template<class T>
  77 void template_caller() {
  78   kernel5<<<1, 1>>>();
  79   var1 = 1;
  80 }
  81
  82 template void template_caller<int>();