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