[HIP] add __builtin_get_device_side_mangled_name

[yaxunl]

Summary

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.

[tra]

LGTM with a couple of nits.

clang/include/clang/Basic/Builtins.h

39

The issue is common for both CUDA and HIP.

clang/include/clang/Basic/DiagnosticSemaKinds.td

8310

Nit. I'd rephrase it in terms of what makes the argument invalid. E.g. symbol must be a device-side function or global variable.

clang/lib/Basic/Builtins.cpp

78

Please enable it for CUDA, too.

78–84

Wow! The density of negations in this function is impressive.

clang/lib/Sema/SemaChecking.cpp

1980–1983

return D->hasAttr<CUDAGlobalAttr>() || D->hasAttr<CUDADeviceAttr>() ||
        D->hasAttr<CUDAConstantAttr>() || D->hasAttr<HIPManagedAttr>();

[yaxunl]

Revised by Artem's comments.

[yaxunl]

clang/include/clang/Basic/Builtins.h

39

will rename it to CUDA_LANG

clang/include/clang/Basic/DiagnosticSemaKinds.td

8310

will do

clang/lib/Basic/Builtins.cpp

78

will do

clang/lib/Sema/SemaChecking.cpp

1980–1983

will do

clang/include/clang/Basic/Builtins.h

enum LanguageID {
   GNU_LANG = 0x1,     // builtin requires GNU mode.
   C_LANG = 0x2,       // builtin for c only.
   CXX_LANG = 0x4,     // builtin for cplusplus only.
   OBJC_LANG = 0x8,    // builtin for objective-c and objective-c++
   MS_LANG = 0x10,     // builtin requires MS mode.
   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.
+  HIP_LANG = 0x100,   // builtin requires HIP.
   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.
   ALL_OCLC_LANGUAGES = OCLC1X_LANG | OCLC20_LANG // builtin for OCLC languages.
 };
 
 namespace Builtin {
 enum ID {

clang/include/clang/Basic/Builtins.def

 
 // Builtins for os_log/os_trace
 BUILTIN(__builtin_os_log_format_buffer_size, "zcC*.", "p:0:nut")
 BUILTIN(__builtin_os_log_format, "v*v*cC*.", "p:0:nt")
 
 // OpenMP 4.0
 LANGBUILTIN(omp_is_initial_device, "i", "nc", OMP_LANG)
 
+// HIP
+LANGBUILTIN(__builtin_get_device_side_mangled_name, "cC*.", "ncT", HIP_LANG)
+
 // Builtins for XRay
 BUILTIN(__xray_customevent, "vcC*z", "")
 BUILTIN(__xray_typedevent, "vzcC*z", "")
 
 // Win64-compatible va_list functions
 BUILTIN(__builtin_ms_va_start, "vc*&.", "nt")
 BUILTIN(__builtin_ms_va_end, "vc*&", "n")
 BUILTIN(__builtin_ms_va_copy, "vc*&c*&", "n")
 
 #undef BUILTIN
 #undef LIBBUILTIN
 #undef LANGBUILTIN

clang/include/clang/Basic/DiagnosticSemaKinds.td

 
 def err_cuda_device_builtin_surftex_ref_decl : Error<
     "illegal device builtin %select{surface|texture}0 reference "
     "type %1 declared here">;
 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: expect 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 "
   "string was %3">, InGroup<NonPODVarargs>, DefaultError;
 // The arguments to this diagnostic should match the warning above.
 def err_cannot_pass_objc_interface_to_vararg_format : Error<
   "cannot pass object with interface type %1 by value to variadic "
   "%select{function|block|method|constructor}2; expected type from format "

clang/lib/Basic/Builtins.cpp

bool Builtin::Context::builtinIsSupported(const Builtin::Info &BuiltinInfo,
   bool OclC1Unsupported = (LangOpts.OpenCLVersion / 100) != 1 &&
                           (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES ) ==  OCLC1X_LANG;
   bool OclC2Unsupported =
       (LangOpts.OpenCLVersion != 200 && !LangOpts.OpenCLCPlusPlus) &&
       (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES) == OCLC20_LANG;
   bool OclCUnsupported = !LangOpts.OpenCL &&
                          (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES);
   bool OpenMPUnsupported = !LangOpts.OpenMP && BuiltinInfo.Langs == OMP_LANG;
+  bool HIPUnsupported = !LangOpts.HIP && BuiltinInfo.Langs == HIP_LANG;
   bool CPlusPlusUnsupported =
       !LangOpts.CPlusPlus && BuiltinInfo.Langs == CXX_LANG;
   return !BuiltinsUnsupported && !MathBuiltinsUnsupported && !OclCUnsupported &&
          !OclC1Unsupported && !OclC2Unsupported && !OpenMPUnsupported &&
          !GnuModeUnsupported && !MSModeUnsupported && !ObjCUnsupported &&
-         !CPlusPlusUnsupported;
+         !CPlusPlusUnsupported && !HIPUnsupported;
 }
 
 /// initializeBuiltins - Mark the identifiers for all the builtins with their
 /// appropriate builtin ID # and mark any non-portable builtin identifiers as
 /// such.
 void Builtin::Context::initializeBuiltins(IdentifierTable &Table,
                                           const LangOptions& LangOpts) {
   // Step #1: mark all target-independent builtins with their ID's.

clang/lib/CodeGen/CGBuiltin.cpp

 //===---- CGBuiltin.cpp - Emit LLVM Code for builtins ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This contains code to emit Builtin calls as LLVM code.
 //
 //===----------------------------------------------------------------------===//
 
+#include "CGCUDARuntime.h"
 #include "CGCXXABI.h"
 #include "CGObjCRuntime.h"
 #include "CGOpenCLRuntime.h"
 #include "CGRecordLayout.h"
 #include "CodeGenFunction.h"
 #include "CodeGenModule.h"
 #include "ConstantEmitter.h"
 #include "PatternInit.h"
case Builtin::BI__builtin_ms_va_copy: {
     DestAddr = Address(Builder.CreateBitCast(DestAddr.getPointer(), BPP, "cp"),
                        DestAddr.getAlignment());
     SrcAddr = Address(Builder.CreateBitCast(SrcAddr.getPointer(), BPP, "ap"),
                       SrcAddr.getAlignment());
 
     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
   // the call using the normal call path, but using the unmangled
   // version of the function name.
   if (getContext().BuiltinInfo.isLibFunction(BuiltinID))
     return emitLibraryCall(*this, FD, E,
                            CGM.getBuiltinLibFunction(FD, BuiltinID));

clang/lib/CodeGen/CGCUDANV.cpp

 //===----- CGCUDANV.cpp - Interface to NVIDIA CUDA Runtime ----------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This provides a class for CUDA code generation targeting the NVIDIA CUDA
 // runtime library.
 //
 //===----------------------------------------------------------------------===//
 
 #include "CGCUDARuntime.h"
+#include "CGCXXABI.h"
 #include "CodeGenFunction.h"
 #include "CodeGenModule.h"
 #include "clang/AST/Decl.h"
 #include "clang/Basic/Cuda.h"
 #include "clang/CodeGen/CodeGenABITypes.h"
 #include "clang/CodeGen/ConstantInitBuilder.h"
 #include "llvm/IR/BasicBlock.h"
 #include "llvm/IR/Constants.h"

 std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) {
   GlobalDecl GD;
   // D could be either a kernel or a variable.
   if (auto *FD = dyn_cast<FunctionDecl>(ND))
     GD = GlobalDecl(FD, KernelReferenceKind::Kernel);
   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());
 
   // Make unique name for device side static file-scope variable for HIP.
   if (CGM.getContext().shouldExternalizeStaticVar(ND) &&
       CGM.getLangOpts().GPURelocatableDeviceCode &&
       !CGM.getLangOpts().CUID.empty()) {

clang/lib/Sema/SemaChecking.cpp

#include "clang/Basic/Builtins.def"
   case Builtin::BI__builtin_matrix_transpose:
     return SemaBuiltinMatrixTranspose(TheCall, TheCallResult);
 
   case Builtin::BI__builtin_matrix_column_major_load:
     return SemaBuiltinMatrixColumnMajorLoad(TheCall, TheCallResult);
 
   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;
+      if (!D->hasAttr<CUDAGlobalAttr>() && !D->hasAttr<CUDADeviceAttr>() &&
+          !D->hasAttr<CUDAConstantAttr>() && !D->hasAttr<HIPManagedAttr>())
+        return false;
+      return true;
+    };
+    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
   // of the arch we are compiling for.
   if (Context.BuiltinInfo.isTSBuiltin(BuiltinID)) {
     if (Context.BuiltinInfo.isAuxBuiltinID(BuiltinID)) {
       assert(Context.getAuxTargetInfo() &&
              "Aux Target Builtin, but not an aux target?");

clang/test/CodeGenCUDA/builtin-mangled-name.cu

+// 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);
+}

clang/test/SemaCUDA/builtin-mangled-name.cu

+// 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: expect a device-side function or global variable}}
+  p = __builtin_get_device_side_mangled_name(kern1, kern1);
+  // expected-error@-1 {{invalid argument: expect a device-side function or global variable}}
+  p = __builtin_get_device_side_mangled_name(1);
+  // expected-error@-1 {{invalid argument: expect a device-side function or global variable}}
+  p = __builtin_get_device_side_mangled_name(x);
+  // expected-error@-1 {{invalid argument: expect a device-side function or global variable}}
+  p = __builtin_get_device_side_mangled_name(fun1);
+  // expected-error@-1 {{invalid argument: expect a device-side function or global variable}}
+  p = __builtin_get_device_side_mangled_name(y);
+  // expected-error@-1 {{invalid argument: expect a device-side function or global variable}}
+}