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 = {{.*}} @var2
  28 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @var3
  29 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @ext_shvar
  30 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @shvar
  31 // NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel1v
  32 // NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel4v
  33 // NORDC-NOT: @__clang_gpu_used_external = {{.*}} @var1
  34 // HOST-NEG-NOT: call void @__hipRegisterVar({{.*}}, ptr @ext_shvar
  35 // HOST-NEG-NOT: call void @__hipRegisterVar({{.*}}, ptr @shvar
  36 __global__ void kernel1();
  37
  38 // kernel2 is not marked as used since it is a definition.
  39 __global__ void kernel2() {}
  40
  41 // kernel3 is not marked as used since it is not called by host function.
  42 __global__ void kernel3();
  43
  44 // kernel4 is marked as used even though it is not called.
  45 __global__ void kernel4();
  46
  47 extern __device__ int var1;
  48
  49 __device__ int var2;
  50
  51 extern __device__ int var3;
  52
  53 void use(int *p);
  54
  55 void test() {
  56   kernel1<<<1, 1>>>();
  57   void *p = (void*)kernel4;
  58   use(&var1);
  59 }
  60
  61 __global__ void test_lambda_using_extern_shared() {
  62   extern __shared__ int ext_shvar[];
  63   __shared__ int shvar[10];
  64   auto lambda = [&]() {
  65     ext_shvar[0] = 1;
  66     shvar[0] = 2;
  67   };
  68   lambda();
  69 }