[CUDA][HIP] Allow comdat for kernels

[yaxunl]

Summary

Two identical instantiations of a template function can be emitted by two TU's
with linkonce_odr linkage without causing duplicate symbols in linker. MSVC
also requires these symbols be in comdat sections. Linux does not require
the symbols in comdat sections to be merged by linker but by default
clang puts them in comdat sections.

If a template kernel is instantiated identically in two TU's. MSVC requires
that them to be in comdat sections, otherwise MSVC linker will diagnose them as
duplicate symbols. However, currently clang does not put instantiated template
kernels in comdat sections, which causes link error for MSVC.

This patch allows putting instantiated template kernels into comdat sections.

[yaxunl]

ping

[tra]

As phrased, the summary would likely be rather confusing for anyone other than you and me.

Currently Visual Studio 2019 has a linker issue which causes linking error
when a template kernel is instantiated in different compilation units.

It's not clear what exactly is the issue and what causes it.

On the other hand, it is unnecessary to prefix kernel stub for MSVC
target since the host and device compilation uses different mangling
ABI.

This could use more details on why different mangling matters here. IIRC, on Linux where both host and device use the same mangling and HIP needed a way to tell apart the GPU-side kernels and their host-side stub. Different mangling makes it a non-issue.

This patch let clang not emit kernel handle for MSVC target to work around the linker issue.

Again, without the back-story the jump from linking error to mangling differences to "let's not emit a handle" does not make much sense.

I'd restructure it along the line of:

• we emit host-side handles to match GPU-side kernels
• the handles cause linking issues on windows because of X/Y/Z.
• handles are not necessary on Windows, because of the different host/device mangling
• not generating the handles avoids the linking issue on Windows.

This prompts the question -- should/could handle generation be improved instead? Having identical behavior on all platforms would arguably be better than a platform-specific workaround.

[yaxunl]

In D112492#3101090, @tra wrote:

As phrased, the summary would likely be rather confusing for anyone other than you and me.

Currently Visual Studio 2019 has a linker issue which causes linking error
when a template kernel is instantiated in different compilation units.

It's not clear what exactly is the issue and what causes it.

Currently, the identical instantiation of a template kernel in different TU causes linking error about duplicate symbols on Windows.

In the beginning, I thought it was due to a bug in MSVC linker. However, further investigation shows that it is not a MSVC linker bug, but a bug of clang.

Basically, it is not sufficient to set linkonce_odr linkage to let MSVC linker merge symbols. The symbol needs to be in comdat sections. This is not required on Linux. However, it would be more consistent to always put kernel stubs and kernel handles with linkonce_odr linkage into comdat if the target supports it.

I will update this patch to fix the comdat for kernel stub and kernel handle.

On the other hand, it is unnecessary to prefix kernel stub for MSVC
target since the host and device compilation uses different mangling
ABI.

This could use more details on why different mangling matters here. IIRC, on Linux where both host and device use the same mangling and HIP needed a way to tell apart the GPU-side kernels and their host-side stub. Different mangling makes it a non-issue.

This patch let clang not emit kernel handle for MSVC target to work around the linker issue.

Again, without the back-story the jump from linking error to mangling differences to "let's not emit a handle" does not make much sense.

I'd restructure it along the line of:

• we emit host-side handles to match GPU-side kernels
• the handles cause linking issues on windows because of X/Y/Z.
• handles are not necessary on Windows, because of the different host/device mangling
• not generating the handles avoids the linking issue on Windows.

This prompts the question -- should/could handle generation be improved instead? Having identical behavior on all platforms would arguably be better than a platform-specific workaround.

With the fix of comdat issue, we should be able to use a consistent kernel launching mechanism for Linux and Windows. Since the debugger requests kernel stub to have a different name than the kernel, we have to let the kernel stub and kernel handle have different names. That mechanism will stay unchanged.

I also found that MSVC name mangling currently does not add device_stub prefix to the mangled name of kernel stubs. I will update this patch to fix that.

With these changes, we should have consistent name mangling for kernel stubs and kernel launching mechanism on Linux and Windows.

[tra]

With these changes, we should have consistent name mangling for kernel stubs and kernel launching mechanism on Linux and Windows.

Nice! Thank you for figuring out the root causes.

[yaxunl]

fix comdat instead

[tra]

clang/lib/CodeGen/CodeGenModule.cpp

