[HIP] add -fhip-kernel-arg-name

[yaxunl]

Summary

Add option -fhip-kernel-arg-name to emit kernel argument
name metadata, which is needed for certain HIP applications.

[tra]

clang/lib/CodeGen/CodeGenModule.cpp

1845–1846

Should we consolidate both options into -fkernel-arg-info and make -cl-kernel-arg-info an alias to it?

1845–1846

Also, this check is odd. For some reason only *arg name* metadata is set conditionally, but for whatever reason OpenCL sets other arg metadata unconditionally.

Now I'm really curious what's so special about "kernel_arg_name" vs the other arg metadata.

[yaxunl]

clang/lib/CodeGen/CodeGenModule.cpp

1845–1846

Should we consolidate both options into -fkernel-arg-info and make -cl-kernel-arg-info an alias to it?

-cl-kernel-arg-info is an OpenCL option defined in OpenCL spec, therefore is made OpenCL only option. It would be confusing to allow it with other languages.

1845–1846

Also, this check is odd. For some reason only *arg name* metadata is set conditionally, but for whatever reason OpenCL sets other arg metadata unconditionally.

Now I'm really curious what's so special about "kernel_arg_name" vs the other arg metadata.

The other metadata are mandatory because they are necessary for OpenCL runtime to set kernel argument.

The kernel argument name is emitted only with -cl-kernel-arg-info since it is only used to support clGetKernelArgInfo which requires -cl-kernel-arg-info to work.

[tra]

clang/lib/CodeGen/CodeGenModule.cpp

1845–1846

It would be confusing to allow it with other languages.

On the other hand having two options that control exactly the same functionality also looks odd to me.

The way I see it is that an entity may have more than one name. OpenCL standard requires that that particular functionality must be enabled via -cl-kernel-arg-info and I'm not proposing to change that. The OpenCL standard has no say whether that flag may have a different name in addition to the standard-required one.

This is similar to how over time we've been transitioning what used to be CUDA-only options into their generic -gpu and -offload variants, only in this case OpenCL functionality becomes useful outside of OpenCL.

[yaxunl]

clang/lib/CodeGen/CodeGenModule.cpp

1845–1846

It would be confusing to allow it with other languages.

On the other hand having two options that control exactly the same functionality also looks odd to me.

The way I see it is that an entity may have more than one name. OpenCL standard requires that that particular functionality must be enabled via -cl-kernel-arg-info and I'm not proposing to change that. The OpenCL standard has no say whether that flag may have a different name in addition to the standard-required one.

This is similar to how over time we've been transitioning what used to be CUDA-only options into their generic -gpu and -offload variants, only in this case OpenCL functionality becomes useful outside of OpenCL.

If we introduce a generic option -fkernel-arg-info and make -cl-kernel-arg-info alias to it, do we want to make it available to all languages or limit it to HIP and OpenCL only?

[b-sumner]

clang/lib/CodeGen/CodeGenModule.cpp

1845–1846

It would be confusing to allow it with other languages.

On the other hand having two options that control exactly the same functionality also looks odd to me.

The way I see it is that an entity may have more than one name. OpenCL standard requires that that particular functionality must be enabled via -cl-kernel-arg-info and I'm not proposing to change that. The OpenCL standard has no say whether that flag may have a different name in addition to the standard-required one.

This is similar to how over time we've been transitioning what used to be CUDA-only options into their generic -gpu and -offload variants, only in this case OpenCL functionality becomes useful outside of OpenCL.

If we introduce a generic option -fkernel-arg-info and make -cl-kernel-arg-info alias to it, do we want to make it available to all languages or limit it to HIP and OpenCL only?

I think that would be fine.

[tra]

clang/lib/CodeGen/CodeGenModule.cpp

1845–1846

If we introduce a generic option -fkernel-arg-info and make -cl-kernel-arg-info alias to it, do we want to make it available to all languages or limit it to HIP and OpenCL only?

ShouldParseIf appears to support expressions, so restricting it to hip + OCL may be as simple as

ShouldParseIf<!strconcat(open_cl.KeyPath, "||", hip.KeyPath)>;

