1 // RUN: echo "GPU binary would be here" > %t 2 3 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \ 4 // RUN: -fcuda-include-gpubinary %t -o - -x hip\ 5 // RUN: | FileCheck %s 6 7 #include "Inputs/cuda.h" 8 9 // Kernel handles 10 11 // CHECK: @[[HCKERN:ckernel]] = constant void ()* @__device_stub__ckernel, align 8 12 // CHECK: @[[HNSKERN:_ZN2ns8nskernelEv]] = constant void ()* @_ZN2ns23__device_stub__nskernelEv, align 8 13 // CHECK: @[[HTKERN:_Z10kernelfuncIiEvv]] = linkonce_odr constant void ()* @_Z25__device_stub__kernelfuncIiEvv, align 8 14 // CHECK: @[[HDKERN:_Z11kernel_declv]] = external constant void ()*, align 8 15 16 extern "C" __global__ void ckernel() {} 17 18 namespace ns { 19 __global__ void nskernel() {} 20 } // namespace ns 21 22 template<class T> 23 __global__ void kernelfunc() {} 24 25 __global__ void kernel_decl(); 26 27 void (*kernel_ptr)(); 28 void *void_ptr; 29 30 void launch(void *kern); 31 32 // Device side kernel names 33 34 // CHECK: @[[CKERN:[0-9]*]] = {{.*}} c"ckernel\00" 35 // CHECK: @[[NSKERN:[0-9]*]] = {{.*}} c"_ZN2ns8nskernelEv\00" 36 // CHECK: @[[TKERN:[0-9]*]] = {{.*}} c"_Z10kernelfuncIiEvv\00" 37 38 // Non-template kernel stub functions 39 40 // CHECK: define{{.*}}@[[CSTUB:__device_stub__ckernel]] 41 // CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HCKERN]] 42 // CHECK: define{{.*}}@[[NSSTUB:_ZN2ns23__device_stub__nskernelEv]] 43 // CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HNSKERN]] 44 45 46 // Check kernel stub is used for triple chevron 47 48 // CHECK-LABEL: define{{.*}}@_Z4fun1v() 49 // CHECK: call void @[[CSTUB]]() 50 // CHECK: call void @[[NSSTUB]]() 51 // CHECK: call void @[[TSTUB:_Z25__device_stub__kernelfuncIiEvv]]() 52 // CHECK: call void @[[DSTUB:_Z26__device_stub__kernel_declv]]() 53 54 void fun1(void) { 55 ckernel<<<1, 1>>>(); 56 ns::nskernel<<<1, 1>>>(); 57 kernelfunc<int><<<1, 1>>>(); 58 kernel_decl<<<1, 1>>>(); 59 } 60 61 // Template kernel stub functions 62 63 // CHECK: define{{.*}}@[[TSTUB]] 64 // CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HTKERN]] 65 66 // Check declaration of stub function for external kernel. 67 68 // CHECK: declare{{.*}}@[[DSTUB]] 69 70 // Check kernel handle is used for passing the kernel as a function pointer 71 72 // CHECK-LABEL: define{{.*}}@_Z4fun2v() 73 // CHECK: call void @_Z6launchPv({{.*}}[[HCKERN]] 74 // CHECK: call void @_Z6launchPv({{.*}}[[HNSKERN]] 75 // CHECK: call void @_Z6launchPv({{.*}}[[HTKERN]] 76 // CHECK: call void @_Z6launchPv({{.*}}[[HDKERN]] 77 void fun2() { 78 launch((void *)ckernel); 79 launch((void *)ns::nskernel); 80 launch((void *)kernelfunc<int>); 81 launch((void *)kernel_decl); 82 } 83 84 // Check kernel handle is used for assigning a kernel to a function pointer 85 86 // CHECK-LABEL: define{{.*}}@_Z4fun3v() 87 // CHECK: store void ()* bitcast (void ()** @[[HCKERN]] to void ()*), void ()** @kernel_ptr, align 8 88 // CHECK: store void ()* bitcast (void ()** @[[HCKERN]] to void ()*), void ()** @kernel_ptr, align 8 89 // CHECK: store i8* bitcast (void ()** @[[HCKERN]] to i8*), i8** @void_ptr, align 8 90 // CHECK: store i8* bitcast (void ()** @[[HCKERN]] to i8*), i8** @void_ptr, align 8 91 void fun3() { 92 kernel_ptr = ckernel; 93 kernel_ptr = &ckernel; 94 void_ptr = (void *)ckernel; 95 void_ptr = (void *)&ckernel; 96 } 97 98 // Check kernel stub is loaded from kernel handle when function pointer is 99 // used with triple chevron 100 101 // CHECK-LABEL: define{{.*}}@_Z4fun4v() 102 // CHECK: store void ()* bitcast (void ()** @[[HCKERN]] to void ()*), void ()** @kernel_ptr 103 // CHECK: call i32 @_Z16hipConfigureCall4dim3S_mP9hipStream 104 // CHECK: %[[HANDLE:.*]] = load void ()*, void ()** @kernel_ptr, align 8 105 // CHECK: %[[CAST:.*]] = bitcast void ()* %[[HANDLE]] to void ()** 106 // CHECK: %[[STUB:.*]] = load void ()*, void ()** %[[CAST]], align 8 107 // CHECK: call void %[[STUB]]() 108 void fun4() { 109 kernel_ptr = ckernel; 110 kernel_ptr<<<1,1>>>(); 111 } 112 113 // Check kernel handle is passed to a function 114 115 // CHECK-LABEL: define{{.*}}@_Z4fun5v() 116 // CHECK: store void ()* bitcast (void ()** @[[HCKERN]] to void ()*), void ()** @kernel_ptr 117 // CHECK: %[[HANDLE:.*]] = load void ()*, void ()** @kernel_ptr, align 8 118 // CHECK: %[[CAST:.*]] = bitcast void ()* %[[HANDLE]] to i8* 119 // CHECK: call void @_Z6launchPv(i8* %[[CAST]]) 120 void fun5() { 121 kernel_ptr = ckernel; 122 launch((void *)kernel_ptr); 123 } 124 125 // CHECK-LABEL: define{{.*}}@__hip_register_globals 126 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[HCKERN]]{{.*}}@[[CKERN]] 127 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[HNSKERN]]{{.*}}@[[NSKERN]] 128 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[HTKERN]]{{.*}}@[[TKERN]] 129 // CHECK-NOT: call{{.*}}@__hipRegisterFunction{{.*}}@[[HDKERN]]{{.*}}@{{[0-9]*}} 130