1 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -x hip %s -o - \
2 // RUN: | FileCheck %s
3
4 #define __global__ __attribute__((global))
5 // CHECK: @_Z4kern7TempValIjE = constant ptr @_Z19__device_stub__kern7TempValIjE, align 8
6 // CHECK: @0 = private unnamed_addr constant [19 x i8] c"_Z4kern7TempValIjE\00", align 1
7 template <typename type>
8 struct TempVal {
9 type value;
10 };
11
12 __global__ void kern(TempVal<unsigned int> in_val);
13
main(int argc,char ** argv)14 int main(int argc, char ** argv) {
15 auto* fptr = &(kern);
16 // CHECK: store ptr @_Z4kern7TempValIjE, ptr %fptr, align 8
17 return 0;
18 }
19 // CHECK: define dso_local void @_Z19__device_stub__kern7TempValIjE(i32 %in_val.coerce) #1 {
20 // CHECK: %2 = call i32 @hipLaunchByPtr(ptr @_Z4kern7TempValIjE)
21
22 // CHECK: define internal void @__hip_register_globals(ptr %0) {
23 // CHECK: %1 = call i32 @__hipRegisterFunction(ptr %0, ptr @_Z4kern7TempValIjE, ptr @0, ptr @0, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
24
kern(TempVal<unsigned int> in_val)25 __global__ void kern(TempVal<unsigned int> in_val) {
26 }
27
28