Author: Yaxun (Sam) Liu Date: 2021-03-25T15:25:29-04:00 New Revision: cc9477166a53faced47cbd4146ac4adea431ccfd
URL: https://github.com/llvm/llvm-project/commit/cc9477166a53faced47cbd4146ac4adea431ccfd DIFF: https://github.com/llvm/llvm-project/commit/cc9477166a53faced47cbd4146ac4adea431ccfd.diff LOG: [CUDA][HIP] add __builtin_get_device_side_mangled_name Add builtin function __builtin_get_device_side_mangled_name to get device side manged name for functions and global variables, which can be used to get symbol address of kernels or variables by mangled name in dynamically loaded bundled code objects at run time. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D99301 Added: clang/test/CodeGenCUDA/builtin-mangled-name.cu clang/test/SemaCUDA/builtin-mangled-name.cu Modified: clang/include/clang/Basic/Builtins.def clang/include/clang/Basic/Builtins.h clang/include/clang/Basic/DiagnosticSemaKinds.td clang/lib/Basic/Builtins.cpp clang/lib/CodeGen/CGBuiltin.cpp clang/lib/CodeGen/CGCUDANV.cpp clang/lib/Sema/SemaChecking.cpp Removed: ################################################################################ diff --git a/clang/include/clang/Basic/Builtins.def b/clang/include/clang/Basic/Builtins.def index ab1b5866c8a7..153e22f00f52 100644 --- a/clang/include/clang/Basic/Builtins.def +++ b/clang/include/clang/Basic/Builtins.def @@ -1639,6 +1639,9 @@ BUILTIN(__builtin_os_log_format, "v*v*cC*.", "p:0:nt") // OpenMP 4.0 LANGBUILTIN(omp_is_initial_device, "i", "nc", OMP_LANG) +// CUDA/HIP +LANGBUILTIN(__builtin_get_device_side_mangled_name, "cC*.", "ncT", CUDA_LANG) + // Builtins for XRay BUILTIN(__xray_customevent, "vcC*z", "") BUILTIN(__xray_typedevent, "vzcC*z", "") diff --git a/clang/include/clang/Basic/Builtins.h b/clang/include/clang/Basic/Builtins.h index 15bfcf797917..efd6cb897293 100644 --- a/clang/include/clang/Basic/Builtins.h +++ b/clang/include/clang/Basic/Builtins.h @@ -36,6 +36,7 @@ enum LanguageID { OCLC20_LANG = 0x20, // builtin for OpenCL C 2.0 only. OCLC1X_LANG = 0x40, // builtin for OpenCL C 1.x only. OMP_LANG = 0x80, // builtin requires OpenMP. + CUDA_LANG = 0x100, // builtin requires CUDA. ALL_LANGUAGES = C_LANG | CXX_LANG | OBJC_LANG, // builtin for all languages. ALL_GNU_LANGUAGES = ALL_LANGUAGES | GNU_LANG, // builtin requires GNU mode. ALL_MS_LANGUAGES = ALL_LANGUAGES | MS_LANG, // builtin requires MS mode. diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index df2f79a4f344..ad592d552030 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -8303,6 +8303,9 @@ def note_cuda_device_builtin_surftex_should_be_template_class : Note< "%0 needs to be instantiated from a class template with proper " "template arguments">; +def err_hip_invalid_args_builtin_mangled_name : Error< + "invalid argument: symbol must be a device-side function or global variable">; + def warn_non_pod_vararg_with_format_string : Warning< "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic " "%select{function|block|method|constructor}2; expected type from format " diff --git a/clang/lib/Basic/Builtins.cpp b/clang/lib/Basic/Builtins.cpp index 0cd89df41b67..49afaa9ba6a3 100644 --- a/clang/lib/Basic/Builtins.cpp +++ b/clang/lib/Basic/Builtins.cpp @@ -75,12 +75,13 @@ bool Builtin::Context::builtinIsSupported(const Builtin::Info &BuiltinInfo, bool OclCUnsupported = !LangOpts.OpenCL && (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES); bool OpenMPUnsupported = !LangOpts.OpenMP && BuiltinInfo.Langs == OMP_LANG; + bool CUDAUnsupported = !LangOpts.CUDA && BuiltinInfo.Langs == CUDA_LANG; bool CPlusPlusUnsupported = !LangOpts.CPlusPlus && BuiltinInfo.Langs == CXX_LANG; return !BuiltinsUnsupported && !MathBuiltinsUnsupported && !OclCUnsupported && !OclC1Unsupported && !OclC2Unsupported && !OpenMPUnsupported && !GnuModeUnsupported && !MSModeUnsupported && !ObjCUnsupported && - !CPlusPlusUnsupported; + !CPlusPlusUnsupported && !CUDAUnsupported; } /// initializeBuiltins - Mark the identifiers for all the builtins with their diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index f86b7e52c9a9..7d24b6a9342e 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -10,6 +10,7 @@ // //===----------------------------------------------------------------------===// +#include "CGCUDARuntime.h" #include "CGCXXABI.h" #include "CGObjCRuntime.h" #include "CGOpenCLRuntime.h" @@ -5058,6 +5059,17 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID, Value *ArgPtr = Builder.CreateLoad(SrcAddr, "ap.val"); return RValue::get(Builder.CreateStore(ArgPtr, DestAddr)); } + + case Builtin::BI__builtin_get_device_side_mangled_name: { + auto Name = CGM.getCUDARuntime().getDeviceSideName( + cast<DeclRefExpr>(E->getArg(0)->IgnoreImpCasts())->getDecl()); + auto Str = CGM.GetAddrOfConstantCString(Name, ""); + llvm::Constant *Zeros[] = {llvm::ConstantInt::get(SizeTy, 0), + llvm::ConstantInt::get(SizeTy, 0)}; + auto *Ptr = llvm::ConstantExpr::getGetElementPtr(Str.getElementType(), + Str.getPointer(), Zeros); + return RValue::get(Ptr); + } } // If this is an alias for a lib function (e.g. __builtin_sin), emit diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index 3a311ab395e4..d53a623b258c 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -12,6 +12,7 @@ //===----------------------------------------------------------------------===// #include "CGCUDARuntime.h" +#include "CGCXXABI.h" #include "CodeGenFunction.h" #include "CodeGenModule.h" #include "clang/AST/Decl.h" @@ -260,10 +261,15 @@ std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) { else GD = GlobalDecl(ND); std::string DeviceSideName; - if (DeviceMC->shouldMangleDeclName(ND)) { + MangleContext *MC; + if (CGM.getLangOpts().CUDAIsDevice) + MC = &CGM.getCXXABI().getMangleContext(); + else + MC = DeviceMC.get(); + if (MC->shouldMangleDeclName(ND)) { SmallString<256> Buffer; llvm::raw_svector_ostream Out(Buffer); - DeviceMC->mangleName(GD, Out); + MC->mangleName(GD, Out); DeviceSideName = std::string(Out.str()); } else DeviceSideName = std::string(ND->getIdentifier()->getName()); diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp index 0570f61458a2..305fcd574a37 100644 --- a/clang/lib/Sema/SemaChecking.cpp +++ b/clang/lib/Sema/SemaChecking.cpp @@ -1966,6 +1966,26 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID, case Builtin::BI__builtin_matrix_column_major_store: return SemaBuiltinMatrixColumnMajorStore(TheCall, TheCallResult); + + case Builtin::BI__builtin_get_device_side_mangled_name: { + auto Check = [](CallExpr *TheCall) { + if (TheCall->getNumArgs() != 1) + return false; + auto *DRE = dyn_cast<DeclRefExpr>(TheCall->getArg(0)->IgnoreImpCasts()); + if (!DRE) + return false; + auto *D = DRE->getDecl(); + if (!isa<FunctionDecl>(D) && !isa<VarDecl>(D)) + return false; + return D->hasAttr<CUDAGlobalAttr>() || D->hasAttr<CUDADeviceAttr>() || + D->hasAttr<CUDAConstantAttr>() || D->hasAttr<HIPManagedAttr>(); + }; + if (!Check(TheCall)) { + Diag(TheCall->getBeginLoc(), + diag::err_hip_invalid_args_builtin_mangled_name); + return ExprError(); + } + } } // Since the target specific builtins for each arch overlap, only check those diff --git a/clang/test/CodeGenCUDA/builtin-mangled-name.cu b/clang/test/CodeGenCUDA/builtin-mangled-name.cu new file mode 100644 index 000000000000..e9dca5680155 --- /dev/null +++ b/clang/test/CodeGenCUDA/builtin-mangled-name.cu @@ -0,0 +1,28 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-gnu-linux -aux-triple amdgcn-amd-amdhsa \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=CHECK,LNX %s +// RUN: %clang_cc1 -triple x86_64-unknown-windows-msvc -aux-triple amdgcn-amd-amdhsa \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=CHECK,MSVC %s + +#include "Inputs/cuda.h" + +namespace X { + __global__ void kern1(int *x); + __device__ int var1; +} + +// CHECK: @[[STR1:.*]] = {{.*}} c"_ZN1X5kern1EPi\00" +// CHECK: @[[STR2:.*]] = {{.*}} c"_ZN1X4var1E\00" + +// LNX-LABEL: define {{.*}}@_Z4fun1v() +// MSVC-LABEL: define {{.*}} @"?fun1@@YAPEBDXZ"() +// CHECK: ret i8* getelementptr inbounds ({{.*}} @[[STR1]], i64 0, i64 0) +const char *fun1() { + return __builtin_get_device_side_mangled_name(X::kern1); +} + +// LNX-LABEL: define {{.*}}@_Z4fun2v() +// MSVC-LABEL: define {{.*}}@"?fun2@@YAPEBDXZ"() +// CHECK: ret i8* getelementptr inbounds ({{.*}} @[[STR2]], i64 0, i64 0) +__host__ __device__ const char *fun2() { + return __builtin_get_device_side_mangled_name(X::var1); +} diff --git a/clang/test/SemaCUDA/builtin-mangled-name.cu b/clang/test/SemaCUDA/builtin-mangled-name.cu new file mode 100644 index 000000000000..6ca85083d717 --- /dev/null +++ b/clang/test/SemaCUDA/builtin-mangled-name.cu @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-gnu-linux -aux-triple amdgcn-amd-amdhsa \ +// RUN: -verify -fsyntax-only -x hip %s + +#include "Inputs/cuda.h" + +__global__ void kern1(); +int y; + +void fun1() { + int x; + const char *p; + p = __builtin_get_device_side_mangled_name(); + // expected-error@-1 {{invalid argument: symbol must be a device-side function or global variable}} + p = __builtin_get_device_side_mangled_name(kern1, kern1); + // expected-error@-1 {{invalid argument: symbol must be a device-side function or global variable}} + p = __builtin_get_device_side_mangled_name(1); + // expected-error@-1 {{invalid argument: symbol must be a device-side function or global variable}} + p = __builtin_get_device_side_mangled_name(x); + // expected-error@-1 {{invalid argument: symbol must be a device-side function or global variable}} + p = __builtin_get_device_side_mangled_name(fun1); + // expected-error@-1 {{invalid argument: symbol must be a device-side function or global variable}} + p = __builtin_get_device_side_mangled_name(y); + // expected-error@-1 {{invalid argument: symbol must be a device-side function or global variable}} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