4290–4293

This was added in D63277 specifically to deal with comdat-related issue on windows.

We do need to have unique addresses for each kernel stub. Placing stubs into comdat may allow them to be merged into one and that would be a problem.
https://docs.microsoft.com/en-us/cpp/build/reference/opt-optimizations?view=vs-2019

@rnk,@kpyzhov -- how do we make sure that identical functions are placed in comdat (needed for templates) but do not get comdat-folded with /OPT:ICF?

[rnk]

clang/lib/CodeGen/CodeGenModule.cpp

4290–4293

These are readonly global variables. I believe MSVC's ICF implementation changed behavior here a few times over the years. I think the current state is that they don't merge such globals. If you are using LLD, you can use safe ICF to avoid merging such globals.

@zequanwu reviewed the state of ICF in LLD in the last year, so he may know more about the current status, and he may be more helpful if you have more questions.

Finally, you could mark these globals as mutable if you really want to block ICF. That will work reliably with any linker.

[rnk]

clang/lib/CodeGen/CodeGenModule.cpp

4290–4293

Broadly, I think this change is correct. If I had reviewed D63277, I probably would have objected to it and not approved it.

[yaxunl]

I think probably it is necessary to merge linkonce_odr symbols for them to work properly.

Consider the following testcase:

// a.cu
template<typename T>
__global__ void foo(T x) {}

void test1() {
    foo<<<1,1>>>(1);
}

// b.cu
template<typename T>
__global__ void foo(T x) {}

void test2() {
    foo<<<1,1>>>(1);
}

// c.cu
template<typename T>
__global__ void foo(T x);

int main() {
    foo<<<1,1>>>(1);
}

Assume a.cu, b.cu, and c.cu are compiled with default -fno-gpu-rdc option and linked together.

Both a.obj and b.obj contain a global symbol foo<int> as the kernel stub function. c.obj contains reference to foo<int>, so it has to resolve to foo<int> in a.obj or b.obj. It only makes sense for linker to merge foo<int> in a.obj and b.obj and let c.obj resolve to the merged symbol. This also requires that the fat binary embedded in a.obj and b.obj must contain the identical definition of kernel foo<int>. That is, if ODR is followed, even though there are two fat binaries containing kernel foo<int>, only one of them will be used (it is fine since they are identical), which corresponds to the merged symbol for the kernel stub foo<int>.

The implication is that, we have to ask users to follow ODR even with the default -fno-gpu-rdc option. And users cannot have different definitions for the same template instantiation (e.g. foo<int>) in different TU's, otherwise there will be UB.

Considering ODR is a fundamental assumption for C++, I think it is justifiable to request users to follow that no matter whether -fgpu-rdc or -fno-gpu-rdc.

[tra]

Yes, we do need to merge identical functions with identical names for templates.

The comdat-folding issue is different. IIUIC, it allows merging two functions with identical code and different names, into one function with two names. That will break CUDA as we do need to have each stub to have a unique address as we use it to find the matching GPU-side kernel.

[rnk]

In D112492#3119892, @tra wrote:

Yes, we do need to merge identical functions with identical names for templates.

The comdat-folding issue is different. IIUIC, it allows merging two functions with identical code and different names, into one function with two names. That will break CUDA as we do need to have each stub to have a unique address as we use it to find the matching GPU-side kernel.

Well, yes, ICF breaks function pointer identity. There's no way around that, and it is documented:
https://docs.microsoft.com/en-us/cpp/build/reference/opt-optimizations?view=msvc-170
CUDA users will have to remove /OPT:ICF from their linker flags.

Maybe you could make this work by embedding an ICF-breaking device into all the stubs. Something like a volatile asm blob that takes the current function as an argument and puts it in a register.

[yaxunl]

clang/lib/CodeGen/CodeGenModule.cpp

4290–4293

These symbols are for kernel stub functions. How do I mark them as mutable?

[yaxunl]

I did an experiment regarding the ICF issue and it seems not to affect kernel stub.

#include "hip/hip_runtime.h"

template<typename T>
void bar(T x) { }

template<typename T>
__global__ void foo(T x) {}

int main() {
  foo<<<1,1>>>(1);
  printf("%p\n", foo<int>);
  printf("%p\n", foo<float>);
  printf("%p\n", bar<int>);
  printf("%p\n", bar<float>);
}

