1 // RUN: echo "GPU binary would be here" > %t
2 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
3 // RUN:     -target-sdk-version=8.0 -fcuda-include-gpubinary %t -o - \
4 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s \
5 // RUN:       --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-OLD
6 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
7 // RUN:     -target-sdk-version=8.0  -fcuda-include-gpubinary %t \
8 // RUN:     -o - -DNOGLOBALS \
9 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s \
10 // RUN:     -check-prefixes=NOGLOBALS,CUDANOGLOBALS
11 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
12 // RUN:     -target-sdk-version=8.0 -fgpu-rdc -fcuda-include-gpubinary %t \
13 // RUN:     -o - \
14 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s \
15 // RUN:       --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-OLD
16 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
17 // RUN:     -target-sdk-version=8.0 -o - \
18 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
19 
20 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s       \
21 // RUN:     -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o - \
22 // RUN:   | FileCheck %s -allow-deprecated-dag-overlap \
23 // RUN:       --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-NEW
24 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
25 // RUN:     -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
26 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s \
27 // RUN:       --check-prefixes=NOGLOBALS,CUDANOGLOBALS
28 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
29 // RUN:     -target-sdk-version=9.2 -fgpu-rdc -fcuda-include-gpubinary %t -o - \
30 // RUN:   | FileCheck %s -allow-deprecated-dag-overlap \
31 // RUN:       --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-NEW
32 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -std=c++17 \
33 // RUN:     -target-sdk-version=9.2 -fgpu-rdc -fcuda-include-gpubinary %t -o - \
34 // RUN:   | FileCheck %s -allow-deprecated-dag-overlap \
35 // RUN:       --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-NEW,LNX_17
36 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
37 // RUN:     -target-sdk-version=9.2 -o - \
38 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
39 
40 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
41 // RUN:     -fcuda-include-gpubinary %t -o - -x hip\
42 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,LNX,NORDC,HIP,HIPEF
43 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
44 // RUN:     -fcuda-include-gpubinary %t -o -  -DNOGLOBALS -x hip \
45 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,HIPNOGLOBALS
46 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
47 // RUN:     -fgpu-rdc -fcuda-include-gpubinary %t -o - -x hip \
48 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,LNX,NORDC,HIP,HIPEF
49 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip\
50 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=ALL,LNX,NORDC,HIP,HIPNEF
51 
52 // RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -aux-triple amdgcn -emit-llvm %s \
53 // RUN:     -fcuda-include-gpubinary %t -o - -x hip\
54 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,WIN
55 
56 #include "Inputs/cuda.h"
57 
58 #ifndef NOGLOBALS
59 // LNX-DAG: @device_var = internal global i32
60 // WIN-DAG: @"?device_var@@3HA" = internal global i32
61 __device__ int device_var;
62 
63 // LNX-DAG: @constant_var = internal global i32
64 // WIN-DAG: @"?constant_var@@3HA" = internal global i32
65 __constant__ int constant_var;
66 
67 // LNX-DAG: @shared_var = internal global i32
68 // WIN-DAG: @"?shared_var@@3HA" = internal global i32
69 __shared__ int shared_var;
70 
71 // Make sure host globals don't get internalized...
72 // LNX-DAG: @host_var = global i32
73 // WIN-DAG: @"?host_var@@3HA" = dso_local global i32
74 int host_var;
75 // ... and that extern vars remain external.
76 // LNX-DAG: @ext_host_var = external global i32
77 // WIN-DAG: @"?ext_host_var@@3HA" = external dso_local global i32
78 extern int ext_host_var;
79 
80 // external device-side variables -> extern references to their shadows.
81 // LNX-DAG: @ext_device_var = external global i32
82 // WIN-DAG: @"?ext_device_var@@3HA" = external dso_local global i32
83 extern __device__ int ext_device_var;
84 // LNX-DAG: @ext_device_var = external global i32
85 // WIN-DAG: @"?ext_constant_var@@3HA" = external dso_local global i32
86 extern __constant__ int ext_constant_var;
87 
88 // external device-side variables with definitions should generate
89 // definitions for the shadows.
90 // LNX-DAG: @ext_device_var_def = internal global i32 undef,
91 // WIN-DAG: @"?ext_device_var_def@@3HA" = internal global i32 undef
92 extern __device__ int ext_device_var_def;
93 __device__ int ext_device_var_def = 1;
94 // LNX-DAG: @ext_device_var_def = internal global i32 undef,
95 // WIN-DAG: @"?ext_constant_var_def@@3HA" = internal global i32 undef
96 __constant__ int ext_constant_var_def = 2;
97 
98 #if __cplusplus > 201402L
99 /// FIXME: Reject __device__ constexpr and inline variables in Sema.
100 // LNX_17: @inline_var = internal global i32 undef, comdat, align 4{{$}}
101 // LNX_17: @_ZN1C17member_inline_varE = internal constant i32 undef, comdat, align 4{{$}}
102 __device__ inline int inline_var = 3;
103 struct C {
104   __device__ static constexpr int member_inline_var = 4;
105 };
106 #endif
107 
use_pointers()108 void use_pointers() {
109   const int *p;
110   p = &device_var;
111   p = &constant_var;
112   p = &shared_var;
113   p = &host_var;
114   p = &ext_device_var;
115   p = &ext_constant_var;
116   p = &ext_host_var;
117 #if __cplusplus > 201402L
118   p = &inline_var;
119   p = &C::member_inline_var;
120 #endif
121 }
122 
123 // Make sure that all parts of GPU code init/cleanup are there:
124 // * constant unnamed string with the device-side kernel name to be passed to
125 //   __hipRegisterFunction/__cudaRegisterFunction.
126 // ALL: @0 = private unnamed_addr constant [18 x i8] c"_Z10kernelfunciii\00"
127 // * constant unnamed string with the device-side kernel name to be passed to
128 //   __hipRegisterVar/__cudaRegisterVar.
129 // ALL: @1 = private unnamed_addr constant [11 x i8] c"device_var\00"
130 // ALL: @2 = private unnamed_addr constant [13 x i8] c"constant_var\00"
131 // ALL: @3 = private unnamed_addr constant [19 x i8] c"ext_device_var_def\00"
132 // ALL: @4 = private unnamed_addr constant [21 x i8] c"ext_constant_var_def\00"
133 // * constant unnamed string with GPU binary
134 // CUDA: @[[FATBIN:.*]] = private constant{{.*GPU binary would be here.*}}\00",
135 // HIPEF: @[[FATBIN:.*]] = private constant{{.*GPU binary would be here.*}}\00",{{.*}}align 4096
136 // HIPNEF: @[[FATBIN:__hip_fatbin]] = external constant i8, section ".hip_fatbin"
137 // CUDANORDC-SAME: section ".nv_fatbin", align 8
138 // CUDARDC-SAME: section "__nv_relfatbin", align 8
139 // * constant struct that wraps GPU binary
140 // ALL: @__[[PREFIX:cuda|hip]]_fatbin_wrapper = internal constant
141 // LNX-SAME: { i32, i32, i8*, i8* }
142 // CUDA-SAME: { i32 1180844977, i32 1,
143 // HIP-SAME: { i32 1212764230, i32 1,
144 // CUDA-SAME: i8* getelementptr inbounds ({{.*}}@[[FATBIN]], i64 0, i64 0),
145 // HIPEF-SAME: i8* getelementptr inbounds ({{.*}}@[[FATBIN]], i64 0, i64 0),
146 // HIPNEF-SAME:  i8* @[[FATBIN]],
147 // LNX-SAME: i8* null }
148 // CUDA-SAME: section ".nvFatBinSegment"
149 // HIP-SAME: section ".hipFatBinSegment"
150 // * variable to save GPU binary handle after initialization
151 // CUDANORDC: @__[[PREFIX]]_gpubin_handle = internal global i8** null
152 // HIPNEF: @__[[PREFIX]]_gpubin_handle = linkonce hidden global i8** null
153 // * constant unnamed string with NVModuleID
154 // RDC: [[MODULE_ID_GLOBAL:@.*]] = private constant
155 // CUDARDC-SAME: c"[[MODULE_ID:.+]]\00", section "__nv_module_id", align 32
156 // * Make sure our constructor was added to global ctor list.
157 // LNX: @llvm.global_ctors = appending global {{.*}}@__[[PREFIX]]_module_ctor
158 // * Alias to global symbol containing the NVModuleID.
159 // RDC: @__fatbinwrap[[MODULE_ID]] = alias { i32, i32, i8*, i8* }
160 // RDC-SAME: { i32, i32, i8*, i8* }* @__[[PREFIX]]_fatbin_wrapper
161 
162 // Test that we build the correct number of calls to cudaSetupArgument followed
163 // by a call to cudaLaunch.
164 
165 // LNX: define{{.*}}kernelfunc
166 
167 // New launch sequence stores arguments into local buffer and passes array of
168 // pointers to them directly to cudaLaunchKernel
169 // CUDA-NEW: alloca
170 // CUDA-NEW: store
171 // CUDA-NEW: store
172 // CUDA-NEW: store
173 // CUDA-NEW: call{{.*}}__cudaPopCallConfiguration
174 // CUDA-NEW: call{{.*}}cudaLaunchKernel
175 
176 // Legacy style launch sequence sets up arguments by passing them to
177 // [cuda|hip]SetupArgument.
178 // CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument
179 // CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument
180 // CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument
181 // CUDA-OLD: call{{.*}}[[PREFIX]]Launch
182 
183 // HIP: call{{.*}}[[PREFIX]]SetupArgument
184 // HIP: call{{.*}}[[PREFIX]]SetupArgument
185 // HIP: call{{.*}}[[PREFIX]]SetupArgument
186 // HIP: call{{.*}}[[PREFIX]]Launch
kernelfunc(int i,int j,int k)187 __global__ void kernelfunc(int i, int j, int k) {}
188 
189 // Test that we've built correct kernel launch sequence.
190 // LNX: define{{.*}}hostfunc
191 // CUDA-OLD: call{{.*}}[[PREFIX]]ConfigureCall
192 // CUDA-NEW: call{{.*}}__cudaPushCallConfiguration
193 // HIP: call{{.*}}[[PREFIX]]ConfigureCall
194 // LNX: call{{.*}}kernelfunc
hostfunc(void)195 void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
196 #endif
197 
198 // Test that we've built a function to register kernels and global vars.
199 // ALL: define internal void @__[[PREFIX]]_register_globals
200 // ALL: call{{.*}}[[PREFIX]]RegisterFunction(i8** %0, {{.*}}kernelfunc{{[^,]*}}, {{[^@]*}}@0
201 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}device_var{{[^,]*}}, {{[^@]*}}@1, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
202 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}constant_var{{[^,]*}}, {{[^@]*}}@2, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
203 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var_def{{[^,]*}}, {{[^@]*}}@3, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
204 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var_def{{[^,]*}}, {{[^@]*}}@4, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
205 // LNX_17-NOT: [[PREFIX]]RegisterVar(i8** %0, {{.*}}inline_var
206 // ALL: ret void
207 
208 // Test that we've built a constructor.
209 // LNX: define internal void @__[[PREFIX]]_module_ctor
210 
211 // In separate mode it calls __[[PREFIX]]RegisterFatBinary(&__[[PREFIX]]_fatbin_wrapper)
212 // HIP only register fat binary once.
213 // HIP: load i8**, i8*** @__hip_gpubin_handle
214 // HIP-NEXT: icmp eq i8** {{.*}}, null
215 // HIP-NEXT: br i1 {{.*}}, label %if, label %exit
216 // HIP: if:
217 // NORDC: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
218 //   .. stores return value in __[[PREFIX]]_gpubin_handle
219 // NORDC-NEXT: store{{.*}}__[[PREFIX]]_gpubin_handle
220 //   .. and then calls __[[PREFIX]]_register_globals
221 // HIP-NEXT: br label %exit
222 // HIP: exit:
223 // HIP-NEXT: load i8**, i8*** @__hip_gpubin_handle
224 // NORDC-NEXT: call void @__[[PREFIX]]_register_globals
225 // * In separate mode we also register a destructor.
226 // NORDC-NEXT: call i32 @atexit(void (i8*)* @__[[PREFIX]]_module_dtor)
227 
228 // With relocatable device code we call __[[PREFIX]]RegisterLinkedBinary%NVModuleID%
229 // RDC: call{{.*}}__[[PREFIX]]RegisterLinkedBinary[[MODULE_ID]](
230 // RDC-SAME: __[[PREFIX]]_register_globals, {{.*}}__[[PREFIX]]_fatbin_wrapper
231 // RDC-SAME: [[MODULE_ID_GLOBAL]]
232 
233 // Test that we've created destructor.
234 // NORDC: define internal void @__[[PREFIX]]_module_dtor
235 // NORDC: load{{.*}}__[[PREFIX]]_gpubin_handle
236 // CUDANORDC-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
237 // HIP-NEXT: icmp ne i8** {{.*}}, null
238 // HIP-NEXT: br i1 {{.*}}, label %if, label %exit
239 // HIP: if:
240 // HIP-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
241 // HIP-NEXT: store i8** null, i8*** @__hip_gpubin_handle
242 // HIP-NEXT: br label %exit
243 // HIP: exit:
244 
245 // There should be no __[[PREFIX]]_register_globals if we have no
246 // device-side globals, but we still need to register GPU binary.
247 // Skip GPU binary string first.
248 // CUDANOGLOBALS-NOT: @{{.*}} = private constant{{.*}}
249 // HIPNOGLOBALS-NOT: @{{.*}} = internal constant{{.*}}
250 // NOGLOBALS-NOT: define internal void @__{{.*}}_register_globals
251 // NOGLOBALS-NOT: define internal void @__[[PREFIX:cuda|hip]]_module_ctor
252 // NOGLOBALS-NOT: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
253 // NOGLOBALS-NOT: call void @__[[PREFIX]]_register_globals
254 // NOGLOBALS-NOT: define internal void @__[[PREFIX]]_module_dtor
255 // NOGLOBALS-NOT: call void @__[[PREFIX]]UnregisterFatBinary
256 
257 // There should be no constructors/destructors if we have no GPU binary.
258 // NOGPUBIN-NOT: define internal void @__[[PREFIX]]_register_globals
259 // NOGPUBIN-NOT: define internal void @__[[PREFIX]]_module_ctor
260 // NOGPUBIN-NOT: define internal void @__[[PREFIX]]_module_dtor
261