clang/test/CodeGenCUDA/kernel-stub-name.cu

   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 -check-prefixes=CHECK,GNU %s
   6
   7 // RUN: %clang_cc1 -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 -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 -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 ptr @[[CSTUB:__device_stub__ckernel]], align 8
  26 // GNU: @[[HNSKERN:_ZN2ns8nskernelEv]] = constant ptr @[[NSSTUB:_ZN2ns23__device_stub__nskernelEv]], align 8
  27 // GNU: @[[HTKERN:_Z10kernelfuncIiEvv]] = linkonce_odr constant ptr @[[TSTUB:_Z25__device_stub__kernelfuncIiEvv]], comdat, align 8
  28 // GNU: @[[HDKERN:_Z11kernel_declv]] = external constant ptr, align 8
  29 // GNU: @[[HTDKERN:_Z20template_kernel_declIiEvT_]] = external constant ptr, align 8
  30
  31 // MSVC: @[[HCKERN:ckernel]] = dso_local constant ptr @[[CSTUB:__device_stub__ckernel]], align 8
  32 // MSVC: @[[HNSKERN:"\?nskernel@ns@@YAXXZ.*"]] = dso_local constant ptr @[[NSSTUB:"\?__device_stub__nskernel@ns@@YAXXZ"]], align 8
  33 // MSVC: @[[HTKERN:"\?\?\$kernelfunc@H@@YAXXZ.*"]] = linkonce_odr dso_local constant ptr @[[TSTUB:"\?\?\$__device_stub__kernelfunc@H@@YAXXZ.*"]], comdat, align 8
  34 // MSVC: @[[HDKERN:"\?kernel_decl@@YAXXZ.*"]] = external dso_local constant ptr, align 8
  35 // MSVC: @[[HTDKERN:"\?\?\$template_kernel_decl@H@@YAXH.*"]] = external dso_local constant ptr, align 8
  36 extern "C" __global__ void ckernel() {}
  37
  38 namespace ns {
  39 __global__ void nskernel() {}
  40 } // namespace ns
  41
  42 template<class T>
  43 __global__ void kernelfunc() {}
  44
  45 __global__ void kernel_decl();
  46
  47 template<class T>
  48 __global__ void template_kernel_decl(T x);
  49
  50 extern "C" void (*kernel_ptr)();
  51 extern "C" void *void_ptr;
  52
  53 extern "C" void launch(void *kern);
  54
  55 // Device side kernel names
  56
  57 // CHECK: @[[CKERN:[0-9]*]] = {{.*}} c"ckernel\00"
  58 // CHECK: @[[NSKERN:[0-9]*]] = {{.*}} c"_ZN2ns8nskernelEv\00"
  59 // CHECK: @[[TKERN:[0-9]*]] = {{.*}} c"_Z10kernelfuncIiEvv\00"
  60
  61 // Non-template kernel stub functions
  62
  63 // CHECK: define{{.*}}@[[CSTUB]]
  64 // CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HCKERN]]
  65
  66 // CHECK: define{{.*}}@[[NSSTUB]]
  67 // CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HNSKERN]]
  68
  69 // Check kernel stub is called for triple chevron.
  70
  71 // CHECK-LABEL: define{{.*}}@fun1()
  72 // CHECK: call void @[[CSTUB]]()
  73 // CHECK: call void @[[NSSTUB]]()
  74 // CHECK: call void @[[TSTUB]]()
  75 // GNU: call void @[[DSTUB:_Z26__device_stub__kernel_declv]]()
  76 // GNU: call void @[[TDSTUB:_Z35__device_stub__template_kernel_declIiEvT_]](
  77 // MSVC: call void @[[DSTUB:"\?__device_stub__kernel_decl@@YAXXZ"]]()
  78 // MSVC: call void @[[TDSTUB:"\?\?\$__device_stub__template_kernel_decl@H@@YAXH@Z"]](
  79
  80 extern "C" void fun1(void) {
  81   ckernel<<<1, 1>>>();
  82   ns::nskernel<<<1, 1>>>();
  83   kernelfunc<int><<<1, 1>>>();
  84   kernel_decl<<<1, 1>>>();
  85   template_kernel_decl<<<1, 1>>>(1);
  86 }
  87
  88 // Template kernel stub functions
  89
  90 // CHECK: define{{.*}}@[[TSTUB]]
  91 // CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HTKERN]]
  92
  93 // Check declaration of stub function for external kernel.
  94
  95 // CHECK: declare{{.*}}@[[DSTUB]]
  96 // CHECK: declare{{.*}}@[[TDSTUB]]
  97
  98 // Check kernel handle is used for passing the kernel as a function pointer.
  99
 100 // CHECK-LABEL: define{{.*}}@fun2()
 101 // CHECK: call void @launch({{.*}}[[HCKERN]]
 102 // CHECK: call void @launch({{.*}}[[HNSKERN]]
 103 // CHECK: call void @launch({{.*}}[[HTKERN]]
 104 // CHECK: call void @launch({{.*}}[[HDKERN]]
 105 // CHECK: call void @launch({{.*}}[[HTDKERN]]
 106 extern "C" void fun2() {
 107   launch((void *)ckernel);
 108   launch((void *)ns::nskernel);
 109   launch((void *)kernelfunc<int>);
 110   launch((void *)kernel_decl);
 111   launch((void *)template_kernel_decl<int>);
 112 }
 113
 114 // Check kernel handle is used for assigning a kernel to a function pointer.
 115
 116 // CHECK-LABEL: define{{.*}}@fun3()
 117 // CHECK:  store ptr @[[HCKERN]], ptr @kernel_ptr, align 8
 118 // CHECK:  store ptr @[[HCKERN]], ptr @kernel_ptr, align 8
 119 // CHECK:  store ptr @[[HCKERN]], ptr @void_ptr, align 8
 120 // CHECK:  store ptr @[[HCKERN]], ptr @void_ptr, align 8
 121 extern "C" void fun3() {
 122   kernel_ptr = ckernel;
 123   kernel_ptr = &ckernel;
 124   void_ptr = (void *)ckernel;
 125   void_ptr = (void *)&ckernel;
 126 }
 127
 128 // Check kernel stub is loaded from kernel handle when function pointer is
 129 // used with triple chevron.
 130
 131 // CHECK-LABEL: define{{.*}}@fun4()
 132 // CHECK:  store ptr @[[HCKERN]], ptr @kernel_ptr
 133 // CHECK:  call noundef i32 @{{.*hipConfigureCall}}
 134 // CHECK:  %[[HANDLE:.*]] = load ptr, ptr @kernel_ptr, align 8
 135 // CHECK:  %[[STUB:.*]] = load ptr, ptr %[[HANDLE]], align 8
 136 // CHECK:  call void %[[STUB]]()
 137 extern "C" void fun4() {
 138   kernel_ptr = ckernel;
 139   kernel_ptr<<<1,1>>>();
 140 }
 141
 142 // Check kernel handle is passed to a function.
 143
 144 // CHECK-LABEL: define{{.*}}@fun5()
 145 // CHECK:  store ptr @[[HCKERN]], ptr @kernel_ptr
 146 // CHECK:  %[[HANDLE:.*]] = load ptr, ptr @kernel_ptr, align 8
 147 // CHECK:  call void @launch(ptr noundef %[[HANDLE]])
 148 extern "C" void fun5() {
 149   kernel_ptr = ckernel;
 150   launch((void *)kernel_ptr);
 151 }
 152
 153 // Check kernel handle is registered.
 154
 155 // CHECK-LABEL: define{{.*}}@__hip_register_globals
 156 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[HCKERN]]{{.*}}@[[CKERN]]
 157 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[HNSKERN]]{{.*}}@[[NSKERN]]
 158 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[HTKERN]]{{.*}}@[[TKERN]]
 159 // NEG-NOT: call{{.*}}@__hipRegisterFunction{{.*}}__device_stub
 160 // NEG-NOT: call{{.*}}@__hipRegisterFunction{{.*}}kernel_decl
 161 // NEG-NOT: call{{.*}}@__hipRegisterFunction{{.*}}template_kernel_decl