[clang] 8ad4c6e - [HIP] add -fhip-kernel-arg-name
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 24 08:15:52 PDT 2022
Author: Yaxun (Sam) Liu
Date: 2022-06-24T11:15:36-04:00
New Revision: 8ad4c6e4b1299d599c0b6defe6a9e90a417c7ba8
URL: https://github.com/llvm/llvm-project/commit/8ad4c6e4b1299d599c0b6defe6a9e90a417c7ba8
DIFF: https://github.com/llvm/llvm-project/commit/8ad4c6e4b1299d599c0b6defe6a9e90a417c7ba8.diff
LOG: [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
Added:
clang/test/CodeGenCUDA/kernel-arg-name-metadata.cu
Modified:
clang/include/clang/Basic/CodeGenOptions.def
clang/include/clang/Driver/Options.td
clang/lib/CodeGen/CGDeclCXX.cpp
clang/lib/CodeGen/CodeGenFunction.cpp
clang/lib/CodeGen/CodeGenFunction.h
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/CodeGen/CodeGenModule.h
clang/lib/Driver/ToolChains/Clang.cpp
clang/test/Driver/hip-options.hip
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/CodeGenOptions.def b/clang/include/clang/Basic/CodeGenOptions.def
index 8e89106993c26..72b0e5d8eb41b 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 ee09e82ae4215..e998612d32e2e 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<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.">;
diff --git a/clang/lib/CodeGen/CGDeclCXX.cpp b/clang/lib/CodeGen/CGDeclCXX.cpp
index b4991f14ba193..de5cb913220a0 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 2745b0c898ff7..4255f1ca9759c 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<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();
@@ -919,9 +921,10 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy,
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
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index a7de87b552d28..daf26d54641f0 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -1968,8 +1968,7 @@ class CodeGenFunction : public CodeGenTypeCache {
/// 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 ff3480a9ac840..f7bac66c484b8 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 79e9a462a3d72..f5cbdafc8db62 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -1460,7 +1460,7 @@ class CodeGenModule : public CodeGenTypeCache {
/// \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 9f3199b33adf9..bcb34d0ffd561 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 0000000000000..f4b00757da0af
--- /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 c4f436669b0b7..2d6ed77cf4d7f 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"
More information about the cfe-commits
mailing list