Skip to content

Commit 4874ff0

Browse files
niconico
committed
Revert "[hip][cuda] Enable extended lambda support on Windows."
This reverts commit a2fdf9d. Slightly speculative, seeing several cuda tests fail on this Windows bot: http://45.33.8.238/win/32620/step_7.txt
1 parent 26ca503 commit 4874ff0

18 files changed

+17
-131
lines changed

‎clang/include/clang/AST/ASTContext.h

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -538,9 +538,6 @@ class ASTContext : public RefCountedBase<ASTContext> {
538538
/// need them (like static local vars).
539539
llvm::MapVector<const NamedDecl *, unsigned> MangleNumbers;
540540
llvm::MapVector<const VarDecl *, unsigned> StaticLocalNumbers;
541-
/// Mapping the associated device lambda mangling number if present.
542-
mutable llvm::DenseMap<const CXXRecordDecl *, unsigned>
543-
DeviceLambdaManglingNumbers;
544541

545542
/// Mapping that stores parameterIndex values for ParmVarDecls when
546543
/// that value exceeds the bitfield size of ParmVarDeclBits.ParameterIndex.

‎clang/include/clang/AST/DeclCXX.h

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1735,12 +1735,6 @@ class CXXRecordDecl : public RecordDecl {
17351735
getLambdaData().HasKnownInternalLinkage = HasKnownInternalLinkage;
17361736
}
17371737

1738-
/// Set the device side mangling number.
1739-
void setDeviceLambdaManglingNumber(unsigned Num) const;
1740-
1741-
/// Retrieve the device side mangling number.
1742-
unsigned getDeviceLambdaManglingNumber() const;
1743-
17441738
/// Returns the inheritance model used for this record.
17451739
MSInheritanceModel getMSInheritanceModel() const;
17461740

‎clang/include/clang/AST/Mangle.h

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -107,9 +107,6 @@ class MangleContext {
107107
virtual bool shouldMangleCXXName(const NamedDecl *D) = 0;
108108
virtual bool shouldMangleStringLiteral(const StringLiteral *SL) = 0;
109109

110-
virtual bool isDeviceMangleContext() const { return false; }
111-
virtual void setDeviceMangleContext(bool) {}
112-
113110
// FIXME: consider replacing raw_ostream & with something like SmallString &.
114111
void mangleName(GlobalDecl GD, raw_ostream &);
115112
virtual void mangleCXXName(GlobalDecl GD, raw_ostream &) = 0;

‎clang/include/clang/AST/MangleNumberingContext.h

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -52,11 +52,6 @@ class MangleNumberingContext {
5252
/// this context.
5353
virtual unsigned getManglingNumber(const TagDecl *TD,
5454
unsigned MSLocalManglingNumber) = 0;
55-
56-
/// Retrieve the mangling number of a new lambda expression with the
57-
/// given call operator within the device context. No device number is
58-
/// assigned if there's no device numbering context is associated.
59-
virtual unsigned getDeviceManglingNumber(const CXXMethodDecl *) { return 0; }
6055
};
6156

6257
} // end namespace clang

‎clang/include/clang/Sema/Sema.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6558,7 +6558,7 @@ class Sema final {
65586558
/// Number lambda for linkage purposes if necessary.
65596559
void handleLambdaNumbering(
65606560
CXXRecordDecl *Class, CXXMethodDecl *Method,
6561-
Optional<std::tuple<bool, unsigned, unsigned, Decl *>> Mangling = None);
6561+
Optional<std::tuple<unsigned, bool, Decl *>> Mangling = None);
65626562

65636563
/// Endow the lambda scope info with the relevant properties.
65646564
void buildLambdaScope(sema::LambdaScopeInfo *LSI,

‎clang/lib/AST/ASTImporter.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2848,8 +2848,6 @@ ExpectedDecl ASTNodeImporter::VisitRecordDecl(RecordDecl *D) {
28482848
return CDeclOrErr.takeError();
28492849
D2CXX->setLambdaMangling(DCXX->getLambdaManglingNumber(), *CDeclOrErr,
28502850
DCXX->hasKnownLambdaInternalLinkage());
2851-
D2CXX->setDeviceLambdaManglingNumber(
2852-
DCXX->getDeviceLambdaManglingNumber());
28532851
} else if (DCXX->isInjectedClassName()) {
28542852
// We have to be careful to do a similar dance to the one in
28552853
// Sema::ActOnStartCXXMemberDeclarations

‎clang/lib/AST/CXXABI.h

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -22,9 +22,8 @@ class ASTContext;
2222
class CXXConstructorDecl;
2323
class DeclaratorDecl;
2424
class Expr;
25-
class MangleContext;
26-
class MangleNumberingContext;
2725
class MemberPointerType;
26+
class MangleNumberingContext;
2827

2928
/// Implements C++ ABI-specific semantic analysis functions.
3029
class CXXABI {
@@ -76,8 +75,6 @@ class CXXABI {
7675
/// Creates an instance of a C++ ABI class.
7776
CXXABI *CreateItaniumCXXABI(ASTContext &Ctx);
7877
CXXABI *CreateMicrosoftCXXABI(ASTContext &Ctx);
79-
std::unique_ptr<MangleNumberingContext>
80-
createItaniumNumberingContext(MangleContext *);
8178
}
8279

8380
#endif

‎clang/lib/AST/DeclCXX.cpp

Lines changed: 0 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1593,20 +1593,6 @@ Decl *CXXRecordDecl::getLambdaContextDecl() const {
15931593
return getLambdaData().ContextDecl.get(Source);
15941594
}
15951595

1596-
void CXXRecordDecl::setDeviceLambdaManglingNumber(unsigned Num) const {
1597-
assert(isLambda() && "Not a lambda closure type!");
1598-
if (Num)
1599-
getASTContext().DeviceLambdaManglingNumbers[this] = Num;
1600-
}
1601-
1602-
unsigned CXXRecordDecl::getDeviceLambdaManglingNumber() const {
1603-
assert(isLambda() && "Not a lambda closure type!");
1604-
auto I = getASTContext().DeviceLambdaManglingNumbers.find(this);
1605-
if (I != getASTContext().DeviceLambdaManglingNumbers.end())
1606-
return I->second;
1607-
return 0;
1608-
}
1609-
16101596
static CanQualType GetConversionType(ASTContext &Context, NamedDecl *Conv) {
16111597
QualType T =
16121598
cast<CXXConversionDecl>(Conv->getUnderlyingDecl()->getAsFunction())

‎clang/lib/AST/ItaniumCXXABI.cpp

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -258,9 +258,3 @@ class ItaniumCXXABI : public CXXABI {
258258
CXXABI *clang::CreateItaniumCXXABI(ASTContext &Ctx) {
259259
return new ItaniumCXXABI(Ctx);
260260
}
261-
262-
std::unique_ptr<MangleNumberingContext>
263-
clang::createItaniumNumberingContext(MangleContext *Mangler) {
264-
return std::make_unique<ItaniumNumberingContext>(
265-
cast<ItaniumMangleContext>(Mangler));
266-
}

‎clang/lib/AST/ItaniumMangle.cpp

Lines changed: 1 addition & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -125,8 +125,6 @@ class ItaniumMangleContextImpl : public ItaniumMangleContext {
125125
llvm::DenseMap<DiscriminatorKeyTy, unsigned> Discriminator;
126126
llvm::DenseMap<const NamedDecl*, unsigned> Uniquifier;
127127

128-
bool IsDevCtx = false;
129-
130128
public:
131129
explicit ItaniumMangleContextImpl(ASTContext &Context,
132130
DiagnosticsEngine &Diags)
@@ -139,10 +137,6 @@ class ItaniumMangleContextImpl : public ItaniumMangleContext {
139137
bool shouldMangleStringLiteral(const StringLiteral *) override {
140138
return false;
141139
}
142-
143-
bool isDeviceMangleContext() const override { return IsDevCtx; }
144-
void setDeviceMangleContext(bool IsDev) override { IsDevCtx = IsDev; }
145-
146140
void mangleCXXName(GlobalDecl GD, raw_ostream &) override;
147141
void mangleThunk(const CXXMethodDecl *MD, const ThunkInfo &Thunk,
148142
raw_ostream &) override;
@@ -1882,15 +1876,7 @@ void CXXNameMangler::mangleLambda(const CXXRecordDecl *Lambda) {
18821876
// (in lexical order) with that same <lambda-sig> and context.
18831877
//
18841878
// The AST keeps track of the number for us.
1885-
//
1886-
// In CUDA/HIP, to ensure the consistent lamba numbering between the device-
1887-
// and host-side compilations, an extra device mangle context may be created
1888-
// if the host-side CXX ABI has different numbering for lambda. In such case,
1889-
// if the mangle context is that device-side one, use the device-side lambda
1890-
// mangling number for this lambda.
1891-
unsigned Number = Context.isDeviceMangleContext()
1892-
? Lambda->getDeviceLambdaManglingNumber()
1893-
: Lambda->getLambdaManglingNumber();
1879+
unsigned Number = Lambda->getLambdaManglingNumber();
18941880
assert(Number > 0 && "Lambda should be mangled as an unnamed class");
18951881
if (Number > 1)
18961882
mangleNumber(Number - 2);

‎clang/lib/AST/MicrosoftCXXABI.cpp

Lines changed: 2 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,6 @@
1616
#include "clang/AST/Attr.h"
1717
#include "clang/AST/CXXInheritance.h"
1818
#include "clang/AST/DeclCXX.h"
19-
#include "clang/AST/Mangle.h"
2019
#include "clang/AST/MangleNumberingContext.h"
2120
#include "clang/AST/RecordLayout.h"
2221
#include "clang/AST/Type.h"
@@ -65,19 +64,6 @@ class MicrosoftNumberingContext : public MangleNumberingContext {
6564
}
6665
};
6766

68-
class MSHIPNumberingContext : public MicrosoftNumberingContext {
69-
std::unique_ptr<MangleNumberingContext> DeviceCtx;
70-
71-
public:
72-
MSHIPNumberingContext(MangleContext *Mangler) {
73-
DeviceCtx = createItaniumNumberingContext(Mangler);
74-
}
75-
76-
unsigned getDeviceManglingNumber(const CXXMethodDecl *CallOperator) override {
77-
return DeviceCtx->getManglingNumber(CallOperator);
78-
}
79-
};
80-
8167
class MicrosoftCXXABI : public CXXABI {
8268
ASTContext &Context;
8369
llvm::SmallDenseMap<CXXRecordDecl *, CXXConstructorDecl *> RecordToCopyCtor;
@@ -87,19 +73,8 @@ class MicrosoftCXXABI : public CXXABI {
8773
llvm::SmallDenseMap<TagDecl *, TypedefNameDecl *>
8874
UnnamedTagDeclToTypedefNameDecl;
8975

90-
// MangleContext for device numbering context, which is based on Itanium C++
91-
// ABI.
92-
std::unique_ptr<MangleContext> Mangler;
93-
9476
public:
95-
MicrosoftCXXABI(ASTContext &Ctx) : Context(Ctx) {
96-
if (Context.getLangOpts().CUDA) {
97-
assert(Context.getTargetInfo().getCXXABI().isMicrosoft() &&
98-
Context.getAuxTargetInfo()->getCXXABI().isItaniumFamily() &&
99-
"Unexpected combination of C++ ABIs.");
100-
Mangler.reset(Context.createMangleContext(Context.getAuxTargetInfo()));
101-
}
102-
}
77+
MicrosoftCXXABI(ASTContext &Ctx) : Context(Ctx) { }
10378

10479
MemberPointerInfo
10580
getMemberPointerInfo(const MemberPointerType *MPT) const override;
@@ -158,8 +133,6 @@ class MicrosoftCXXABI : public CXXABI {
158133

159134
std::unique_ptr<MangleNumberingContext>
160135
createMangleNumberingContext() const override {
161-
if (Context.getLangOpts().CUDA)
162-
return std::make_unique<MSHIPNumberingContext>(Mangler.get());
163136
return std::make_unique<MicrosoftNumberingContext>();
164137
}
165138
};
@@ -293,3 +266,4 @@ CXXABI::MemberPointerInfo MicrosoftCXXABI::getMemberPointerInfo(
293266
CXXABI *clang::CreateMicrosoftCXXABI(ASTContext &Ctx) {
294267
return new MicrosoftCXXABI(Ctx);
295268
}
269+

‎clang/lib/CodeGen/CGCUDANV.cpp

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -190,12 +190,6 @@ CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
190190
CharPtrTy = llvm::PointerType::getUnqual(Types.ConvertType(Ctx.CharTy));
191191
VoidPtrTy = cast<llvm::PointerType>(Types.ConvertType(Ctx.VoidPtrTy));
192192
VoidPtrPtrTy = VoidPtrTy->getPointerTo();
193-
// If the host and device have different C++ ABIs, mark it as the device
194-
// mangle context so that the mangling needs to retrieve the additonal device
195-
// lambda mangling number instead of the regular host one.
196-
DeviceMC->setDeviceMangleContext(
197-
CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() &&
198-
CGM.getContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily());
199193
}
200194

201195
llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const {

‎clang/lib/Sema/SemaLambda.cpp

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -429,16 +429,15 @@ CXXMethodDecl *Sema::startLambdaDefinition(CXXRecordDecl *Class,
429429

430430
void Sema::handleLambdaNumbering(
431431
CXXRecordDecl *Class, CXXMethodDecl *Method,
432-
Optional<std::tuple<bool, unsigned, unsigned, Decl *>> Mangling) {
432+
Optional<std::tuple<unsigned, bool, Decl *>> Mangling) {
433433
if (Mangling) {
434+
unsigned ManglingNumber;
434435
bool HasKnownInternalLinkage;
435-
unsigned ManglingNumber, DeviceManglingNumber;
436436
Decl *ManglingContextDecl;
437-
std::tie(HasKnownInternalLinkage, ManglingNumber, DeviceManglingNumber,
438-
ManglingContextDecl) = Mangling.getValue();
437+
std::tie(ManglingNumber, HasKnownInternalLinkage, ManglingContextDecl) =
438+
Mangling.getValue();
439439
Class->setLambdaMangling(ManglingNumber, ManglingContextDecl,
440440
HasKnownInternalLinkage);
441-
Class->setDeviceLambdaManglingNumber(DeviceManglingNumber);
442441
return;
443442
}
444443

@@ -474,7 +473,6 @@ void Sema::handleLambdaNumbering(
474473
unsigned ManglingNumber = MCtx->getManglingNumber(Method);
475474
Class->setLambdaMangling(ManglingNumber, ManglingContextDecl,
476475
HasKnownInternalLinkage);
477-
Class->setDeviceLambdaManglingNumber(MCtx->getDeviceManglingNumber(Method));
478476
}
479477
}
480478

‎clang/lib/Sema/TreeTransform.h

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -12505,11 +12505,10 @@ TreeTransform<Derived>::TransformLambdaExpr(LambdaExpr *E) {
1250512505
E->getCaptureDefault());
1250612506
getDerived().transformedLocalDecl(OldClass, {Class});
1250712507

12508-
Optional<std::tuple<bool, unsigned, unsigned, Decl *>> Mangling;
12508+
Optional<std::tuple<unsigned, bool, Decl *>> Mangling;
1250912509
if (getDerived().ReplacingOriginal())
12510-
Mangling = std::make_tuple(OldClass->hasKnownLambdaInternalLinkage(),
12511-
OldClass->getLambdaManglingNumber(),
12512-
OldClass->getDeviceLambdaManglingNumber(),
12510+
Mangling = std::make_tuple(OldClass->getLambdaManglingNumber(),
12511+
OldClass->hasKnownLambdaInternalLinkage(),
1251312512
OldClass->getLambdaContextDecl());
1251412513

1251512514
// Build the call operator.

‎clang/lib/Serialization/ASTReaderDecl.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1748,7 +1748,6 @@ void ASTDeclReader::ReadCXXDefinitionData(
17481748
Lambda.NumExplicitCaptures = Record.readInt();
17491749
Lambda.HasKnownInternalLinkage = Record.readInt();
17501750
Lambda.ManglingNumber = Record.readInt();
1751-
D->setDeviceLambdaManglingNumber(Record.readInt());
17521751
Lambda.ContextDecl = readDeclID();
17531752
Lambda.Captures = (Capture *)Reader.getContext().Allocate(
17541753
sizeof(Capture) * Lambda.NumCaptures);

‎clang/lib/Serialization/ASTWriter.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5663,7 +5663,6 @@ void ASTRecordWriter::AddCXXDefinitionData(const CXXRecordDecl *D) {
56635663
Record->push_back(Lambda.NumExplicitCaptures);
56645664
Record->push_back(Lambda.HasKnownInternalLinkage);
56655665
Record->push_back(Lambda.ManglingNumber);
5666-
Record->push_back(D->getDeviceLambdaManglingNumber());
56675666
AddDeclRef(D->getLambdaContextDecl());
56685667
AddTypeSourceInfo(Lambda.MethodTyInfo);
56695668
for (unsigned I = 0, N = Lambda.NumCaptures; I != N; ++I) {

‎clang/test/CodeGenCUDA/ms-linker-options.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2,12 +2,12 @@
22
// RUN: -fno-autolink -triple amdgcn-amd-amdhsa \
33
// RUN: | FileCheck -check-prefix=DEV %s
44
// RUN: %clang_cc1 -emit-llvm -o - -fms-extensions -x hip %s -triple \
5-
// RUN: x86_64-pc-windows-msvc -aux-triple amdgcn | FileCheck -check-prefix=HOST %s
5+
// RUN: x86_64-pc-windows-msvc | FileCheck -check-prefix=HOST %s
66
// RUN: %clang_cc1 -emit-llvm -o - -fcuda-is-device -fms-extensions %s \
77
// RUN: -fno-autolink -triple amdgcn-amd-amdhsa \
88
// RUN: | FileCheck -check-prefix=DEV %s
99
// RUN: %clang_cc1 -emit-llvm -o - -fms-extensions %s -triple \
10-
// RUN: x86_64-pc-windows-msvc -aux-triple amdgcn | FileCheck -check-prefix=HOST %s
10+
// RUN: x86_64-pc-windows-msvc | FileCheck -check-prefix=HOST %s
1111

1212
// DEV-NOT: llvm.linker.options
1313
// DEV-NOT: llvm.dependent-libraries

‎clang/test/CodeGenCUDA/unnamed-types.cu

Lines changed: 3 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -1,39 +1,24 @@
11
// RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
2-
// RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-pc-windows-msvc -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=MSVC
32
// RUN: %clang_cc1 -std=c++11 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
43

54
#include "Inputs/cuda.h"
65

76
// HOST: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1
8-
// HOST: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
9-
// Check that, on MSVC, the same device kernel mangling name is generated.
10-
// MSVC: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1
11-
// MSVC: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
127

138
__device__ float d0(float x) {
14-
return [](float x) { return x + 1.f; }(x);
9+
return [](float x) { return x + 2.f; }(x);
1510
}
1611

1712
__device__ float d1(float x) {
1813
return [](float x) { return x * 2.f; }(x);
1914
}
2015

2116
// DEVICE: amdgpu_kernel void @_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_(
22-
// DEVICE: define internal float @_ZZZ2f1PfENKUlS_E_clES_ENKUlfE_clEf(
2317
template <typename F>
2418
__global__ void k0(float *p, F f) {
2519
p[0] = f(p[0]) + d0(p[1]) + d1(p[2]);
2620
}
2721

28-
// DEVICE: amdgpu_kernel void @_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_(
29-
// DEVICE: define internal float @_ZZ2f1PfENKUlfE_clEf(
30-
// DEVICE: define internal float @_ZZ2f1PfENKUlffE_clEff(
31-
// DEVICE: define internal float @_ZZ2f1PfENKUlfE0_clEf(
32-
template <typename F0, typename F1, typename F2>
33-
__global__ void k1(float *p, F0 f0, F1 f1, F2 f2) {
34-
p[0] = f0(p[0]) + f1(p[1], p[2]) + f2(p[3]);
35-
}
36-
3722
void f0(float *p) {
3823
[](float *p) {
3924
*p = 1.f;
@@ -44,17 +29,11 @@ void f0(float *p) {
4429
// linkages are still required to keep the original `internal` linkage.
4530

4631
// HOST: define internal void @_ZZ2f1PfENKUlS_E_clES_(
32+
// DEVICE: define internal float @_ZZZ2f1PfENKUlS_E_clES_ENKUlfE_clEf(
4733
void f1(float *p) {
4834
[](float *p) {
49-
k0<<<1,1>>>(p, [] __device__ (float x) { return x + 3.f; });
35+
k0<<<1,1>>>(p, [] __device__ (float x) { return x + 1.f; });
5036
}(p);
51-
k1<<<1,1>>>(p,
52-
[] __device__ (float x) { return x + 4.f; },
53-
[] __device__ (float x, float y) { return x * y; },
54-
[] __device__ (float x) { return x + 5.f; });
5537
}
5638
// HOST: @__hip_register_globals
5739
// HOST: __hipRegisterFunction{{.*}}@_Z17__device_stub__k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_{{.*}}@0
58-
// HOST: __hipRegisterFunction{{.*}}@_Z17__device_stub__k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_{{.*}}@1
59-
// MSVC: __hipRegisterFunction{{.*}}@"??$k0@V<lambda_1>@?0???R1?0??f1@@YAXPEAM@Z@QEBA@0@Z@@@YAXPEAMV<lambda_1>@?0???R0?0??f1@@YAX0@Z@QEBA@0@Z@@Z{{.*}}@0
60-
// MSVC: __hipRegisterFunction{{.*}}@"??$k1@V<lambda_2>@?0??f1@@YAXPEAM@Z@V<lambda_3>@?0??2@YAX0@Z@V<lambda_4>@?0??2@YAX0@Z@@@YAXPEAMV<lambda_2>@?0??f1@@YAX0@Z@V<lambda_3>@?0??1@YAX0@Z@V<lambda_4>@?0??1@YAX0@Z@@Z{{.*}}@1

0 commit comments

Comments
 (0)