From 8ad4c6e4b1299d599c0b6defe6a9e90a417c7ba8 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Thu, 16 Jun 2022 21:26:33 -0400 Subject: [PATCH] [HIP] add -fhip-kernel-arg-name Add option -fhip-kernel-arg-name to emit kernel argument name metadata, which is needed for certain HIP applications. Reviewed by: Artem Belevich, Fangrui Song, Brian Sumner Differential Revision: https://reviews.llvm.org/D128022 --- clang/include/clang/Basic/CodeGenOptions.def | 1 + clang/include/clang/Driver/Options.td | 6 ++++ clang/lib/CodeGen/CGDeclCXX.cpp | 2 +- clang/lib/CodeGen/CodeGenFunction.cpp | 17 ++++++----- clang/lib/CodeGen/CodeGenFunction.h | 3 +- clang/lib/CodeGen/CodeGenModule.cpp | 35 ++++++++++++---------- clang/lib/CodeGen/CodeGenModule.h | 2 +- clang/lib/Driver/ToolChains/Clang.cpp | 2 ++ clang/test/CodeGenCUDA/kernel-arg-name-metadata.cu | 16 ++++++++++ clang/test/Driver/hip-options.hip | 10 +++++++ 10 files changed, 68 insertions(+), 26 deletions(-) create mode 100644 clang/test/CodeGenCUDA/kernel-arg-name-metadata.cu diff --git a/clang/include/clang/Basic/CodeGenOptions.def b/clang/include/clang/Basic/CodeGenOptions.def index 8e89106..72b0e5d 100644 --- a/clang/include/clang/Basic/CodeGenOptions.def +++ b/clang/include/clang/Basic/CodeGenOptions.def @@ -187,6 +187,7 @@ CODEGENOPT(NoImplicitFloat , 1, 0) ///< Set when -mno-implicit-float is enable 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 index ee09e82..e998612 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -1007,6 +1007,12 @@ defm hip_fp32_correctly_rounded_divide_sqrt : BoolFOption<"hip-fp32-correctly-ro 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 index b4991f1..de5cb91 100644 --- a/clang/lib/CodeGen/CGDeclCXX.cpp +++ b/clang/lib/CodeGen/CGDeclCXX.cpp @@ -707,7 +707,7 @@ CodeGenModule::EmitCXXGlobalInitFunc() { // 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.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 2745b0c..4255f1c 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -596,15 +596,17 @@ CodeGenFunction::DecodeAddrUsedInPrologue(llvm::Value *F, "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 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, 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/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h index a7de87b..daf26d5 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -1968,8 +1968,7 @@ private: /// 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/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index ff3480a..f7bac66 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1697,7 +1697,7 @@ static unsigned ArgInfoAddressSpace(LangAS AS) { } } -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 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, 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 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, } 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 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, 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/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index 79e9a46..f5cbdaf 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1460,7 +1460,7 @@ public: /// \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/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 9f3199b..bcb34d0 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -6279,6 +6279,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, 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 index 0000000..f4b0075 --- /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 index c4f4366..2d6ed77 100644 --- 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" -- 2.7.4