[llvm-branch-commits] [clang] 5ae2b97 - Recommit of a2fdf9d4d734732a6fa9288f1ffdf12bf8618123.
Tom Stellard via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Feb 15 17:11:24 PST 2021
Author: Michael Liao
Date: 2021-02-15T17:10:57-08:00
New Revision: 5ae2b9726f27b571de71542ef4d59ebceee8aca7
URL: https://github.com/llvm/llvm-project/commit/5ae2b9726f27b571de71542ef4d59ebceee8aca7
DIFF: https://github.com/llvm/llvm-project/commit/5ae2b9726f27b571de71542ef4d59ebceee8aca7.diff
LOG: Recommit of a2fdf9d4d734732a6fa9288f1ffdf12bf8618123.
- The failures are all cc1-based tests due to the missing `-aux-triple` options,
which is always prepared by the driver in CUDA/HIP compilation.
- Add extra check on the missing aux-targetinfo to prevent crashing.
[hip][cuda] Enable extended lambda support on Windows.
- On Windows, extended lambda has extra issues due to the numbering
schemes are different between the host compilation (Microsoft C++ ABI)
and the device compilation (Itanium C++ ABI. Additional device side
lambda number is required per lambda for the host compilation to
correctly mangle the device-side lambda name.
- A hybrid numbering context `MSHIPNumberingContext` is introduced to
number a lambda for both host- and device-compilations.
Reviewed By: rnk
Differential Revision: https://reviews.llvm.org/D69322
This reverts commit 4874ff02417916cc9ff994b34abcb5e563056546.
(cherry picked from commit 01bf529db2cf465b029e29e537807576bfcbc452)
Added:
Modified:
clang/include/clang/AST/ASTContext.h
clang/include/clang/AST/DeclCXX.h
clang/include/clang/AST/Mangle.h
clang/include/clang/AST/MangleNumberingContext.h
clang/include/clang/Sema/Sema.h
clang/lib/AST/ASTImporter.cpp
clang/lib/AST/CXXABI.h
clang/lib/AST/DeclCXX.cpp
clang/lib/AST/ItaniumCXXABI.cpp
clang/lib/AST/ItaniumMangle.cpp
clang/lib/AST/MicrosoftCXXABI.cpp
clang/lib/CodeGen/CGCUDANV.cpp
clang/lib/Sema/SemaLambda.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReaderDecl.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/CodeGenCUDA/unnamed-types.cu
Removed:
################################################################################
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index ce47d54e44b0..ae69a68608b7 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -538,6 +538,9 @@ class ASTContext : public RefCountedBase<ASTContext> {
/// need them (like static local vars).
llvm::MapVector<const NamedDecl *, unsigned> MangleNumbers;
llvm::MapVector<const VarDecl *, unsigned> StaticLocalNumbers;
+ /// Mapping the associated device lambda mangling number if present.
+ mutable llvm::DenseMap<const CXXRecordDecl *, unsigned>
+ DeviceLambdaManglingNumbers;
/// Mapping that stores parameterIndex values for ParmVarDecls when
/// that value exceeds the bitfield size of ParmVarDeclBits.ParameterIndex.
diff --git a/clang/include/clang/AST/DeclCXX.h b/clang/include/clang/AST/DeclCXX.h
index e32101bb2276..89006b1cfa7f 100644
--- a/clang/include/clang/AST/DeclCXX.h
+++ b/clang/include/clang/AST/DeclCXX.h
@@ -1735,6 +1735,12 @@ class CXXRecordDecl : public RecordDecl {
getLambdaData().HasKnownInternalLinkage = HasKnownInternalLinkage;
}
+ /// Set the device side mangling number.
+ void setDeviceLambdaManglingNumber(unsigned Num) const;
+
+ /// Retrieve the device side mangling number.
+ unsigned getDeviceLambdaManglingNumber() const;
+
/// Returns the inheritance model used for this record.
MSInheritanceModel getMSInheritanceModel() const;
diff --git a/clang/include/clang/AST/Mangle.h b/clang/include/clang/AST/Mangle.h
index 0e8d6dd53d8a..7b6495d85eb6 100644
--- a/clang/include/clang/AST/Mangle.h
+++ b/clang/include/clang/AST/Mangle.h
@@ -96,6 +96,9 @@ class MangleContext {
virtual bool shouldMangleCXXName(const NamedDecl *D) = 0;
virtual bool shouldMangleStringLiteral(const StringLiteral *SL) = 0;
+ virtual bool isDeviceMangleContext() const { return false; }
+ virtual void setDeviceMangleContext(bool) {}
+
// FIXME: consider replacing raw_ostream & with something like SmallString &.
void mangleName(GlobalDecl GD, raw_ostream &);
virtual void mangleCXXName(GlobalDecl GD, raw_ostream &) = 0;
diff --git a/clang/include/clang/AST/MangleNumberingContext.h b/clang/include/clang/AST/MangleNumberingContext.h
index f1ca6a05dbaf..eb33759682d6 100644
--- a/clang/include/clang/AST/MangleNumberingContext.h
+++ b/clang/include/clang/AST/MangleNumberingContext.h
@@ -52,6 +52,11 @@ class MangleNumberingContext {
/// this context.
virtual unsigned getManglingNumber(const TagDecl *TD,
unsigned MSLocalManglingNumber) = 0;
+
+ /// Retrieve the mangling number of a new lambda expression with the
+ /// given call operator within the device context. No device number is
+ /// assigned if there's no device numbering context is associated.
+ virtual unsigned getDeviceManglingNumber(const CXXMethodDecl *) { return 0; }
};
} // end namespace clang
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 42814f6ba8f6..2530a2776373 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -6558,7 +6558,7 @@ class Sema final {
/// Number lambda for linkage purposes if necessary.
void handleLambdaNumbering(
CXXRecordDecl *Class, CXXMethodDecl *Method,
- Optional<std::tuple<unsigned, bool, Decl *>> Mangling = None);
+ Optional<std::tuple<bool, unsigned, unsigned, Decl *>> Mangling = None);
/// Endow the lambda scope info with the relevant properties.
void buildLambdaScope(sema::LambdaScopeInfo *LSI,
diff --git a/clang/lib/AST/ASTImporter.cpp b/clang/lib/AST/ASTImporter.cpp
index 085c50c0667b..0d723fbbcd8c 100644
--- a/clang/lib/AST/ASTImporter.cpp
+++ b/clang/lib/AST/ASTImporter.cpp
@@ -2848,6 +2848,8 @@ ExpectedDecl ASTNodeImporter::VisitRecordDecl(RecordDecl *D) {
return CDeclOrErr.takeError();
D2CXX->setLambdaMangling(DCXX->getLambdaManglingNumber(), *CDeclOrErr,
DCXX->hasKnownLambdaInternalLinkage());
+ D2CXX->setDeviceLambdaManglingNumber(
+ DCXX->getDeviceLambdaManglingNumber());
} else if (DCXX->isInjectedClassName()) {
// We have to be careful to do a similar dance to the one in
// Sema::ActOnStartCXXMemberDeclarations
diff --git a/clang/lib/AST/CXXABI.h b/clang/lib/AST/CXXABI.h
index 31cb36918726..ca9424bcb7a4 100644
--- a/clang/lib/AST/CXXABI.h
+++ b/clang/lib/AST/CXXABI.h
@@ -22,8 +22,9 @@ class ASTContext;
class CXXConstructorDecl;
class DeclaratorDecl;
class Expr;
-class MemberPointerType;
+class MangleContext;
class MangleNumberingContext;
+class MemberPointerType;
/// Implements C++ ABI-specific semantic analysis functions.
class CXXABI {
@@ -75,6 +76,8 @@ class CXXABI {
/// Creates an instance of a C++ ABI class.
CXXABI *CreateItaniumCXXABI(ASTContext &Ctx);
CXXABI *CreateMicrosoftCXXABI(ASTContext &Ctx);
+std::unique_ptr<MangleNumberingContext>
+createItaniumNumberingContext(MangleContext *);
}
#endif
diff --git a/clang/lib/AST/DeclCXX.cpp b/clang/lib/AST/DeclCXX.cpp
index 0368ada0b81c..0375f9b4432e 100644
--- a/clang/lib/AST/DeclCXX.cpp
+++ b/clang/lib/AST/DeclCXX.cpp
@@ -1593,6 +1593,20 @@ Decl *CXXRecordDecl::getLambdaContextDecl() const {
return getLambdaData().ContextDecl.get(Source);
}
+void CXXRecordDecl::setDeviceLambdaManglingNumber(unsigned Num) const {
+ assert(isLambda() && "Not a lambda closure type!");
+ if (Num)
+ getASTContext().DeviceLambdaManglingNumbers[this] = Num;
+}
+
+unsigned CXXRecordDecl::getDeviceLambdaManglingNumber() const {
+ assert(isLambda() && "Not a lambda closure type!");
+ auto I = getASTContext().DeviceLambdaManglingNumbers.find(this);
+ if (I != getASTContext().DeviceLambdaManglingNumbers.end())
+ return I->second;
+ return 0;
+}
+
static CanQualType GetConversionType(ASTContext &Context, NamedDecl *Conv) {
QualType T =
cast<CXXConversionDecl>(Conv->getUnderlyingDecl()->getAsFunction())
diff --git a/clang/lib/AST/ItaniumCXXABI.cpp b/clang/lib/AST/ItaniumCXXABI.cpp
index 069add8464ae..be10258a2d77 100644
--- a/clang/lib/AST/ItaniumCXXABI.cpp
+++ b/clang/lib/AST/ItaniumCXXABI.cpp
@@ -258,3 +258,9 @@ class ItaniumCXXABI : public CXXABI {
CXXABI *clang::CreateItaniumCXXABI(ASTContext &Ctx) {
return new ItaniumCXXABI(Ctx);
}
+
+std::unique_ptr<MangleNumberingContext>
+clang::createItaniumNumberingContext(MangleContext *Mangler) {
+ return std::make_unique<ItaniumNumberingContext>(
+ cast<ItaniumMangleContext>(Mangler));
+}
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index 4420f6a2c1c3..5cad84a96845 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -125,6 +125,8 @@ class ItaniumMangleContextImpl : public ItaniumMangleContext {
llvm::DenseMap<DiscriminatorKeyTy, unsigned> Discriminator;
llvm::DenseMap<const NamedDecl*, unsigned> Uniquifier;
+ bool IsDevCtx = false;
+
public:
explicit ItaniumMangleContextImpl(ASTContext &Context,
DiagnosticsEngine &Diags)
@@ -137,6 +139,10 @@ class ItaniumMangleContextImpl : public ItaniumMangleContext {
bool shouldMangleStringLiteral(const StringLiteral *) override {
return false;
}
+
+ bool isDeviceMangleContext() const override { return IsDevCtx; }
+ void setDeviceMangleContext(bool IsDev) override { IsDevCtx = IsDev; }
+
void mangleCXXName(GlobalDecl GD, raw_ostream &) override;
void mangleThunk(const CXXMethodDecl *MD, const ThunkInfo &Thunk,
raw_ostream &) override;
@@ -1846,7 +1852,15 @@ void CXXNameMangler::mangleLambda(const CXXRecordDecl *Lambda) {
// (in lexical order) with that same <lambda-sig> and context.
//
// The AST keeps track of the number for us.
- unsigned Number = Lambda->getLambdaManglingNumber();
+ //
+ // In CUDA/HIP, to ensure the consistent lamba numbering between the device-
+ // and host-side compilations, an extra device mangle context may be created
+ // if the host-side CXX ABI has
diff erent numbering for lambda. In such case,
+ // if the mangle context is that device-side one, use the device-side lambda
+ // mangling number for this lambda.
+ unsigned Number = Context.isDeviceMangleContext()
+ ? Lambda->getDeviceLambdaManglingNumber()
+ : Lambda->getLambdaManglingNumber();
assert(Number > 0 && "Lambda should be mangled as an unnamed class");
if (Number > 1)
mangleNumber(Number - 2);
diff --git a/clang/lib/AST/MicrosoftCXXABI.cpp b/clang/lib/AST/MicrosoftCXXABI.cpp
index f9f9fe985b6f..166aa3b3bd60 100644
--- a/clang/lib/AST/MicrosoftCXXABI.cpp
+++ b/clang/lib/AST/MicrosoftCXXABI.cpp
@@ -16,6 +16,7 @@
#include "clang/AST/Attr.h"
#include "clang/AST/CXXInheritance.h"
#include "clang/AST/DeclCXX.h"
+#include "clang/AST/Mangle.h"
#include "clang/AST/MangleNumberingContext.h"
#include "clang/AST/RecordLayout.h"
#include "clang/AST/Type.h"
@@ -64,6 +65,19 @@ class MicrosoftNumberingContext : public MangleNumberingContext {
}
};
+class MSHIPNumberingContext : public MicrosoftNumberingContext {
+ std::unique_ptr<MangleNumberingContext> DeviceCtx;
+
+public:
+ MSHIPNumberingContext(MangleContext *DeviceMangler) {
+ DeviceCtx = createItaniumNumberingContext(DeviceMangler);
+ }
+
+ unsigned getDeviceManglingNumber(const CXXMethodDecl *CallOperator) override {
+ return DeviceCtx->getManglingNumber(CallOperator);
+ }
+};
+
class MicrosoftCXXABI : public CXXABI {
ASTContext &Context;
llvm::SmallDenseMap<CXXRecordDecl *, CXXConstructorDecl *> RecordToCopyCtor;
@@ -73,8 +87,20 @@ class MicrosoftCXXABI : public CXXABI {
llvm::SmallDenseMap<TagDecl *, TypedefNameDecl *>
UnnamedTagDeclToTypedefNameDecl;
+ // MangleContext for device numbering context, which is based on Itanium C++
+ // ABI.
+ std::unique_ptr<MangleContext> DeviceMangler;
+
public:
- MicrosoftCXXABI(ASTContext &Ctx) : Context(Ctx) { }
+ MicrosoftCXXABI(ASTContext &Ctx) : Context(Ctx) {
+ if (Context.getLangOpts().CUDA && Context.getAuxTargetInfo()) {
+ assert(Context.getTargetInfo().getCXXABI().isMicrosoft() &&
+ Context.getAuxTargetInfo()->getCXXABI().isItaniumFamily() &&
+ "Unexpected combination of C++ ABIs.");
+ DeviceMangler.reset(
+ Context.createMangleContext(Context.getAuxTargetInfo()));
+ }
+ }
MemberPointerInfo
getMemberPointerInfo(const MemberPointerType *MPT) const override;
@@ -133,6 +159,10 @@ class MicrosoftCXXABI : public CXXABI {
std::unique_ptr<MangleNumberingContext>
createMangleNumberingContext() const override {
+ if (Context.getLangOpts().CUDA && Context.getAuxTargetInfo()) {
+ assert(DeviceMangler && "Missing device mangler");
+ return std::make_unique<MSHIPNumberingContext>(DeviceMangler.get());
+ }
return std::make_unique<MicrosoftNumberingContext>();
}
};
@@ -266,4 +296,3 @@ CXXABI::MemberPointerInfo MicrosoftCXXABI::getMemberPointerInfo(
CXXABI *clang::CreateMicrosoftCXXABI(ASTContext &Ctx) {
return new MicrosoftCXXABI(Ctx);
}
-
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 33a2d6f4483e..e03631a7243a 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -184,6 +184,14 @@ CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
CharPtrTy = llvm::PointerType::getUnqual(Types.ConvertType(Ctx.CharTy));
VoidPtrTy = cast<llvm::PointerType>(Types.ConvertType(Ctx.VoidPtrTy));
VoidPtrPtrTy = VoidPtrTy->getPointerTo();
+ if (CGM.getContext().getAuxTargetInfo()) {
+ // If the host and device have
diff erent C++ ABIs, mark it as the device
+ // mangle context so that the mangling needs to retrieve the additonal
+ // device lambda mangling number instead of the regular host one.
+ DeviceMC->setDeviceMangleContext(
+ CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() &&
+ CGM.getContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily());
+ }
}
llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const {
diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index af61c82c2002..c1c6a4bf5c68 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -432,15 +432,16 @@ CXXMethodDecl *Sema::startLambdaDefinition(CXXRecordDecl *Class,
void Sema::handleLambdaNumbering(
CXXRecordDecl *Class, CXXMethodDecl *Method,
- Optional<std::tuple<unsigned, bool, Decl *>> Mangling) {
+ Optional<std::tuple<bool, unsigned, unsigned, Decl *>> Mangling) {
if (Mangling) {
- unsigned ManglingNumber;
bool HasKnownInternalLinkage;
+ unsigned ManglingNumber, DeviceManglingNumber;
Decl *ManglingContextDecl;
- std::tie(ManglingNumber, HasKnownInternalLinkage, ManglingContextDecl) =
- Mangling.getValue();
+ std::tie(HasKnownInternalLinkage, ManglingNumber, DeviceManglingNumber,
+ ManglingContextDecl) = Mangling.getValue();
Class->setLambdaMangling(ManglingNumber, ManglingContextDecl,
HasKnownInternalLinkage);
+ Class->setDeviceLambdaManglingNumber(DeviceManglingNumber);
return;
}
@@ -476,6 +477,7 @@ void Sema::handleLambdaNumbering(
unsigned ManglingNumber = MCtx->getManglingNumber(Method);
Class->setLambdaMangling(ManglingNumber, ManglingContextDecl,
HasKnownInternalLinkage);
+ Class->setDeviceLambdaManglingNumber(MCtx->getDeviceManglingNumber(Method));
}
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 0a596e50658b..3c68f9458e58 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -12504,10 +12504,11 @@ TreeTransform<Derived>::TransformLambdaExpr(LambdaExpr *E) {
E->getCaptureDefault());
getDerived().transformedLocalDecl(OldClass, {Class});
- Optional<std::tuple<unsigned, bool, Decl *>> Mangling;
+ Optional<std::tuple<bool, unsigned, unsigned, Decl *>> Mangling;
if (getDerived().ReplacingOriginal())
- Mangling = std::make_tuple(OldClass->getLambdaManglingNumber(),
- OldClass->hasKnownLambdaInternalLinkage(),
+ Mangling = std::make_tuple(OldClass->hasKnownLambdaInternalLinkage(),
+ OldClass->getLambdaManglingNumber(),
+ OldClass->getDeviceLambdaManglingNumber(),
OldClass->getLambdaContextDecl());
// Build the call operator.
diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index 6bfb9bd783b5..18ab4666a7d8 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -1748,6 +1748,7 @@ void ASTDeclReader::ReadCXXDefinitionData(
Lambda.NumExplicitCaptures = Record.readInt();
Lambda.HasKnownInternalLinkage = Record.readInt();
Lambda.ManglingNumber = Record.readInt();
+ D->setDeviceLambdaManglingNumber(Record.readInt());
Lambda.ContextDecl = readDeclID();
Lambda.Captures = (Capture *)Reader.getContext().Allocate(
sizeof(Capture) * Lambda.NumCaptures);
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 6bfa7b0e7d6d..40900af6f9e0 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -5667,6 +5667,7 @@ void ASTRecordWriter::AddCXXDefinitionData(const CXXRecordDecl *D) {
Record->push_back(Lambda.NumExplicitCaptures);
Record->push_back(Lambda.HasKnownInternalLinkage);
Record->push_back(Lambda.ManglingNumber);
+ Record->push_back(D->getDeviceLambdaManglingNumber());
AddDeclRef(D->getLambdaContextDecl());
AddTypeSourceInfo(Lambda.MethodTyInfo);
for (unsigned I = 0, N = Lambda.NumCaptures; I != N; ++I) {
diff --git a/clang/test/CodeGenCUDA/unnamed-types.cu b/clang/test/CodeGenCUDA/unnamed-types.cu
index 59bfa6d7a18f..f598117d969d 100644
--- a/clang/test/CodeGenCUDA/unnamed-types.cu
+++ b/clang/test/CodeGenCUDA/unnamed-types.cu
@@ -1,12 +1,17 @@
// 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
+// 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
// RUN: %clang_cc1 -std=c++11 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
#include "Inputs/cuda.h"
// HOST: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1
+// HOST: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
+// Check that, on MSVC, the same device kernel mangling name is generated.
+// MSVC: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1
+// MSVC: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
__device__ float d0(float x) {
- return [](float x) { return x + 2.f; }(x);
+ return [](float x) { return x + 1.f; }(x);
}
__device__ float d1(float x) {
@@ -14,11 +19,21 @@ __device__ float d1(float x) {
}
// DEVICE: amdgpu_kernel void @_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_(
+// DEVICE: define internal float @_ZZZ2f1PfENKUlS_E_clES_ENKUlfE_clEf(
template <typename F>
__global__ void k0(float *p, F f) {
p[0] = f(p[0]) + d0(p[1]) + d1(p[2]);
}
+// DEVICE: amdgpu_kernel void @_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_(
+// DEVICE: define internal float @_ZZ2f1PfENKUlfE_clEf(
+// DEVICE: define internal float @_ZZ2f1PfENKUlffE_clEff(
+// DEVICE: define internal float @_ZZ2f1PfENKUlfE0_clEf(
+template <typename F0, typename F1, typename F2>
+__global__ void k1(float *p, F0 f0, F1 f1, F2 f2) {
+ p[0] = f0(p[0]) + f1(p[1], p[2]) + f2(p[3]);
+}
+
void f0(float *p) {
[](float *p) {
*p = 1.f;
@@ -29,11 +44,17 @@ void f0(float *p) {
// linkages are still required to keep the original `internal` linkage.
// HOST: define internal void @_ZZ2f1PfENKUlS_E_clES_(
-// DEVICE: define internal float @_ZZZ2f1PfENKUlS_E_clES_ENKUlfE_clEf(
void f1(float *p) {
[](float *p) {
- k0<<<1,1>>>(p, [] __device__ (float x) { return x + 1.f; });
+ k0<<<1,1>>>(p, [] __device__ (float x) { return x + 3.f; });
}(p);
+ k1<<<1,1>>>(p,
+ [] __device__ (float x) { return x + 4.f; },
+ [] __device__ (float x, float y) { return x * y; },
+ [] __device__ (float x) { return x + 5.f; });
}
// HOST: @__hip_register_globals
// HOST: __hipRegisterFunction{{.*}}@_Z17__device_stub__k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_{{.*}}@0
+// HOST: __hipRegisterFunction{{.*}}@_Z17__device_stub__k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_{{.*}}@1
+// MSVC: __hipRegisterFunction{{.*}}@"??$k0 at V<lambda_1>@?0???R1?0??f1@@YAXPEAM at Z@QEBA at 0@Z@@@YAXPEAMV<lambda_1>@?0???R0?0??f1@@YAX0 at Z@QEBA at 0@Z@@Z{{.*}}@0
+// MSVC: __hipRegisterFunction{{.*}}@"??$k1 at V<lambda_2>@?0??f1@@YAXPEAM at Z@V<lambda_3>@?0??2 at YAX0@Z at V<lambda_4>@?0??2 at YAX0@Z@@@YAXPEAMV<lambda_2>@?0??f1@@YAX0 at Z@V<lambda_3>@?0??1 at YAX0@Z at V<lambda_4>@?0??1 at YAX0@Z@@Z{{.*}}@1
More information about the llvm-branch-commits
mailing list