[clang] 3f4b589 - [AMDGPU] Add option -munsafe-fp-atomics
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 16 18:53:04 PST 2020
Author: Yaxun (Sam) Liu
Date: 2020-11-16T21:52:12-05:00
New Revision: 3f4b5893efed620d93015896d79eb276628286f8
URL: https://github.com/llvm/llvm-project/commit/3f4b5893efed620d93015896d79eb276628286f8
DIFF: https://github.com/llvm/llvm-project/commit/3f4b5893efed620d93015896d79eb276628286f8.diff
LOG: [AMDGPU] Add option -munsafe-fp-atomics
Add an option -munsafe-fp-atomics for AMDGPU target.
When enabled, clang adds function attribute "amdgpu-unsafe-fp-atomics"
to any functions for amdgpu target. This allows amdgpu backend to use
unsafe fp atomic instructions in these functions.
Differential Revision: https://reviews.llvm.org/D91546
Added:
clang/test/CodeGenCUDA/amdgpu-func-attrs.cu
Modified:
clang/include/clang/Basic/TargetInfo.h
clang/include/clang/Basic/TargetOptions.h
clang/include/clang/Driver/Options.td
clang/lib/Basic/TargetInfo.cpp
clang/lib/Basic/Targets/AMDGPU.cpp
clang/lib/CodeGen/TargetInfo.cpp
clang/lib/Driver/ToolChains/Clang.cpp
clang/lib/Frontend/CompilerInvocation.cpp
clang/test/Driver/hip-options.hip
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h
index 26dc6eacb204..698964b94ee2 100644
--- a/clang/include/clang/Basic/TargetInfo.h
+++ b/clang/include/clang/Basic/TargetInfo.h
@@ -218,6 +218,8 @@ class TargetInfo : public virtual TransferrableTargetInfo,
unsigned HasAArch64SVETypes : 1;
+ unsigned AllowAMDGPUUnsafeFPAtomics : 1;
+
unsigned ARMCDECoprocMask : 8;
unsigned MaxOpenCLWorkGroupSize;
@@ -857,6 +859,10 @@ class TargetInfo : public virtual TransferrableTargetInfo,
/// available on this target.
bool hasAArch64SVETypes() const { return HasAArch64SVETypes; }
+ /// Returns whether or not the AMDGPU unsafe floating point atomics are
+ /// allowed.
+ bool allowAMDGPUUnsafeFPAtomics() const { return AllowAMDGPUUnsafeFPAtomics; }
+
/// For ARM targets returns a mask defining which coprocessors are configured
/// as Custom Datapath.
uint32_t getARMCDECoprocMask() const { return ARMCDECoprocMask; }
diff --git a/clang/include/clang/Basic/TargetOptions.h b/clang/include/clang/Basic/TargetOptions.h
index d1cc024957da..f81c150b7d0a 100644
--- a/clang/include/clang/Basic/TargetOptions.h
+++ b/clang/include/clang/Basic/TargetOptions.h
@@ -75,6 +75,9 @@ class TargetOptions {
/// address space.
bool NVPTXUseShortPointers = false;
+ /// \brief If enabled, allow AMDGPU unsafe floating point atomics.
+ bool AllowAMDGPUUnsafeFPAtomics = false;
+
// The code model to be used as specified by the user. Corresponds to
// CodeModel::Model enum defined in include/llvm/Support/CodeGen.h, plus
// "default" for the case when the user has not explicitly specified a
diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
index ec86c5e07ab6..0168d7000737 100644
--- a/clang/include/clang/Driver/Options.td
+++ b/clang/include/clang/Driver/Options.td
@@ -2546,6 +2546,11 @@ def mxnack : Flag<["-"], "mxnack">, Group<m_amdgpu_Features_Group>,
HelpText<"Specify XNACK mode (AMDGPU only)">;
def mno_xnack : Flag<["-"], "mno-xnack">, Group<m_amdgpu_Features_Group>;
+def munsafe_fp_atomics : Flag<["-"], "munsafe-fp-atomics">, Group<m_Group>,
+ HelpText<"Enable unsafe floating point atomic instructions (AMDGPU only)">,
+ Flags<[CC1Option]>;
+def mno_unsafe_fp_atomics : Flag<["-"], "mno-unsafe-fp-atomics">, Group<m_Group>;
+
def faltivec : Flag<["-"], "faltivec">, Group<f_Group>, Flags<[NoXarchOption]>;
def fno_altivec : Flag<["-"], "fno-altivec">, Group<f_Group>, Flags<[NoXarchOption]>;
def maltivec : Flag<["-"], "maltivec">, Group<m_ppc_Features_Group>;
diff --git a/clang/lib/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp
index eccdc21d724a..642ee753d224 100644
--- a/clang/lib/Basic/TargetInfo.cpp
+++ b/clang/lib/Basic/TargetInfo.cpp
@@ -115,6 +115,7 @@ TargetInfo::TargetInfo(const llvm::Triple &T) : TargetOpts(), Triple(T) {
HasBuiltinMSVaList = false;
IsRenderScriptTarget = false;
HasAArch64SVETypes = false;
+ AllowAMDGPUUnsafeFPAtomics = false;
ARMCDECoprocMask = 0;
// Default to no types using fpret.
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index 4d6a9a5e0b51..9b88dff7c4af 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -323,6 +323,7 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple,
HasLegalHalfType = true;
HasFloat16 = true;
WavefrontSize = GPUFeatures & llvm::AMDGPU::FEATURE_WAVE32 ? 32 : 64;
+ AllowAMDGPUUnsafeFPAtomics = Opts.AllowAMDGPUUnsafeFPAtomics;
// Set pointer width and alignment for target address space 0.
PointerWidth = PointerAlign = DataLayout->getPointerSizeInBits();
diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp
index 1e5920322ecd..a98e4095b074 100644
--- a/clang/lib/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CodeGen/TargetInfo.cpp
@@ -9080,6 +9080,9 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes(
if (NumVGPR != 0)
F->addFnAttr("amdgpu-num-vgpr", llvm::utostr(NumVGPR));
}
+
+ if (M.getContext().getTargetInfo().allowAMDGPUUnsafeFPAtomics())
+ F->addFnAttr("amdgpu-unsafe-fp-atomics", "true");
}
unsigned AMDGPUTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 9ef408f42eff..ae9e1ce61d11 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -6217,6 +6217,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
}
HandleAmdgcnLegacyOptions(D, Args, CmdArgs);
+ if (Triple.isAMDGPU()) {
+ if (Args.hasFlag(options::OPT_munsafe_fp_atomics,
+ options::OPT_mno_unsafe_fp_atomics))
+ CmdArgs.push_back("-munsafe-fp-atomics");
+ }
// For all the host OpenMP offloading compile jobs we need to pass the targets
// information using -fopenmp-targets= option.
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index c4133ec1bcd8..506423057476 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -3747,6 +3747,9 @@ static void ParseTargetArgs(TargetOptions &Opts, ArgList &Args,
Opts.ForceEnableInt128 = Args.hasArg(OPT_fforce_enable_int128);
Opts.NVPTXUseShortPointers = Args.hasFlag(
options::OPT_fcuda_short_ptr, options::OPT_fno_cuda_short_ptr, false);
+ Opts.AllowAMDGPUUnsafeFPAtomics =
+ Args.hasFlag(options::OPT_munsafe_fp_atomics,
+ options::OPT_mno_unsafe_fp_atomics, false);
if (Arg *A = Args.getLastArg(options::OPT_target_sdk_version_EQ)) {
llvm::VersionTuple Version;
if (Version.tryParse(A->getValue()))
diff --git a/clang/test/CodeGenCUDA/amdgpu-func-attrs.cu b/clang/test/CodeGenCUDA/amdgpu-func-attrs.cu
new file mode 100644
index 000000000000..6a798c67f038
--- /dev/null
+++ b/clang/test/CodeGenCUDA/amdgpu-func-attrs.cu
@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
+// RUN: | FileCheck -check-prefixes=NO-UNSAFE-FP-ATOMICS %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
+// RUN: -munsafe-fp-atomics \
+// RUN: | FileCheck -check-prefixes=UNSAFE-FP-ATOMICS %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
+// RUN: -o - -x hip %s -munsafe-fp-atomics \
+// RUN: | FileCheck -check-prefix=NO-UNSAFE-FP-ATOMICS %s
+
+#include "Inputs/cuda.h"
+
+__device__ void test() {
+// UNSAFE-FP-ATOMICS: define void @_Z4testv() [[ATTR:#[0-9]+]]
+}
+
+
+// Make sure this is silently accepted on other targets.
+// NO-UNSAFE-FP-ATOMICS-NOT: "amdgpu-unsafe-fp-atomics"
+
+// UNSAFE-FP-ATOMICS-DAG: attributes [[ATTR]] = {{.*}}"amdgpu-unsafe-fp-atomics"="true"
diff --git a/clang/test/Driver/hip-options.hip b/clang/test/Driver/hip-options.hip
index fa7b019e5762..46cfe0a531f6 100644
--- a/clang/test/Driver/hip-options.hip
+++ b/clang/test/Driver/hip-options.hip
@@ -31,3 +31,7 @@
// HOST-NOT: clang{{.*}} "-fcuda-is-device" {{.*}} "-debug-info-kind={{.*}}"
// HOST-NOT: clang{{.*}} "-fcuda-is-device" {{.*}} "-debug-info-kind={{.*}}"
// HOST: clang{{.*}} "-debug-info-kind={{.*}}"
+
+// RUN: %clang -### -nogpuinc -nogpulib -munsafe-fp-atomics \
+// RUN: --cuda-gpu-arch=gfx906 %s 2>&1 | FileCheck -check-prefix=UNSAFE-FP-ATOMICS %s
+// UNSAFE-FP-ATOMICS: clang{{.*}} "-triple" "amdgcn-amd-amdhsa" {{.*}} "-munsafe-fp-atomics"
More information about the cfe-commits
mailing list