I'm not sure what to do about -cl-kernel-arg-info being in Group<opencl_Group>. AFAICT, option currently may be only in one group. It's not a problem for now as only OpenCL arg needs to be in a group, but that is a point against merging the options. That makes my suggestion a wash -- a bit better here, possibly worse there.

I withdraw the suggestion.

[MaskRay]

clang/lib/Driver/ToolChains/Clang.cpp

6353

Use addOptInFlag

clang/test/Driver/hip-options.hip

120

This comment is redundant. There is only one CHECK line and the string is clear there.

[yaxunl]

clang/lib/Driver/ToolChains/Clang.cpp

6353

will do when committing

clang/test/Driver/hip-options.hip

120

will remove when commiting

clang/include/clang/Basic/CodeGenOptions.def

 CODEGENOPT(EnableSegmentedStacks , 1, 0) ///< Set when -fsplit-stack is enabled.
 CODEGENOPT(NoInlineLineTables, 1, 0) ///< Whether debug info should contain
                                      ///< inline line tables.
 CODEGENOPT(StackClashProtector, 1, 0) ///< Set when -fstack-clash-protection is enabled.
 CODEGENOPT(NoImplicitFloat   , 1, 0) ///< Set when -mno-implicit-float is enabled.
 CODEGENOPT(NullPointerIsValid , 1, 0) ///< Assume Null pointer deference is defined.
 CODEGENOPT(OpenCLCorrectlyRoundedDivSqrt, 1, 0) ///< -cl-fp32-correctly-rounded-divide-sqrt
 CODEGENOPT(HIPCorrectlyRoundedDivSqrt, 1, 1) ///< -fno-hip-fp32-correctly-rounded-divide-sqrt
+CODEGENOPT(HIPSaveKernelArgName, 1, 0) ///< Set when -fhip-kernel-arg-name is enabled.
 CODEGENOPT(UniqueInternalLinkageNames, 1, 0) ///< Internal Linkage symbols get unique names.
 CODEGENOPT(SplitMachineFunctions, 1, 0) ///< Split machine functions using profile information.
 
 /// When false, this attempts to generate code as if the result of an
 /// overflowing conversion matches the overflowing behavior of a target's native
 /// float-to-int conversion instructions.
 CODEGENOPT(StrictFloatCastOverflow, 1, 1)
 

clang/include/clang/Driver/Options.td

defm hip_new_launch_api : BoolFOption<"hip-new-launch-api",
   BothFlags<[], " new kernel launching API for HIP">>;
 defm hip_fp32_correctly_rounded_divide_sqrt : BoolFOption<"hip-fp32-correctly-rounded-divide-sqrt",
   CodeGenOpts<"HIPCorrectlyRoundedDivSqrt">, DefaultTrue,
   PosFlag<SetTrue, [], "Specify">,
   NegFlag<SetFalse, [CC1Option], "Don't specify">,
   BothFlags<[], " that single precision floating-point divide and sqrt used in "
   "the program source are correctly rounded (HIP device compilation only)">>,
   ShouldParseIf<hip.KeyPath>;
+defm hip_kernel_arg_name : BoolFOption<"hip-kernel-arg-name",
+  CodeGenOpts<"HIPSaveKernelArgName">, DefaultFalse,
+  PosFlag<SetTrue, [CC1Option], "Specify">,
+  NegFlag<SetFalse, [], "Don't specify">,
+  BothFlags<[], " that kernel argument names are preserved (HIP only)">>,
+  ShouldParseIf<hip.KeyPath>;
 def hipspv_pass_plugin_EQ : Joined<["--"], "hipspv-pass-plugin=">,
   Group<Link_Group>, MetaVarName<"<dsopath>">,
   HelpText<"path to a pass plugin for HIP to SPIR-V passes.">;
 defm gpu_allow_device_init : BoolFOption<"gpu-allow-device-init",
   LangOpts<"GPUAllowDeviceInit">, DefaultFalse,
   PosFlag<SetTrue, [CC1Option], "Allow">, NegFlag<SetFalse, [], "Don't allow">,
   BothFlags<[], " device side init function in HIP (experimental)">>,
   ShouldParseIf<hip.KeyPath>;

clang/lib/CodeGen/CGDeclCXX.cpp

