tra updated this revision to Diff 141932.
tra added a comment.
@echristo convinced me that this functionality can be implemented without
growing a target-specific hook for custom interpretation of constraints used in
TARGET_BUILTIN. Instead, we can hide unwieldy feature lists behind a macro.
https://reviews.llvm.org/D45061
Files:
clang/include/clang/Basic/BuiltinsNVPTX.def
clang/lib/Basic/Targets/NVPTX.cpp
clang/lib/Basic/Targets/NVPTX.h
clang/test/CodeGen/builtins-nvptx-ptx50.cu
clang/test/CodeGen/builtins-nvptx.c
llvm/lib/Target/NVPTX/NVPTX.td
llvm/lib/Target/NVPTX/NVPTXSubtarget.h
Index: llvm/lib/Target/NVPTX/NVPTXSubtarget.h
===================================================================
--- llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -48,10 +48,6 @@
// FrameLowering class because TargetFrameLowering is abstract.
NVPTXFrameLowering FrameLowering;
-protected:
- // Processor supports scoped atomic operations.
- bool HasAtomScope;
-
public:
/// This constructor initializes the data members to match that
/// of the specified module.
@@ -74,7 +70,7 @@
}
bool hasAtomAddF64() const { return SmVersion >= 60; }
- bool hasAtomScope() const { return HasAtomScope; }
+ bool hasAtomScope() const { return SmVersion >= 60; }
bool hasAtomBitwise64() const { return SmVersion >= 32; }
bool hasAtomMinMax64() const { return SmVersion >= 32; }
bool hasLDG() const { return SmVersion >= 32; }
Index: llvm/lib/Target/NVPTX/NVPTX.td
===================================================================
--- llvm/lib/Target/NVPTX/NVPTX.td
+++ llvm/lib/Target/NVPTX/NVPTX.td
@@ -53,9 +53,6 @@
def SM70 : SubtargetFeature<"sm_70", "SmVersion", "70",
"Target SM 7.0">;
-def SATOM : SubtargetFeature<"satom", "HasAtomScope", "true",
- "Atomic operations with scope">;
-
// PTX Versions
def PTX32 : SubtargetFeature<"ptx32", "PTXVersion", "32",
"Use PTX version 3.2">;
@@ -88,10 +85,10 @@
def : Proc<"sm_50", [SM50, PTX40]>;
def : Proc<"sm_52", [SM52, PTX41]>;
def : Proc<"sm_53", [SM53, PTX42]>;
-def : Proc<"sm_60", [SM60, PTX50, SATOM]>;
-def : Proc<"sm_61", [SM61, PTX50, SATOM]>;
-def : Proc<"sm_62", [SM62, PTX50, SATOM]>;
-def : Proc<"sm_70", [SM70, PTX60, SATOM]>;
+def : Proc<"sm_60", [SM60, PTX50]>;
+def : Proc<"sm_61", [SM61, PTX50]>;
+def : Proc<"sm_62", [SM62, PTX50]>;
+def : Proc<"sm_70", [SM70, PTX60]>;
def NVPTXInstrInfo : InstrInfo {
}
Index: clang/test/CodeGen/builtins-nvptx.c
===================================================================
--- clang/test/CodeGen/builtins-nvptx.c
+++ clang/test/CodeGen/builtins-nvptx.c
@@ -5,6 +5,9 @@
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_60 \
// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_61 \
+// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \
+// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s
// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_53 \
// RUN: -DERROR_CHECK -fcuda-is-device -S -o /dev/null -x cuda -verify %s
@@ -292,245 +295,245 @@
#if ERROR_CHECK || __CUDA_ARCH__ >= 600
// CHECK: call i32 @llvm.nvvm.atomic.add.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_add_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_add_gen_i' needs target feature sm_60}}
__nvvm_atom_cta_add_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.add.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.add.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_add_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_add_gen_l' needs target feature sm_60}}
__nvvm_atom_cta_add_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.add.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_add_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_add_gen_ll' needs target feature sm_60}}
__nvvm_atom_cta_add_gen_ll(&sll, ll);
// CHECK: call i32 @llvm.nvvm.atomic.add.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_add_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_add_gen_i' needs target feature sm_60}}
__nvvm_atom_sys_add_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.add.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.add.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_add_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_add_gen_l' needs target feature sm_60}}
__nvvm_atom_sys_add_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.add.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_add_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_add_gen_ll' needs target feature sm_60}}
__nvvm_atom_sys_add_gen_ll(&sll, ll);
// CHECK: call float @llvm.nvvm.atomic.add.gen.f.cta.f32.p0f32
- // expected-error@+1 {{'__nvvm_atom_cta_add_gen_f' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_add_gen_f' needs target feature sm_60}}
__nvvm_atom_cta_add_gen_f(fp, f);
// CHECK: call double @llvm.nvvm.atomic.add.gen.f.cta.f64.p0f64
- // expected-error@+1 {{'__nvvm_atom_cta_add_gen_d' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_add_gen_d' needs target feature sm_60}}
__nvvm_atom_cta_add_gen_d(dfp, df);
// CHECK: call float @llvm.nvvm.atomic.add.gen.f.sys.f32.p0f32
- // expected-error@+1 {{'__nvvm_atom_sys_add_gen_f' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_add_gen_f' needs target feature sm_60}}
__nvvm_atom_sys_add_gen_f(fp, f);
// CHECK: call double @llvm.nvvm.atomic.add.gen.f.sys.f64.p0f64
- // expected-error@+1 {{'__nvvm_atom_sys_add_gen_d' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_add_gen_d' needs target feature sm_60}}
__nvvm_atom_sys_add_gen_d(dfp, df);
// CHECK: call i32 @llvm.nvvm.atomic.exch.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_i' needs target feature sm_60}}
__nvvm_atom_cta_xchg_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.exch.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.exch.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_l' needs target feature sm_60}}
__nvvm_atom_cta_xchg_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_ll' needs target feature sm_60}}
__nvvm_atom_cta_xchg_gen_ll(&sll, ll);
// CHECK: call i32 @llvm.nvvm.atomic.exch.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_i' needs target feature sm_60}}
__nvvm_atom_sys_xchg_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.exch.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.exch.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_l' needs target feature sm_60}}
__nvvm_atom_sys_xchg_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_ll' needs target feature sm_60}}
__nvvm_atom_sys_xchg_gen_ll(&sll, ll);
// CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_max_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_max_gen_i' needs target feature sm_60}}
__nvvm_atom_cta_max_gen_i(ip, i);
// CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ui' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ui' needs target feature sm_60}}
__nvvm_atom_cta_max_gen_ui((unsigned int *)ip, i);
// LP32: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_max_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_max_gen_l' needs target feature sm_60}}
__nvvm_atom_cta_max_gen_l(&dl, l);
// LP32: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ul' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ul' needs target feature sm_60}}
__nvvm_atom_cta_max_gen_ul((unsigned long *)lp, l);
// CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ll' needs target feature sm_60}}
__nvvm_atom_cta_max_gen_ll(&sll, ll);
// CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ull' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ull' needs target feature sm_60}}
__nvvm_atom_cta_max_gen_ull((unsigned long long *)llp, ll);
// CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_max_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_max_gen_i' needs target feature sm_60}}
__nvvm_atom_sys_max_gen_i(ip, i);
// CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ui' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ui' needs target feature sm_60}}
__nvvm_atom_sys_max_gen_ui((unsigned int *)ip, i);
// LP32: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_max_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_max_gen_l' needs target feature sm_60}}
__nvvm_atom_sys_max_gen_l(&dl, l);
// LP32: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ul' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ul' needs target feature sm_60}}
__nvvm_atom_sys_max_gen_ul((unsigned long *)lp, l);
// CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ll' needs target feature sm_60}}
__nvvm_atom_sys_max_gen_ll(&sll, ll);
// CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ull' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ull' needs target feature sm_60}}
__nvvm_atom_sys_max_gen_ull((unsigned long long *)llp, ll);
// CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_min_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_min_gen_i' needs target feature sm_60}}
__nvvm_atom_cta_min_gen_i(ip, i);
// CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ui' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ui' needs target feature sm_60}}
__nvvm_atom_cta_min_gen_ui((unsigned int *)ip, i);
// LP32: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_min_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_min_gen_l' needs target feature sm_60}}
__nvvm_atom_cta_min_gen_l(&dl, l);
// LP32: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ul' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ul' needs target feature sm_60}}
__nvvm_atom_cta_min_gen_ul((unsigned long *)lp, l);
// CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ll' needs target feature sm_60}}
__nvvm_atom_cta_min_gen_ll(&sll, ll);
// CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ull' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ull' needs target feature sm_60}}
__nvvm_atom_cta_min_gen_ull((unsigned long long *)llp, ll);
// CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_min_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_min_gen_i' needs target feature sm_60}}
__nvvm_atom_sys_min_gen_i(ip, i);
// CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ui' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ui' needs target feature sm_60}}
__nvvm_atom_sys_min_gen_ui((unsigned int *)ip, i);
// LP32: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_min_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_min_gen_l' needs target feature sm_60}}
__nvvm_atom_sys_min_gen_l(&dl, l);
// LP32: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ul' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ul' needs target feature sm_60}}
__nvvm_atom_sys_min_gen_ul((unsigned long *)lp, l);
// CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ll' needs target feature sm_60}}
__nvvm_atom_sys_min_gen_ll(&sll, ll);
// CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ull' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ull' needs target feature sm_60}}
__nvvm_atom_sys_min_gen_ull((unsigned long long *)llp, ll);
// CHECK: call i32 @llvm.nvvm.atomic.inc.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_inc_gen_ui' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_inc_gen_ui' needs target feature sm_60}}
__nvvm_atom_cta_inc_gen_ui((unsigned int *)ip, i);
// CHECK: call i32 @llvm.nvvm.atomic.inc.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_inc_gen_ui' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_inc_gen_ui' needs target feature sm_60}}
__nvvm_atom_sys_inc_gen_ui((unsigned int *)ip, i);
// CHECK: call i32 @llvm.nvvm.atomic.dec.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_dec_gen_ui' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_dec_gen_ui' needs target feature sm_60}}
__nvvm_atom_cta_dec_gen_ui((unsigned int *)ip, i);
// CHECK: call i32 @llvm.nvvm.atomic.dec.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_dec_gen_ui' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_dec_gen_ui' needs target feature sm_60}}
__nvvm_atom_sys_dec_gen_ui((unsigned int *)ip, i);
// CHECK: call i32 @llvm.nvvm.atomic.and.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_and_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_and_gen_i' needs target feature sm_60}}
__nvvm_atom_cta_and_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.and.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.and.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_and_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_and_gen_l' needs target feature sm_60}}
__nvvm_atom_cta_and_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.and.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_and_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_and_gen_ll' needs target feature sm_60}}
__nvvm_atom_cta_and_gen_ll(&sll, ll);
// CHECK: call i32 @llvm.nvvm.atomic.and.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_and_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_and_gen_i' needs target feature sm_60}}
__nvvm_atom_sys_and_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.and.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.and.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_and_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_and_gen_l' needs target feature sm_60}}
__nvvm_atom_sys_and_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.and.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_and_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_and_gen_ll' needs target feature sm_60}}
__nvvm_atom_sys_and_gen_ll(&sll, ll);
// CHECK: call i32 @llvm.nvvm.atomic.or.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_or_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_or_gen_i' needs target feature sm_60}}
__nvvm_atom_cta_or_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.or.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.or.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_or_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_or_gen_l' needs target feature sm_60}}
__nvvm_atom_cta_or_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.or.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_or_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_or_gen_ll' needs target feature sm_60}}
__nvvm_atom_cta_or_gen_ll(&sll, ll);
// CHECK: call i32 @llvm.nvvm.atomic.or.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_or_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_or_gen_i' needs target feature sm_60}}
__nvvm_atom_sys_or_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.or.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.or.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_or_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_or_gen_l' needs target feature sm_60}}
__nvvm_atom_sys_or_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.or.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_or_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_or_gen_ll' needs target feature sm_60}}
__nvvm_atom_sys_or_gen_ll(&sll, ll);
// CHECK: call i32 @llvm.nvvm.atomic.xor.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_i' needs target feature sm_60}}
__nvvm_atom_cta_xor_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.xor.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.xor.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_l' needs target feature sm_60}}
__nvvm_atom_cta_xor_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.xor.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_ll' needs target feature sm_60}}
__nvvm_atom_cta_xor_gen_ll(&sll, ll);
// CHECK: call i32 @llvm.nvvm.atomic.xor.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_i' needs target feature sm_60}}
__nvvm_atom_sys_xor_gen_i(ip, i);
// LP32: call i32 @llvm.nvvm.atomic.xor.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.xor.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_l' needs target feature sm_60}}
__nvvm_atom_sys_xor_gen_l(&dl, l);
// CHECK: call i64 @llvm.nvvm.atomic.xor.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_ll' needs target feature sm_60}}
__nvvm_atom_sys_xor_gen_ll(&sll, ll);
// CHECK: call i32 @llvm.nvvm.atomic.cas.gen.i.cta.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_i' needs target feature sm_60}}
__nvvm_atom_cta_cas_gen_i(ip, i, 0);
// LP32: call i32 @llvm.nvvm.atomic.cas.gen.i.cta.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.cas.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_l' needs target feature sm_60}}
__nvvm_atom_cta_cas_gen_l(&dl, l, 0);
// CHECK: call i64 @llvm.nvvm.atomic.cas.gen.i.cta.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_ll' needs target feature sm_60}}
__nvvm_atom_cta_cas_gen_ll(&sll, ll, 0);
// CHECK: call i32 @llvm.nvvm.atomic.cas.gen.i.sys.i32.p0i32
- // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_i' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_i' needs target feature sm_60}}
__nvvm_atom_sys_cas_gen_i(ip, i, 0);
// LP32: call i32 @llvm.nvvm.atomic.cas.gen.i.sys.i32.p0i32
// LP64: call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_l' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_l' needs target feature sm_60}}
__nvvm_atom_sys_cas_gen_l(&dl, l, 0);
// CHECK: call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0i64
- // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_ll' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_ll' needs target feature sm_60}}
__nvvm_atom_sys_cas_gen_ll(&sll, ll, 0);
#endif
Index: clang/test/CodeGen/builtins-nvptx-ptx50.cu
===================================================================
--- clang/test/CodeGen/builtins-nvptx-ptx50.cu
+++ clang/test/CodeGen/builtins-nvptx-ptx50.cu
@@ -18,6 +18,6 @@
// CHECK-LABEL: test_fn
__device__ void test_fn(double d, double* double_ptr) {
// CHECK: call double @llvm.nvvm.atomic.load.add.f64.p0f64
- // expected-error@+1 {{'__nvvm_atom_add_gen_d' needs target feature satom}}
+ // expected-error@+1 {{'__nvvm_atom_add_gen_d' needs target feature sm_60}}
__nvvm_atom_add_gen_d(double_ptr, d);
}
Index: clang/lib/Basic/Targets/NVPTX.h
===================================================================
--- clang/lib/Basic/Targets/NVPTX.h
+++ clang/lib/Basic/Targets/NVPTX.h
@@ -40,6 +40,7 @@
static const char *const GCCRegNames[];
static const Builtin::Info BuiltinInfo[];
CudaArch GPU;
+ uint32_t PTXVersion;
std::unique_ptr<TargetInfo> HostTarget;
public:
@@ -55,7 +56,8 @@
initFeatureMap(llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags,
StringRef CPU,
const std::vector<std::string> &FeaturesVec) const override {
- Features["satom"] = GPU >= CudaArch::SM_60;
+ Features[CudaArchToString(GPU)] = true;
+ Features["ptx" + std::to_string(PTXVersion)] = true;
return TargetInfo::initFeatureMap(Features, Diags, CPU, FeaturesVec);
}
Index: clang/lib/Basic/Targets/NVPTX.cpp
===================================================================
--- clang/lib/Basic/Targets/NVPTX.cpp
+++ clang/lib/Basic/Targets/NVPTX.cpp
@@ -40,6 +40,22 @@
assert((TargetPointerWidth == 32 || TargetPointerWidth == 64) &&
"NVPTX only supports 32- and 64-bit modes.");
+ PTXVersion = 32;
+ for (const StringRef Feature : Opts.FeaturesAsWritten) {
+ if (!Feature.startswith("+ptx"))
+ continue;
+ PTXVersion = llvm::StringSwitch<unsigned>(Feature)
+ .Case("+ptx61", 61)
+ .Case("+ptx60", 60)
+ .Case("+ptx50", 50)
+ .Case("+ptx43", 43)
+ .Case("+ptx42", 42)
+ .Case("+ptx41", 41)
+ .Case("+ptx40", 40)
+ .Case("+ptx32", 32)
+ .Default(32);
+ }
+
TLSSupported = false;
VLASupported = false;
AddrSpaceMap = &NVPTXAddrSpaceMap;
@@ -145,7 +161,6 @@
bool NVPTXTargetInfo::hasFeature(StringRef Feature) const {
return llvm::StringSwitch<bool>(Feature)
.Cases("ptx", "nvptx", true)
- .Case("satom", GPU >= CudaArch::SM_60) // Atomics w/ scope.
.Default(false);
}
Index: clang/include/clang/Basic/BuiltinsNVPTX.def
===================================================================
--- clang/include/clang/Basic/BuiltinsNVPTX.def
+++ clang/include/clang/Basic/BuiltinsNVPTX.def
@@ -18,6 +18,12 @@
# define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
#endif
+#pragma push_macro("SM_60")
+#define SM_60 "sm_60|sm_61|sm_62|sm_70|sm_71"
+
+#pragma push_macro("PTX60")
+#define PTX60 "ptx60|ptx61"
+
// Special Registers
BUILTIN(__nvvm_read_ptx_sreg_tid_x, "i", "nc")
@@ -372,18 +378,18 @@
BUILTIN(__nvvm_bitcast_d2ll, "LLid", "")
// FNS
-TARGET_BUILTIN(__nvvm_fns, "UiUiUii", "n", "ptx60")
+TARGET_BUILTIN(__nvvm_fns, "UiUiUii", "n", PTX60)
// Sync
BUILTIN(__syncthreads, "v", "")
BUILTIN(__nvvm_bar0_popc, "ii", "")
BUILTIN(__nvvm_bar0_and, "ii", "")
BUILTIN(__nvvm_bar0_or, "ii", "")
BUILTIN(__nvvm_bar_sync, "vi", "n")
-TARGET_BUILTIN(__nvvm_bar_warp_sync, "vUi", "n", "ptx60")
-TARGET_BUILTIN(__nvvm_barrier_sync, "vUi", "n", "ptx60")
-TARGET_BUILTIN(__nvvm_barrier_sync_cnt, "vUiUi", "n", "ptx60")
+TARGET_BUILTIN(__nvvm_bar_warp_sync, "vUi", "n", PTX60)
+TARGET_BUILTIN(__nvvm_barrier_sync, "vUi", "n", PTX60)
+TARGET_BUILTIN(__nvvm_barrier_sync_cnt, "vUiUi", "n", PTX60)
// Shuffle
@@ -396,32 +402,32 @@
BUILTIN(__nvvm_shfl_idx_i32, "iiii", "")
BUILTIN(__nvvm_shfl_idx_f32, "ffii", "")
-TARGET_BUILTIN(__nvvm_shfl_sync_down_i32, "iUiiii", "", "ptx60")
-TARGET_BUILTIN(__nvvm_shfl_sync_down_f32, "fUifii", "", "ptx60")
-TARGET_BUILTIN(__nvvm_shfl_sync_up_i32, "iUiiii", "", "ptx60")
-TARGET_BUILTIN(__nvvm_shfl_sync_up_f32, "fUifii", "", "ptx60")
-TARGET_BUILTIN(__nvvm_shfl_sync_bfly_i32, "iUiiii", "", "ptx60")
-TARGET_BUILTIN(__nvvm_shfl_sync_bfly_f32, "fUifii", "", "ptx60")
-TARGET_BUILTIN(__nvvm_shfl_sync_idx_i32, "iUiiii", "", "ptx60")
-TARGET_BUILTIN(__nvvm_shfl_sync_idx_f32, "fUifii", "", "ptx60")
+TARGET_BUILTIN(__nvvm_shfl_sync_down_i32, "iUiiii", "", PTX60)
+TARGET_BUILTIN(__nvvm_shfl_sync_down_f32, "fUifii", "", PTX60)
+TARGET_BUILTIN(__nvvm_shfl_sync_up_i32, "iUiiii", "", PTX60)
+TARGET_BUILTIN(__nvvm_shfl_sync_up_f32, "fUifii", "", PTX60)
+TARGET_BUILTIN(__nvvm_shfl_sync_bfly_i32, "iUiiii", "", PTX60)
+TARGET_BUILTIN(__nvvm_shfl_sync_bfly_f32, "fUifii", "", PTX60)
+TARGET_BUILTIN(__nvvm_shfl_sync_idx_i32, "iUiiii", "", PTX60)
+TARGET_BUILTIN(__nvvm_shfl_sync_idx_f32, "fUifii", "", PTX60)
// Vote
BUILTIN(__nvvm_vote_all, "bb", "")
BUILTIN(__nvvm_vote_any, "bb", "")
BUILTIN(__nvvm_vote_uni, "bb", "")
BUILTIN(__nvvm_vote_ballot, "Uib", "")
-TARGET_BUILTIN(__nvvm_vote_all_sync, "bUib", "", "ptx60")
-TARGET_BUILTIN(__nvvm_vote_any_sync, "bUib", "", "ptx60")
-TARGET_BUILTIN(__nvvm_vote_uni_sync, "bUib", "", "ptx60")
-TARGET_BUILTIN(__nvvm_vote_ballot_sync, "UiUib", "", "ptx60")
+TARGET_BUILTIN(__nvvm_vote_all_sync, "bUib", "", PTX60)
+TARGET_BUILTIN(__nvvm_vote_any_sync, "bUib", "", PTX60)
+TARGET_BUILTIN(__nvvm_vote_uni_sync, "bUib", "", PTX60)
+TARGET_BUILTIN(__nvvm_vote_ballot_sync, "UiUib", "", PTX60)
// Match
-TARGET_BUILTIN(__nvvm_match_any_sync_i32, "UiUiUi", "", "ptx60")
-TARGET_BUILTIN(__nvvm_match_any_sync_i64, "WiUiWi", "", "ptx60")
+TARGET_BUILTIN(__nvvm_match_any_sync_i32, "UiUiUi", "", PTX60)
+TARGET_BUILTIN(__nvvm_match_any_sync_i64, "WiUiWi", "", PTX60)
// These return a pair {value, predicate}, which requires custom lowering.
-TARGET_BUILTIN(__nvvm_match_all_sync_i32p, "UiUiUii*", "", "ptx60")
-TARGET_BUILTIN(__nvvm_match_all_sync_i64p, "WiUiWii*", "", "ptx60")
+TARGET_BUILTIN(__nvvm_match_all_sync_i32p, "UiUiUii*", "", PTX60)
+TARGET_BUILTIN(__nvvm_match_all_sync_i64p, "WiUiWii*", "", PTX60)
// Membar
@@ -465,28 +471,28 @@
BUILTIN(__nvvm_atom_add_g_i, "iiD*1i", "n")
BUILTIN(__nvvm_atom_add_s_i, "iiD*3i", "n")
BUILTIN(__nvvm_atom_add_gen_i, "iiD*i", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_add_gen_i, "iiD*i", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_add_gen_i, "iiD*i", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_add_gen_i, "iiD*i", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_add_gen_i, "iiD*i", "n", SM_60)
BUILTIN(__nvvm_atom_add_g_l, "LiLiD*1Li", "n")
BUILTIN(__nvvm_atom_add_s_l, "LiLiD*3Li", "n")
BUILTIN(__nvvm_atom_add_gen_l, "LiLiD*Li", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_add_gen_l, "LiLiD*Li", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_add_gen_l, "LiLiD*Li", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_add_gen_l, "LiLiD*Li", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_add_gen_l, "LiLiD*Li", "n", SM_60)
BUILTIN(__nvvm_atom_add_g_ll, "LLiLLiD*1LLi", "n")
BUILTIN(__nvvm_atom_add_s_ll, "LLiLLiD*3LLi", "n")
BUILTIN(__nvvm_atom_add_gen_ll, "LLiLLiD*LLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_add_gen_ll, "LLiLLiD*LLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_add_gen_ll, "LLiLLiD*LLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_add_gen_ll, "LLiLLiD*LLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_add_gen_ll, "LLiLLiD*LLi", "n", SM_60)
BUILTIN(__nvvm_atom_add_g_f, "ffD*1f", "n")
BUILTIN(__nvvm_atom_add_s_f, "ffD*3f", "n")
BUILTIN(__nvvm_atom_add_gen_f, "ffD*f", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_add_gen_f, "ffD*f", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_add_gen_f, "ffD*f", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_add_gen_f, "ffD*f", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_add_gen_f, "ffD*f", "n", SM_60)
BUILTIN(__nvvm_atom_add_g_d, "ddD*1d", "n")
BUILTIN(__nvvm_atom_add_s_d, "ddD*3d", "n")
-TARGET_BUILTIN(__nvvm_atom_add_gen_d, "ddD*d", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_cta_add_gen_d, "ddD*d", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_add_gen_d, "ddD*d", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_add_gen_d, "ddD*d", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_cta_add_gen_d, "ddD*d", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_add_gen_d, "ddD*d", "n", SM_60)
BUILTIN(__nvvm_atom_sub_g_i, "iiD*1i", "n")
BUILTIN(__nvvm_atom_sub_s_i, "iiD*3i", "n")
@@ -501,155 +507,155 @@
BUILTIN(__nvvm_atom_xchg_g_i, "iiD*1i", "n")
BUILTIN(__nvvm_atom_xchg_s_i, "iiD*3i", "n")
BUILTIN(__nvvm_atom_xchg_gen_i, "iiD*i", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_i, "iiD*i", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_i, "iiD*i", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_i, "iiD*i", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_i, "iiD*i", "n", SM_60)
BUILTIN(__nvvm_atom_xchg_g_l, "LiLiD*1Li", "n")
BUILTIN(__nvvm_atom_xchg_s_l, "LiLiD*3Li", "n")
BUILTIN(__nvvm_atom_xchg_gen_l, "LiLiD*Li", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_l, "LiLiD*Li", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_l, "LiLiD*Li", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_l, "LiLiD*Li", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_l, "LiLiD*Li", "n", SM_60)
BUILTIN(__nvvm_atom_xchg_g_ll, "LLiLLiD*1LLi", "n")
BUILTIN(__nvvm_atom_xchg_s_ll, "LLiLLiD*3LLi", "n")
BUILTIN(__nvvm_atom_xchg_gen_ll, "LLiLLiD*LLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_ll, "LLiLLiD*LLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_ll, "LLiLLiD*LLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_60)
BUILTIN(__nvvm_atom_max_g_i, "iiD*1i", "n")
BUILTIN(__nvvm_atom_max_s_i, "iiD*3i", "n")
BUILTIN(__nvvm_atom_max_gen_i, "iiD*i", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_max_gen_i, "iiD*i", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_max_gen_i, "iiD*i", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_max_gen_i, "iiD*i", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_max_gen_i, "iiD*i", "n", SM_60)
BUILTIN(__nvvm_atom_max_g_ui, "UiUiD*1Ui", "n")
BUILTIN(__nvvm_atom_max_s_ui, "UiUiD*3Ui", "n")
BUILTIN(__nvvm_atom_max_gen_ui, "UiUiD*Ui", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ui, "UiUiD*Ui", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ui, "UiUiD*Ui", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ui, "UiUiD*Ui", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ui, "UiUiD*Ui", "n", SM_60)
BUILTIN(__nvvm_atom_max_g_l, "LiLiD*1Li", "n")
BUILTIN(__nvvm_atom_max_s_l, "LiLiD*3Li", "n")
BUILTIN(__nvvm_atom_max_gen_l, "LiLiD*Li", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_max_gen_l, "LiLiD*Li", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_max_gen_l, "LiLiD*Li", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_max_gen_l, "LiLiD*Li", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_max_gen_l, "LiLiD*Li", "n", SM_60)
BUILTIN(__nvvm_atom_max_g_ul, "ULiULiD*1ULi", "n")
BUILTIN(__nvvm_atom_max_s_ul, "ULiULiD*3ULi", "n")
BUILTIN(__nvvm_atom_max_gen_ul, "ULiULiD*ULi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ul, "ULiULiD*ULi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ul, "ULiULiD*ULi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ul, "ULiULiD*ULi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ul, "ULiULiD*ULi", "n", SM_60)
BUILTIN(__nvvm_atom_max_g_ll, "LLiLLiD*1LLi", "n")
BUILTIN(__nvvm_atom_max_s_ll, "LLiLLiD*3LLi", "n")
BUILTIN(__nvvm_atom_max_gen_ll, "LLiLLiD*LLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ll, "LLiLLiD*LLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ll, "LLiLLiD*LLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ll, "LLiLLiD*LLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ll, "LLiLLiD*LLi", "n", SM_60)
BUILTIN(__nvvm_atom_max_g_ull, "ULLiULLiD*1ULLi", "n")
BUILTIN(__nvvm_atom_max_s_ull, "ULLiULLiD*3ULLi", "n")
BUILTIN(__nvvm_atom_max_gen_ull, "ULLiULLiD*ULLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ull, "ULLiULLiD*ULLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ull, "ULLiULLiD*ULLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ull, "ULLiULLiD*ULLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ull, "ULLiULLiD*ULLi", "n", SM_60)
BUILTIN(__nvvm_atom_min_g_i, "iiD*1i", "n")
BUILTIN(__nvvm_atom_min_s_i, "iiD*3i", "n")
BUILTIN(__nvvm_atom_min_gen_i, "iiD*i", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_min_gen_i, "iiD*i", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_min_gen_i, "iiD*i", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_min_gen_i, "iiD*i", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_min_gen_i, "iiD*i", "n", SM_60)
BUILTIN(__nvvm_atom_min_g_ui, "UiUiD*1Ui", "n")
BUILTIN(__nvvm_atom_min_s_ui, "UiUiD*3Ui", "n")
BUILTIN(__nvvm_atom_min_gen_ui, "UiUiD*Ui", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ui, "UiUiD*Ui", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ui, "UiUiD*Ui", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ui, "UiUiD*Ui", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ui, "UiUiD*Ui", "n", SM_60)
BUILTIN(__nvvm_atom_min_g_l, "LiLiD*1Li", "n")
BUILTIN(__nvvm_atom_min_s_l, "LiLiD*3Li", "n")
BUILTIN(__nvvm_atom_min_gen_l, "LiLiD*Li", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_min_gen_l, "LiLiD*Li", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_min_gen_l, "LiLiD*Li", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_min_gen_l, "LiLiD*Li", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_min_gen_l, "LiLiD*Li", "n", SM_60)
BUILTIN(__nvvm_atom_min_g_ul, "ULiULiD*1ULi", "n")
BUILTIN(__nvvm_atom_min_s_ul, "ULiULiD*3ULi", "n")
BUILTIN(__nvvm_atom_min_gen_ul, "ULiULiD*ULi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ul, "ULiULiD*ULi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ul, "ULiULiD*ULi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ul, "ULiULiD*ULi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ul, "ULiULiD*ULi", "n", SM_60)
BUILTIN(__nvvm_atom_min_g_ll, "LLiLLiD*1LLi", "n")
BUILTIN(__nvvm_atom_min_s_ll, "LLiLLiD*3LLi", "n")
BUILTIN(__nvvm_atom_min_gen_ll, "LLiLLiD*LLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ll, "LLiLLiD*LLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ll, "LLiLLiD*LLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ll, "LLiLLiD*LLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ll, "LLiLLiD*LLi", "n", SM_60)
BUILTIN(__nvvm_atom_min_g_ull, "ULLiULLiD*1ULLi", "n")
BUILTIN(__nvvm_atom_min_s_ull, "ULLiULLiD*3ULLi", "n")
BUILTIN(__nvvm_atom_min_gen_ull, "ULLiULLiD*ULLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ull, "ULLiULLiD*ULLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ull, "ULLiULLiD*ULLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ull, "ULLiULLiD*ULLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ull, "ULLiULLiD*ULLi", "n", SM_60)
BUILTIN(__nvvm_atom_inc_g_ui, "UiUiD*1Ui", "n")
BUILTIN(__nvvm_atom_inc_s_ui, "UiUiD*3Ui", "n")
BUILTIN(__nvvm_atom_inc_gen_ui, "UiUiD*Ui", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_inc_gen_ui, "UiUiD*Ui", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_inc_gen_ui, "UiUiD*Ui", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_inc_gen_ui, "UiUiD*Ui", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_inc_gen_ui, "UiUiD*Ui", "n", SM_60)
BUILTIN(__nvvm_atom_dec_g_ui, "UiUiD*1Ui", "n")
BUILTIN(__nvvm_atom_dec_s_ui, "UiUiD*3Ui", "n")
BUILTIN(__nvvm_atom_dec_gen_ui, "UiUiD*Ui", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_dec_gen_ui, "UiUiD*Ui", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_dec_gen_ui, "UiUiD*Ui", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_dec_gen_ui, "UiUiD*Ui", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_dec_gen_ui, "UiUiD*Ui", "n", SM_60)
BUILTIN(__nvvm_atom_and_g_i, "iiD*1i", "n")
BUILTIN(__nvvm_atom_and_s_i, "iiD*3i", "n")
BUILTIN(__nvvm_atom_and_gen_i, "iiD*i", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_and_gen_i, "iiD*i", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_and_gen_i, "iiD*i", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_and_gen_i, "iiD*i", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_and_gen_i, "iiD*i", "n", SM_60)
BUILTIN(__nvvm_atom_and_g_l, "LiLiD*1Li", "n")
BUILTIN(__nvvm_atom_and_s_l, "LiLiD*3Li", "n")
BUILTIN(__nvvm_atom_and_gen_l, "LiLiD*Li", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_and_gen_l, "LiLiD*Li", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_and_gen_l, "LiLiD*Li", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_and_gen_l, "LiLiD*Li", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_and_gen_l, "LiLiD*Li", "n", SM_60)
BUILTIN(__nvvm_atom_and_g_ll, "LLiLLiD*1LLi", "n")
BUILTIN(__nvvm_atom_and_s_ll, "LLiLLiD*3LLi", "n")
BUILTIN(__nvvm_atom_and_gen_ll, "LLiLLiD*LLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_and_gen_ll, "LLiLLiD*LLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_and_gen_ll, "LLiLLiD*LLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_and_gen_ll, "LLiLLiD*LLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_and_gen_ll, "LLiLLiD*LLi", "n", SM_60)
BUILTIN(__nvvm_atom_or_g_i, "iiD*1i", "n")
BUILTIN(__nvvm_atom_or_s_i, "iiD*3i", "n")
BUILTIN(__nvvm_atom_or_gen_i, "iiD*i", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_or_gen_i, "iiD*i", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_or_gen_i, "iiD*i", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_or_gen_i, "iiD*i", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_or_gen_i, "iiD*i", "n", SM_60)
BUILTIN(__nvvm_atom_or_g_l, "LiLiD*1Li", "n")
BUILTIN(__nvvm_atom_or_s_l, "LiLiD*3Li", "n")
BUILTIN(__nvvm_atom_or_gen_l, "LiLiD*Li", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_or_gen_l, "LiLiD*Li", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_or_gen_l, "LiLiD*Li", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_or_gen_l, "LiLiD*Li", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_or_gen_l, "LiLiD*Li", "n", SM_60)
BUILTIN(__nvvm_atom_or_g_ll, "LLiLLiD*1LLi", "n")
BUILTIN(__nvvm_atom_or_s_ll, "LLiLLiD*3LLi", "n")
BUILTIN(__nvvm_atom_or_gen_ll, "LLiLLiD*LLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_or_gen_ll, "LLiLLiD*LLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_or_gen_ll, "LLiLLiD*LLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_or_gen_ll, "LLiLLiD*LLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_or_gen_ll, "LLiLLiD*LLi", "n", SM_60)
BUILTIN(__nvvm_atom_xor_g_i, "iiD*1i", "n")
BUILTIN(__nvvm_atom_xor_s_i, "iiD*3i", "n")
BUILTIN(__nvvm_atom_xor_gen_i, "iiD*i", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_i, "iiD*i", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_i, "iiD*i", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_i, "iiD*i", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_i, "iiD*i", "n", SM_60)
BUILTIN(__nvvm_atom_xor_g_l, "LiLiD*1Li", "n")
BUILTIN(__nvvm_atom_xor_s_l, "LiLiD*3Li", "n")
BUILTIN(__nvvm_atom_xor_gen_l, "LiLiD*Li", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_l, "LiLiD*Li", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_l, "LiLiD*Li", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_l, "LiLiD*Li", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_l, "LiLiD*Li", "n", SM_60)
BUILTIN(__nvvm_atom_xor_g_ll, "LLiLLiD*1LLi", "n")
BUILTIN(__nvvm_atom_xor_s_ll, "LLiLLiD*3LLi", "n")
BUILTIN(__nvvm_atom_xor_gen_ll, "LLiLLiD*LLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_ll, "LLiLLiD*LLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_ll, "LLiLLiD*LLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_ll, "LLiLLiD*LLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_ll, "LLiLLiD*LLi", "n", SM_60)
BUILTIN(__nvvm_atom_cas_g_i, "iiD*1ii", "n")
BUILTIN(__nvvm_atom_cas_s_i, "iiD*3ii", "n")
BUILTIN(__nvvm_atom_cas_gen_i, "iiD*ii", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_i, "iiD*ii", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_i, "iiD*ii", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_i, "iiD*ii", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_i, "iiD*ii", "n", SM_60)
BUILTIN(__nvvm_atom_cas_g_l, "LiLiD*1LiLi", "n")
BUILTIN(__nvvm_atom_cas_s_l, "LiLiD*3LiLi", "n")
BUILTIN(__nvvm_atom_cas_gen_l, "LiLiD*LiLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_l, "LiLiD*LiLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_l, "LiLiD*LiLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_l, "LiLiD*LiLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_l, "LiLiD*LiLi", "n", SM_60)
BUILTIN(__nvvm_atom_cas_g_ll, "LLiLLiD*1LLiLLi", "n")
BUILTIN(__nvvm_atom_cas_s_ll, "LLiLLiD*3LLiLLi", "n")
BUILTIN(__nvvm_atom_cas_gen_ll, "LLiLLiD*LLiLLi", "n")
-TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_ll, "LLiLLiD*LLiLLi", "n", "satom")
-TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_ll, "LLiLLiD*LLiLLi", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_60)
+TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_60)
// Compiler Error Warn
BUILTIN(__nvvm_compiler_error, "vcC*4", "n")
@@ -692,17 +698,19 @@
BUILTIN(__nvvm_ldg_d2, "E2dE2dC*", "")
// Builtins to support WMMA instructions on sm_70
-TARGET_BUILTIN(__hmma_m16n16k16_ld_a, "vi*iC*UiIi", "", "ptx60")
-TARGET_BUILTIN(__hmma_m16n16k16_ld_b, "vi*iC*UiIi", "", "ptx60")
-TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f16, "vi*iC*UiIi", "", "ptx60")
-TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f32, "vf*fC*UiIi", "", "ptx60")
-TARGET_BUILTIN(__hmma_m16n16k16_st_c_f16, "vi*i*UiIi", "", "ptx60")
-TARGET_BUILTIN(__hmma_m16n16k16_st_c_f32, "vf*f*UiIi", "", "ptx60")
+TARGET_BUILTIN(__hmma_m16n16k16_ld_a, "vi*iC*UiIi", "", PTX60)
+TARGET_BUILTIN(__hmma_m16n16k16_ld_b, "vi*iC*UiIi", "", PTX60)
+TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f16, "vi*iC*UiIi", "", PTX60)
+TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f32, "vf*fC*UiIi", "", PTX60)
+TARGET_BUILTIN(__hmma_m16n16k16_st_c_f16, "vi*i*UiIi", "", PTX60)
+TARGET_BUILTIN(__hmma_m16n16k16_st_c_f32, "vf*f*UiIi", "", PTX60)
-TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f16, "vi*iC*iC*iC*IiIi", "", "ptx60")
-TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f16, "vf*iC*iC*iC*IiIi", "", "ptx60")
-TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f32, "vf*iC*iC*fC*IiIi", "", "ptx60")
-TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f32, "vi*iC*iC*fC*IiIi", "", "ptx60")
+TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f16, "vi*iC*iC*iC*IiIi", "", PTX60)
+TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f16, "vf*iC*iC*iC*IiIi", "", PTX60)
+TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f32, "vf*iC*iC*fC*IiIi", "", PTX60)
+TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f32, "vi*iC*iC*fC*IiIi", "", PTX60)
#undef BUILTIN
#undef TARGET_BUILTIN
+#pragma pop_macro("SM_60")
+#pragma pop_macro("PTX60")
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
