diff --git a/clang/include/clang/Basic/CodeGenOptions.def b/clang/include/clang/Basic/CodeGenOptions.def --- a/clang/include/clang/Basic/CodeGenOptions.def +++ b/clang/include/clang/Basic/CodeGenOptions.def @@ -187,6 +187,7 @@ 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. diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -1007,6 +1007,12 @@ BothFlags<[], " that single precision floating-point divide and sqrt used in " "the program source are correctly rounded (HIP device compilation only)">>, ShouldParseIf; +defm hip_kernel_arg_name : BoolFOption<"hip-kernel-arg-name", + CodeGenOpts<"HIPSaveKernelArgName">, DefaultFalse, + PosFlag, + NegFlag, + BothFlags<[], " that kernel argument names are preserved (HIP only)">>, + ShouldParseIf; def hipspv_pass_plugin_EQ : Joined<["--"], "hipspv-pass-plugin=">, Group, MetaVarName<"">, HelpText<"path to a pass plugin for HIP to SPIR-V passes.">; diff --git a/clang/lib/CodeGen/CGDeclCXX.cpp b/clang/lib/CodeGen/CGDeclCXX.cpp --- a/clang/lib/CodeGen/CGDeclCXX.cpp +++ b/clang/lib/CodeGen/CGDeclCXX.cpp @@ -707,7 +707,7 @@ // 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); } diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -1968,8 +1968,7 @@ /// 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); diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -596,15 +596,17 @@ "decoded_addr"); } -void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, - llvm::Function *Fn) -{ - if (!FD->hasAttr()) +void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD, + llvm::Function *Fn) { + if (!FD->hasAttr() && !FD->hasAttr()) 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()) { QualType HintQTy = A->getTypeHint(); @@ -919,9 +921,10 @@ if (D && D->hasAttr()) 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 diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1460,7 +1460,7 @@ /// \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); diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1697,7 +1697,7 @@ } } -void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, +void CodeGenModule::GenKernelArgMetadata(llvm::Function *Fn, const FunctionDecl *FD, CodeGenFunction *CGF) { assert(((FD && CGF) || (!FD && !CGF)) && @@ -1729,6 +1729,11 @@ 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; @@ -1747,9 +1752,6 @@ } 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); @@ -1822,17 +1824,20 @@ argTypeQuals.push_back(llvm::MDString::get(VMContext, typeQuals)); } - 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 (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 || + getCodeGenOpts().HIPSaveKernelArgName) Fn->setMetadata("kernel_arg_name", llvm::MDNode::get(VMContext, argNames)); } diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -6279,6 +6279,8 @@ if (Args.hasFlag(options::OPT_fgpu_allow_device_init, options::OPT_fno_gpu_allow_device_init, false)) CmdArgs.push_back("-fgpu-allow-device-init"); + Args.addOptInFlag(CmdArgs, options::OPT_fhip_kernel_arg_name, + options::OPT_fno_hip_kernel_arg_name); } if (IsCuda || IsHIP) { diff --git a/clang/test/CodeGenCUDA/kernel-arg-name-metadata.cu b/clang/test/CodeGenCUDA/kernel-arg-name-metadata.cu new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenCUDA/kernel-arg-name-metadata.cu @@ -0,0 +1,16 @@ +// 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"} diff --git a/clang/test/Driver/hip-options.hip b/clang/test/Driver/hip-options.hip --- a/clang/test/Driver/hip-options.hip +++ b/clang/test/Driver/hip-options.hip @@ -116,3 +116,13 @@ // 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" + +// 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" + +// RUN: %clang -### --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib \ +// RUN: --offload-arch=gfx906 %s 2>&1 \ +// RUN: | FileCheck -check-prefix=KANNEG %s +// KANNEG-NOT: "-fhip-kernel-arg-name"