diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -217,11 +217,6 @@ 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)) diff --git a/clang/test/CodeGenCUDA/unnamed-types.cu b/clang/test/CodeGenCUDA/unnamed-types.cu new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenCUDA/unnamed-types.cu @@ -0,0 +1,24 @@ +// 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 +__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