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