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.
kernel2()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
test()60 void test() {
61 kernel1<<<1, 1>>>();
62 void *p = (void*)kernel4;
63 use(&var1);
64 }
65
test_lambda_using_extern_shared()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>
template_caller()77 void template_caller() {
78 kernel5<<<1, 1>>>();
79 var1 = 1;
80 }
81
82 template void template_caller<int>();
83