[HIP] Remove the assertion on match between host/device names.

[hliao]

Summary

• Under different ABIs, it's obvious that assertion is too strong. Even under the same ABI, once there are unnamed type not required to follow ODR rule, host- and device-side mangling may still get different names. As both the host- and device-side compilation always observe the same AST tree, even with different names, we still could associate the correct pairs, i.e., we don't use (mangled) names to linkage host- and device-side globals. There's no need to have this assertion.

[yaxunl]

LGTM. It seems no reason to assume the mangled name to be same on host and device side once anonymous types are mangled differently in host and device code. On windows, kernel has totally different names on host and device side without issues.

[hliao]

a little explanation of the test case and what's issue is.

[hliao]

clang/test/CodeGenCUDA/unnamed-types.cu

5

device-side mangled name, notice that $_0 refers to the unnamed closure in f1.

24

the registration of host-side stub function to the device-side function name, which is defined in @0. Notice that the host-side stub function has $_1, which refers to the closure in f1 as there's another closure (host-only) in f0.

[hliao]

I took this back. I fab a case where anonymous type IDs mismatch between the device-side name between host-compilation and device-compilation.

clang/lib/CodeGen/CGCUDANV.cpp

if (DeviceMC->shouldMangleDeclName(ND)) {
     DeviceSideName = Out.str();
   } else
     DeviceSideName = ND->getIdentifier()->getName();
   return DeviceSideName;
 }
 
 void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
                                      FunctionArgList &Args) {
-  assert(getDeviceSideName(CGF.CurFuncDecl) == CGF.CurFn->getName() ||
-         getDeviceSideName(CGF.CurFuncDecl) + ".stub" == CGF.CurFn->getName() ||
-         CGF.CGM.getContext().getTargetInfo().getCXXABI() !=
-             CGF.CGM.getContext().getAuxTargetInfo()->getCXXABI());
-
   EmittedKernels.push_back({CGF.CurFn, CGF.CurFuncDecl});
   if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
                          CudaFeature::CUDA_USES_NEW_LAUNCH))
     emitDeviceStubBodyNew(CGF, Args);
   else
     emitDeviceStubBodyLegacy(CGF, Args);
 }
 

clang/test/CodeGenCUDA/unnamed-types.cu

+// RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-linux-gnu -emit-llvm %s -o - | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// CHECK: @0 = private unnamed_addr constant [40 x i8] c"_Z2k0IZZ2f1PfENK3$_0clES0_EUlfE_EvS0_T_\00"
+
+template <typename F>
+__global__ void k0(float *p, F f) {
+  p[0] = f(p[0]);
+}
+
+void f0(float *p) {
+  [](float *p) {
+    *p = 1.f;
+  }(p);
+}
+
+void f1(float *p) {
+  [](float *p) {
+    k0<<<1,1>>>(p, [] __device__ (float x) { return x + 1.f; });
+  }(p);
+}
+// CHECK: @__hip_register_globals
+// CHECK: __hipRegisterFunction{{.*}}_Z2k0IZZ2f1PfENK3$_1clES0_EUlfE_EvS0_T_{{.*}}@0