CodeGenModule::EmitCXXGlobalInitFunc() {
   // be able to launch them from the host.
   // FIXME: Some more work might be needed to handle destructors correctly.
   // Current initialization function makes use of function pointers callbacks.
   // We can't support function pointers especially between host and device.
   // However it seems global destruction has little meaning without any
   // dynamic resource allocation on the device and program scope variables are
   // destroyed by the runtime when program is released.
   if (getLangOpts().OpenCL) {
-    GenOpenCLArgMetadata(Fn);
+    GenKernelArgMetadata(Fn);
     Fn->setCallingConv(llvm::CallingConv::SPIR_KERNEL);
   }
 
   assert(!getLangOpts().CUDA || !getLangOpts().CUDAIsDevice ||
          getLangOpts().GPUAllowDeviceInit);
   if (getLangOpts().HIP && getLangOpts().CUDAIsDevice) {
     Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
     Fn->addFnAttr("device-init");

clang/lib/CodeGen/CodeGenFunction.h

private:
   unsigned LargestVectorWidth = 0;
 
   /// True if we need emit the life-time markers. This is initially set in
   /// the constructor, but could be overwritten to true if this is a coroutine.
   bool ShouldEmitLifetimeMarkers;
 
   /// Add OpenCL kernel arg metadata and the kernel attribute metadata to
   /// the function metadata.
-  void EmitOpenCLKernelMetadata(const FunctionDecl *FD,
-                                llvm::Function *Fn);
+  void EmitKernelMetadata(const FunctionDecl *FD, llvm::Function *Fn);
 
 public:
   CodeGenFunction(CodeGenModule &cgm, bool suppressNewContext=false);
   ~CodeGenFunction();
 
   CodeGenTypes &getTypes() const { return CGM.getTypes(); }
   ASTContext &getContext() const { return CGM.getContext(); }
   CGDebugInfo *getDebugInfo() {

clang/lib/CodeGen/CodeGenFunction.cpp

CodeGenFunction::DecodeAddrUsedInPrologue(llvm::Value *F,
   auto *GOTAsInt = Builder.CreateAdd(PCRelAsInt, FuncAsInt, "global_addr.int");
   auto *GOTAddr = Builder.CreateIntToPtr(GOTAsInt, Int8PtrPtrTy, "global_addr");
 
   // Load the original pointer through the global.
   return Builder.CreateLoad(Address(GOTAddr, Int8PtrTy, getPointerAlign()),
                             "decoded_addr");
 }
 
-void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
-                                               llvm::Function *Fn)
-{
-  if (!FD->hasAttr<OpenCLKernelAttr>())
+void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
+                                         llvm::Function *Fn) {
+  if (!FD->hasAttr<OpenCLKernelAttr>() && !FD->hasAttr<CUDAGlobalAttr>())
     return;
 
   llvm::LLVMContext &Context = getLLVMContext();
 
-  CGM.GenOpenCLArgMetadata(Fn, FD, this);
+  CGM.GenKernelArgMetadata(Fn, FD, this);
+
+  if (!getLangOpts().OpenCL)
+    return;
 
   if (const VecTypeHintAttr *A = FD->getAttr<VecTypeHintAttr>()) {
     QualType HintQTy = A->getTypeHint();
     const ExtVectorType *HintEltQTy = HintQTy->getAs<ExtVectorType>();
     bool IsSignedInteger =
         HintQTy->isSignedIntegerType() ||
         (HintEltQTy && HintEltQTy->getElementType()->isSignedIntegerType());
     llvm::Metadata *AttrMDArgs[] = {
if (!CGM.getCodeGenOpts().SampleProfileFile.empty())
     Fn->addFnAttr("use-sample-profile");
 
   if (D && D->hasAttr<CFICanonicalJumpTableAttr>())
     Fn->addFnAttr("cfi-canonical-jump-table");
 
   if (D && D->hasAttr<NoProfileFunctionAttr>())
     Fn->addFnAttr(llvm::Attribute::NoProfile);
 
-  if (FD && getLangOpts().OpenCL) {
+  if (FD && (getLangOpts().OpenCL ||
+             (getLangOpts().HIP && getLangOpts().CUDAIsDevice))) {
     // Add metadata for a kernel function.
-    EmitOpenCLKernelMetadata(FD, Fn);
+    EmitKernelMetadata(FD, Fn);
   }
 
   // If we are checking function types, emit a function type signature as
   // prologue data.
   if (FD && getLangOpts().CPlusPlus && SanOpts.has(SanitizerKind::Function)) {
     if (llvm::Constant *PrologueSig = getPrologueSignature(CGM, FD)) {
       // Remove any (C++17) exception specifications, to allow calling e.g. a
       // noexcept function through a non-noexcept pointer.

clang/lib/CodeGen/CodeGenModule.h

public:
   /// includes the argument name, its type, the address and access qualifiers
   /// used. This helper can be used to generate metadata for source code kernel
   /// function as well as generated implicitly kernels. If a kernel is generated
   /// implicitly null value has to be passed to the last two parameters,
   /// otherwise all parameters must have valid non-null values.
   /// \param FN is a pointer to IR function being generated.
   /// \param FD is a pointer to function declaration if any.
   /// \param CGF is a pointer to CodeGenFunction that generates this function.
-  void GenOpenCLArgMetadata(llvm::Function *FN,
+  void GenKernelArgMetadata(llvm::Function *FN,
                             const FunctionDecl *FD = nullptr,
                             CodeGenFunction *CGF = nullptr);
 
   /// Get target specific null pointer.
   /// \param T is the LLVM type of the null pointer.
   /// \param QT is the clang QualType of the null pointer.
   llvm::Constant *getNullPointer(llvm::PointerType *T, QualType QT);
 

clang/lib/CodeGen/CodeGenModule.cpp

case LangAS::opencl_global_device:
     return 5;
   case LangAS::opencl_global_host:
     return 6;
   default:
     return 0; // Assume private.
   }
 }
 
-void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
+void CodeGenModule::GenKernelArgMetadata(llvm::Function *Fn,
                                          const FunctionDecl *FD,
                                          CodeGenFunction *CGF) {
   assert(((FD && CGF) || (!FD && !CGF)) &&
          "Incorrect use - FD and CGF should either be both null or not!");
   // Create MDNodes that represent the kernel arg metadata.
   // Each MDNode is a list in the form of "key", N number of values which is
   // the same number of values as their are kernel arguments.
 
void CodeGenModule::GenKernelArgMetadata(llvm::Function *Fn,
   SmallVector<llvm::Metadata *, 8> argTypeQuals;
 
   // MDNode for the kernel argument names.
   SmallVector<llvm::Metadata *, 8> argNames;
 
   if (FD && CGF)
     for (unsigned i = 0, e = FD->getNumParams(); i != e; ++i) {
       const ParmVarDecl *parm = FD->getParamDecl(i);
+      // Get argument name.
+      argNames.push_back(llvm::MDString::get(VMContext, parm->getName()));
+
+      if (!getLangOpts().OpenCL)
+        continue;
       QualType ty = parm->getType();
       std::string typeQuals;
 
       // Get image and pipe access qualifier:
       if (ty->isImageType() || ty->isPipeType()) {
         const Decl *PDecl = parm;
         if (auto *TD = dyn_cast<TypedefType>(ty))
           PDecl = TD->getDecl();
         const OpenCLAccessAttr *A = PDecl->getAttr<OpenCLAccessAttr>();
         if (A && A->isWriteOnly())
           accessQuals.push_back(llvm::MDString::get(VMContext, "write_only"));
         else if (A && A->isReadWrite())
           accessQuals.push_back(llvm::MDString::get(VMContext, "read_write"));
         else
           accessQuals.push_back(llvm::MDString::get(VMContext, "read_only"));
       } else
         accessQuals.push_back(llvm::MDString::get(VMContext, "none"));
 
-      // Get argument name.
-      argNames.push_back(llvm::MDString::get(VMContext, parm->getName()));
-
       auto getTypeSpelling = [&](QualType Ty) {
         auto typeName = Ty.getUnqualifiedType().getAsString(Policy);
 
         if (Ty.isCanonical()) {
           StringRef typeNameRef = typeName;
           // Turn "unsigned type" to "utype"
           if (typeNameRef.consume_front("unsigned "))
             return std::string("u") + typeNameRef.str();
for (unsigned i = 0, e = FD->getNumParams(); i != e; ++i) {
             llvm::MDString::get(VMContext, baseTypeName));
 
         if (isPipe)
           typeQuals = "pipe";
       }
       argTypeQuals.push_back(llvm::MDString::get(VMContext, typeQuals));
     }
 
+  if (getLangOpts().OpenCL) {
     Fn->setMetadata("kernel_arg_addr_space",
                     llvm::MDNode::get(VMContext, addressQuals));
     Fn->setMetadata("kernel_arg_access_qual",
                     llvm::MDNode::get(VMContext, accessQuals));
     Fn->setMetadata("kernel_arg_type",
                     llvm::MDNode::get(VMContext, argTypeNames));
     Fn->setMetadata("kernel_arg_base_type",
                     llvm::MDNode::get(VMContext, argBaseTypeNames));
     Fn->setMetadata("kernel_arg_type_qual",
                     llvm::MDNode::get(VMContext, argTypeQuals));
-  if (getCodeGenOpts().EmitOpenCLArgMetadata)
+  }
+  if (getCodeGenOpts().EmitOpenCLArgMetadata ||
+      getCodeGenOpts().HIPSaveKernelArgName)
     Fn->setMetadata("kernel_arg_name",
                     llvm::MDNode::get(VMContext, argNames));
 }
 
 /// Determines whether the language options require us to model
 /// unwind exceptions.  We treat -fexceptions as mandating this
 /// except under the fragile ObjC ABI with only ObjC exceptions
 /// enabled.  This means, for example, that C with -fexceptions

clang/lib/Driver/ToolChains/Clang.cpp

 
   if (IsHIP) {
     if (Args.hasFlag(options::OPT_fhip_new_launch_api,
                      options::OPT_fno_hip_new_launch_api, true))
       CmdArgs.push_back("-fhip-new-launch-api");
     if (Args.hasFlag(options::OPT_fgpu_allow_device_init,
                      options::OPT_fno_gpu_allow_device_init, false))
       CmdArgs.push_back("-fgpu-allow-device-init");
+    if (Args.hasFlag(options::OPT_fhip_kernel_arg_name,
+                     options::OPT_fno_hip_kernel_arg_name, false))
+      CmdArgs.push_back("-fhip-kernel-arg-name");
   }
 
   if (IsCuda || IsHIP) {
     if (!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false) &&
         Args.hasArg(options::OPT_offload_new_driver))
       D.Diag(diag::err_drv_no_rdc_new_driver);
     if (Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false))
       CmdArgs.push_back("-fgpu-rdc");

clang/test/CodeGenCUDA/kernel-arg-name-metadata.cu

+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fhip-kernel-arg-name \
+// RUN:     -fcuda-is-device -emit-llvm -o - -x hip %s \
+// RUN:     | FileCheck %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN:     -fcuda-is-device -emit-llvm -o - -x hip %s \
+// RUN:     | FileCheck -check-prefix=NEG %s
+
+#include "Inputs/cuda.h"
+
+// CHECK: define{{.*}} amdgpu_kernel void @_Z6kerneliPf({{.*}} !kernel_arg_name [[MD:![0-9]+]]
+// NEG-NOT: define{{.*}} amdgpu_kernel void @_Z6kerneliPf({{.*}} !kernel_arg_name
+__global__ void kernel(int arg1, float *arg2) {
+}
+
+// CHECK: [[MD]] = !{!"arg1", !"arg2"}

clang/test/Driver/hip-options.hip

 
 // RUN: %clang -### --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib \
 // RUN:   --cuda-gpu-arch=gfx906 -fgpu-rdc -Xoffload-linker --build-id=md5 %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=OFL-LINK %s
 // RUN: %clang -### --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib \
 // RUN:   --cuda-gpu-arch=gfx906 -Xoffload-linker --build-id=md5 %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=OFL-LINK %s
 // OFL-LINK: lld{{.*}}"--build-id=md5"
+
+// Check -fhip-kernel-arg-name is passed to clang -cc1.
+
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib \
+// RUN:   --offload-arch=gfx906 -fhip-kernel-arg-name %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=KAN %s
+// KAN: "-cc1"{{.*}} "-triple" "amdgcn-amd-amdhsa" {{.*}} "-fhip-kernel-arg-name"