If I pass -Wl,/opt:noicf, I got

00007FF622A01100
00007FF622A01170
00007FF622A01360
00007FF622A01370

By default, I got

00007FF693521100
00007FF693521170
00007FF693521360
00007FF693521360

This indicates bar<int> and bar<float> are folded but kernel stubs are not folded.

I also tried -Wl,/opt:icf=10, and kernel stubs are still not folded.

For HIP, since the kernel stub passes a unique kernel symbol to the internal kernel launching API, you may think the kernel stubs are not folded because they are not identical.

To imitate the CUDA case, where the address of kernel stub function itself is passed to the internal kernel launching API, I used the original patch of this review, where the kernel stub function passes the address of itself to the internal kernel launching API, therefore in a sense, the kernel stubs are all the same. Still, the kernel stubs are not folded.

Looking at the assembly of the kernel stub function:

; foo<int>
.seh_proc "??$foo@H@@YAXH@Z"
# %bb.0:
        pushq   %rsi
        .seh_pushreg %rsi
        pushq   %rdi
        .seh_pushreg %rdi
        subq    $120, %rsp
        .seh_stackalloc 120
        .seh_endprologue
        movl    %ecx, 60(%rsp)
        leaq    60(%rsp), %rax
        movq    %rax, 64(%rsp)
        leaq    104(%rsp), %rsi
        leaq    88(%rsp), %rdi
        leaq    80(%rsp), %r8
        leaq    72(%rsp), %r9
        movq    %rsi, %rcx
        movq    %rdi, %rdx
        callq   __hipPopCallConfiguration
        movq    80(%rsp), %rax
        movq    72(%rsp), %rcx
        movq    %rcx, 40(%rsp)
        movq    %rax, 32(%rsp)
        leaq    "??$foo@H@@YAXH@Z"(%rip), %rcx
        leaq    64(%rsp), %r9
        movq    %rsi, %rdx
        movq    %rdi, %r8
        callq   hipLaunchKernel
        nop
        addq    $120, %rsp
        popq    %rdi
        popq    %rsi
        retq
        .seh_endproc

; foo<float>
.seh_proc "??$foo@M@@YAXM@Z"
# %bb.0:
        pushq   %rsi
        .seh_pushreg %rsi
        pushq   %rdi
        .seh_pushreg %rdi
        subq    $120, %rsp
        .seh_stackalloc 120
        .seh_endprologue
        movss   %xmm0, 60(%rsp)
        leaq    60(%rsp), %rax
        movq    %rax, 64(%rsp)
        leaq    104(%rsp), %rsi
        leaq    88(%rsp), %rdi
        leaq    80(%rsp), %r8
        leaq    72(%rsp), %r9
        movq    %rsi, %rcx
        movq    %rdi, %rdx
        callq   __hipPopCallConfiguration
        movq    80(%rsp), %rax
        movq    72(%rsp), %rcx
        movq    %rcx, 40(%rsp)
        movq    %rax, 32(%rsp)
        leaq    "??$foo@M@@YAXM@Z"(%rip), %rcx
        leaq    64(%rsp), %r9
        movq    %rsi, %rdx
        movq    %rdi, %r8
        callq   hipLaunchKernel
        nop
        addq    $120, %rsp
        popq    %rdi
        popq    %rsi
        retq
        .seh_endproc

I think they are not folded because link.exe is smart enough to treat them as not identical comdat functions. I think we may stop worrying about the ICF foading kernel stubs.

[rnk]

I think the key is the self-reference in the LEA instruction:

; foo<int>
.seh_proc "??$foo@H@@YAXH@Z"
...
        leaq    "??$foo@H@@YAXH@Z"(%rip), %rcx
...
; foo<float>
.seh_proc "??$foo@M@@YAXM@Z"
...
        leaq    "??$foo@M@@YAXM@Z"(%rip), %rcx

I think they are not folded because link.exe is smart enough to treat them as not identical comdat functions. I think we may stop worrying about the ICF foading kernel stubs.

It sounds like the behavior may have changed since D63277 was landed, and maybe we don't need it anymore.

clang/lib/CodeGen/CGCUDANV.cpp

llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F,
   auto *Var = new llvm::GlobalVariable(
       TheModule, F->getType(), /*isConstant=*/true, F->getLinkage(),
       /*Initializer=*/nullptr,
       CGM.getMangledName(
           GD.getWithKernelReferenceKind(KernelReferenceKind::Kernel)));
   Var->setAlignment(CGM.getPointerAlign().getAsAlign());
   Var->setDSOLocal(F->isDSOLocal());
   Var->setVisibility(F->getVisibility());
+  CGM.maybeSetTrivialComdat(*GD.getDecl(), *Var);
   KernelHandles[F] = Var;
   KernelStubs[Var] = F;
   return Var;
 }

clang/lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::MaybeHandleStaticInExternC(const SomeDecl *D,
   if (!R.second)
     R.first->second = nullptr;
 }
 
 static bool shouldBeInCOMDAT(CodeGenModule &CGM, const Decl &D) {
   if (!CGM.supportsCOMDAT())
     return false;
 
-  // Do not set COMDAT attribute for CUDA/HIP stub functions to prevent
-  // them being "merged" by the COMDAT Folding linker optimization.
-  if (D.hasAttr<CUDAGlobalAttr>())
-    return false;
-
   if (D.hasAttr<SelectAnyAttr>())
     return true;
 
   GVALinkage Linkage;
   if (auto *VD = dyn_cast<VarDecl>(&D))
     Linkage = CGM.getContext().GetGVALinkageForVariable(VD);
   else
     Linkage = CGM.getContext().GetGVALinkageForFunction(cast<FunctionDecl>(&D));

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

 // RUN: echo "GPU binary would be here" > %t
 
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
 // RUN:     -fcuda-include-gpubinary %t -o - -x hip\
-// RUN:   | FileCheck %s
+// RUN:   | FileCheck -check-prefixes=CHECK,GNU %s
+
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN:     -fcuda-include-gpubinary %t -o - -x hip\
+// RUN:   | FileCheck -check-prefix=NEG %s
+
+// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -emit-llvm %s \
+// RUN:     -aux-triple amdgcn-amd-amdhsa -fcuda-include-gpubinary \
+// RUN:     %t -o - -x hip\
+// RUN:   | FileCheck -check-prefixes=CHECK,MSVC %s
+
+// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -emit-llvm %s \
+// RUN:     -aux-triple amdgcn-amd-amdhsa -fcuda-include-gpubinary \
+// RUN:     %t -o - -x hip\
+// RUN:   | FileCheck -check-prefix=NEG %s
 
 #include "Inputs/cuda.h"
 
-// Kernel handles
+// Check kernel handles are emitted for non-MSVC target but not for MSVC target.
 
-// CHECK: @[[HCKERN:ckernel]] = constant void ()* @__device_stub__ckernel, align 8
-// CHECK: @[[HNSKERN:_ZN2ns8nskernelEv]] = constant void ()* @_ZN2ns23__device_stub__nskernelEv, align 8
-// CHECK: @[[HTKERN:_Z10kernelfuncIiEvv]] = linkonce_odr constant void ()* @_Z25__device_stub__kernelfuncIiEvv, align 8
-// CHECK: @[[HDKERN:_Z11kernel_declv]] = external constant void ()*, align 8
+// GNU: @[[HCKERN:ckernel]] = constant void ()* @[[CSTUB:__device_stub__ckernel]], align 8
+// GNU: @[[HNSKERN:_ZN2ns8nskernelEv]] = constant void ()* @[[NSSTUB:_ZN2ns23__device_stub__nskernelEv]], align 8
+// GNU: @[[HTKERN:_Z10kernelfuncIiEvv]] = linkonce_odr constant void ()* @[[TSTUB:_Z25__device_stub__kernelfuncIiEvv]], comdat, align 8
+// GNU: @[[HDKERN:_Z11kernel_declv]] = external constant void ()*, align 8
+
+// MSVC: @[[HCKERN:ckernel]] = dso_local constant void ()* @[[CSTUB:__device_stub__ckernel]], align 8
+// MSVC: @[[HNSKERN:"\?nskernel@ns@@YAXXZ.*"]] = dso_local constant void ()* @[[NSSTUB:"\?nskernel@ns@@YAXXZ"]], align 8
+// MSVC: @[[HTKERN:"\?\?\$kernelfunc@H@@YAXXZ.*"]] = linkonce_odr dso_local constant void ()* @[[TSTUB:"\?\?\$kernelfunc@H@@YAXXZ.*"]], comdat, align 8
+// MSVC: @[[HDKERN:"\?kernel_decl@@YAXXZ.*"]] = external dso_local constant void ()*, align 8
 
 extern "C" __global__ void ckernel() {}
 
 namespace ns {
 __global__ void nskernel() {}
 } // namespace ns
 
 template<class T>
 __global__ void kernelfunc() {}
 
 __global__ void kernel_decl();
 
-void (*kernel_ptr)();
-void *void_ptr;
+extern "C" void (*kernel_ptr)();
+extern "C" void *void_ptr;
 
-void launch(void *kern);
+extern "C" void launch(void *kern);
 
 // Device side kernel names
 
 // CHECK: @[[CKERN:[0-9]*]] = {{.*}} c"ckernel\00"
 // CHECK: @[[NSKERN:[0-9]*]] = {{.*}} c"_ZN2ns8nskernelEv\00"
 // CHECK: @[[TKERN:[0-9]*]] = {{.*}} c"_Z10kernelfuncIiEvv\00"
 
 // Non-template kernel stub functions
 
-// CHECK: define{{.*}}@[[CSTUB:__device_stub__ckernel]]
+// CHECK: define{{.*}}@[[CSTUB]]
 // CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HCKERN]]
-// CHECK: define{{.*}}@[[NSSTUB:_ZN2ns23__device_stub__nskernelEv]]
-// CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HNSKERN]]
 
+// CHECK: define{{.*}}@[[NSSTUB]]
+// CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HNSKERN]]
 
-// Check kernel stub is used for triple chevron
+// Check kernel stub is called for triple chevron.
 
-// CHECK-LABEL: define{{.*}}@_Z4fun1v()
+// CHECK-LABEL: define{{.*}}@fun1()
 // CHECK: call void @[[CSTUB]]()
 // CHECK: call void @[[NSSTUB]]()
-// CHECK: call void @[[TSTUB:_Z25__device_stub__kernelfuncIiEvv]]()
-// CHECK: call void @[[DSTUB:_Z26__device_stub__kernel_declv]]()
+// CHECK: call void @[[TSTUB]]()
+// GNU: call void @[[DSTUB:_Z26__device_stub__kernel_declv]]()
+// MSVC: call void @[[DSTUB:"\?kernel_decl@@YAXXZ"]]()
 
-void fun1(void) {
+extern "C" void fun1(void) {
   ckernel<<<1, 1>>>();
   ns::nskernel<<<1, 1>>>();
   kernelfunc<int><<<1, 1>>>();
   kernel_decl<<<1, 1>>>();
 }
 
 // Template kernel stub functions
 
 // CHECK: define{{.*}}@[[TSTUB]]
 // CHECK: call{{.*}}@hipLaunchByPtr{{.*}}@[[HTKERN]]
 
 // Check declaration of stub function for external kernel.
 
 // CHECK: declare{{.*}}@[[DSTUB]]
 
-// Check kernel handle is used for passing the kernel as a function pointer
+// Check kernel handle is used for passing the kernel as a function pointer.
 
-// CHECK-LABEL: define{{.*}}@_Z4fun2v()
-// CHECK: call void @_Z6launchPv({{.*}}[[HCKERN]]
-// CHECK: call void @_Z6launchPv({{.*}}[[HNSKERN]]
-// CHECK: call void @_Z6launchPv({{.*}}[[HTKERN]]
-// CHECK: call void @_Z6launchPv({{.*}}[[HDKERN]]
-void fun2() {
+// CHECK-LABEL: define{{.*}}@fun2()
+// CHECK: call void @launch({{.*}}[[HCKERN]]
+// CHECK: call void @launch({{.*}}[[HNSKERN]]
+// CHECK: call void @launch({{.*}}[[HTKERN]]
+// CHECK: call void @launch({{.*}}[[HDKERN]]
+extern "C" void fun2() {
   launch((void *)ckernel);
   launch((void *)ns::nskernel);
   launch((void *)kernelfunc<int>);
   launch((void *)kernel_decl);
 }
 
-// Check kernel handle is used for assigning a kernel to a function pointer
+// Check kernel handle is used for assigning a kernel to a function pointer.
 
-// CHECK-LABEL: define{{.*}}@_Z4fun3v()
+// CHECK-LABEL: define{{.*}}@fun3()
 // CHECK:  store void ()* bitcast (void ()** @[[HCKERN]] to void ()*), void ()** @kernel_ptr, align 8
 // CHECK:  store void ()* bitcast (void ()** @[[HCKERN]] to void ()*), void ()** @kernel_ptr, align 8
 // CHECK:  store i8* bitcast (void ()** @[[HCKERN]] to i8*), i8** @void_ptr, align 8
 // CHECK:  store i8* bitcast (void ()** @[[HCKERN]] to i8*), i8** @void_ptr, align 8
-void fun3() {
+extern "C" void fun3() {
   kernel_ptr = ckernel;
   kernel_ptr = &ckernel;
   void_ptr = (void *)ckernel;
   void_ptr = (void *)&ckernel;
 }
 
 // Check kernel stub is loaded from kernel handle when function pointer is
-// used with triple chevron
+// used with triple chevron.
 
-// CHECK-LABEL: define{{.*}}@_Z4fun4v()
+// CHECK-LABEL: define{{.*}}@fun4()
 // CHECK:  store void ()* bitcast (void ()** @[[HCKERN]] to void ()*), void ()** @kernel_ptr
-// CHECK:  call i32 @_Z16hipConfigureCall4dim3S_mP9hipStream
+// CHECK:  call i32 @{{.*hipConfigureCall}}
 // CHECK:  %[[HANDLE:.*]] = load void ()*, void ()** @kernel_ptr, align 8
 // CHECK:  %[[CAST:.*]] = bitcast void ()* %[[HANDLE]] to void ()**
 // CHECK:  %[[STUB:.*]] = load void ()*, void ()** %[[CAST]], align 8
 // CHECK:  call void %[[STUB]]()
-void fun4() {
+extern "C" void fun4() {
   kernel_ptr = ckernel;
   kernel_ptr<<<1,1>>>();
 }
 
-// Check kernel handle is passed to a function
+// Check kernel handle is passed to a function.
 
-// CHECK-LABEL: define{{.*}}@_Z4fun5v()
+// CHECK-LABEL: define{{.*}}@fun5()
 // CHECK:  store void ()* bitcast (void ()** @[[HCKERN]] to void ()*), void ()** @kernel_ptr
 // CHECK:  %[[HANDLE:.*]] = load void ()*, void ()** @kernel_ptr, align 8
 // CHECK:  %[[CAST:.*]] = bitcast void ()* %[[HANDLE]] to i8*
-// CHECK:  call void @_Z6launchPv(i8* %[[CAST]])
-void fun5() {
+// CHECK:  call void @launch(i8* %[[CAST]])
+extern "C" void fun5() {
   kernel_ptr = ckernel;
   launch((void *)kernel_ptr);
 }
 
+// Check kernel handle is registered.
+
 // CHECK-LABEL: define{{.*}}@__hip_register_globals
 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[HCKERN]]{{.*}}@[[CKERN]]
 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[HNSKERN]]{{.*}}@[[NSKERN]]
 // CHECK: call{{.*}}@__hipRegisterFunction{{.*}}@[[HTKERN]]{{.*}}@[[TKERN]]
-// CHECK-NOT: call{{.*}}@__hipRegisterFunction{{.*}}@[[HDKERN]]{{.*}}@{{[0-9]*}}
+// NEG-NOT: call{{.*}}@__hipRegisterFunction{{.*}}__device_stub
+// NEG-NOT: call{{.*}}@__hipRegisterFunction{{.*}}kernel_decl

clang/test/CodeGenCUDA/usual-deallocators.cu

__host__ __device__ void tests_hd(void *t) {
   // HOST:   call void @_ZN6H1H2D2dlEPv(i8* {{.*}})
   test_hd<H1H2D2>(t);
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv
   // COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}})
   test_hd<H1H2D1D2>(t);
 }
 
 // Make sure that we've generated the kernel used by A::~A.
-// DEVICE-LABEL: define dso_local void @_Z1fIiEvT_
+// DEVICE-LABEL: define void @_Z1fIiEvT_
 
 // Make sure we've picked deallocator for the correct side of compilation.
 
 // COMMON-LABEL: define  linkonce_odr void @_ZN4H1D1dlEPv(i8* %0)
 // DEVICE: call void @dev_fn()
 // HOST:   call void @host_fn()
 
 // DEVICE-LABEL: define  linkonce_odr void @_ZN4H1D2dlEPvj(i8* %0, i32 %1)