[llvm] [LLVM][NVPTX] Enable family specific support for a few intrinsics (PR #173268)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Dec 22 07:09:51 PST 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Pradeep Kumar (schwarzschild-radius)
<details>
<summary>Changes</summary>
This commit adds support for family specific support for the following intrinsics:
- ldmatrix
- stmatrix
- mma.block_scale, mma.sp.block_scale
- redux.sync
- cvt.rs
- clusterlaunchcontrol
- setmaxnreg
- tcgen05.mma
Removed `hasTcgen05Instructions` function in the favour of `hasTcgen05InstSupport` Updated wmma.py script with family specific support and added new tests
---
Patch is 222.70 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/173268.diff
25 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+1-2)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+55-39)
- (modified) llvm/lib/Target/NVPTX/NVPTXSubtarget.h (+42-24)
- (modified) llvm/test/CodeGen/NVPTX/clusterlaunchcontrol-multicast.ll (+17)
- (modified) llvm/test/CodeGen/NVPTX/redux-sync-f32.ll (+2)
- (modified) llvm/test/CodeGen/NVPTX/setmaxnreg-sm100a.ll (+6)
- (added) llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-ptx88-aa.ll (+481)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-mma-block-scale-ptx88.ll (+4-473)
- (added) llvm/test/CodeGen/NVPTX/tcgen05-mma-disable-output-lane-i8.ll (+219)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-mma-disable-output-lane.ll (+4-212)
- (added) llvm/test/CodeGen/NVPTX/tcgen05-mma-i8.ll (+166)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-mma-scale-d.ll (+2)
- (added) llvm/test/CodeGen/NVPTX/tcgen05-mma-ws-i8.ll (+168)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-mma-ws.ll (+4-80)
- (modified) llvm/test/CodeGen/NVPTX/tcgen05-mma.ll (+4-158)
- (modified) llvm/test/CodeGen/NVPTX/wmma-ptx78-sm90.py (+1-1)
- (modified) llvm/test/CodeGen/NVPTX/wmma-ptx86-sm100a.py (+1-1)
- (modified) llvm/test/CodeGen/NVPTX/wmma-ptx86-sm101a.py (+1-1)
- (modified) llvm/test/CodeGen/NVPTX/wmma-ptx86-sm120a.py (+1-1)
- (modified) llvm/test/CodeGen/NVPTX/wmma-ptx87-sm120a.py (+1-1)
- (added) llvm/test/CodeGen/NVPTX/wmma-ptx88-sm100f.py (+12)
- (modified) llvm/test/CodeGen/NVPTX/wmma-ptx88-sm120a.py (+1-1)
- (added) llvm/test/CodeGen/NVPTX/wmma-ptx88-sm120f.py (+12)
- (added) llvm/test/CodeGen/NVPTX/wmma-ptx90-sm110f.py (+12)
- (modified) llvm/test/CodeGen/NVPTX/wmma.py (+53-26)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 74a552502ccf2..8c07d96fcd010 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -128,6 +128,7 @@ def hasPTXASUnreachableBug : Predicate<"Subtarget->hasPTXASUnreachableBug()">;
def noPTXASUnreachableBug : Predicate<"!Subtarget->hasPTXASUnreachableBug()">;
def hasOptEnabled : Predicate<"TM.getOptLevel() != CodeGenOptLevel::None">;
def hasArchAccelFeatures : Predicate<"Subtarget->hasArchAccelFeatures()">;
+def hasFamilySpecificFeatures : Predicate<"Subtarget->hasFamilySpecificFeatures()">;
def doF32FTZ : Predicate<"useF32FTZ()">;
def doNoF32FTZ : Predicate<"!useF32FTZ()">;
@@ -137,8 +138,6 @@ def doMADWideOpt : Predicate<"doMADWideOpt()">;
def hasHWROT32 : Predicate<"Subtarget->hasHWROT32()">;
def noHWROT32 : Predicate<"!Subtarget->hasHWROT32()">;
def hasDotInstructions : Predicate<"Subtarget->hasDotInstructions()">;
-def hasTcgen05Instructions : Predicate<"Subtarget->hasTcgen05Instructions()">;
-def hasTcgen05MMAScaleInputDImm : Predicate<"Subtarget->hasTcgen05MMAScaleInputDImm()">;
def hasF32x2Instructions : Predicate<"Subtarget->hasF32x2Instructions()">;
class hasPTX<int version>: Predicate<"Subtarget->getPTXVersion() >= " # version>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 2f6894867c43d..bf227ded5d861 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -340,7 +340,7 @@ multiclass REDUX_SYNC_F<string BinOp, string abs, string NaN> {
(ins B32:$src, B32:$mask),
"redux.sync." # BinOp # abs # NaN # ".f32",
[(set f32:$dst, (!cast<Intrinsic>(intr_name) f32:$src, B32:$mask))]>,
- Requires<[hasPTX<86>, hasSM100a]>;
+ Requires<[callSubtarget<"hasReduxSyncF32">]>;
}
defm REDUX_SYNC_FMIN : REDUX_SYNC_F<"min", "", "">;
@@ -2072,7 +2072,7 @@ let Predicates = [hasPTX<81>, hasSM<80>] in {
def : Pat<(int_nvvm_ff2bf16x2_rz_satfinite f32:$a, f32:$b), (CVT_bf16x2_f32_sf $a, $b, CvtRZ)>;
def : Pat<(int_nvvm_ff2bf16x2_rz_relu_satfinite f32:$a, f32:$b), (CVT_bf16x2_f32_sf $a, $b, CvtRZ_RELU)>;
}
-let Predicates = [hasPTX<87>, hasSM100aOrSM103a] in {
+let Predicates = [callSubtarget<"hasCVTRs">] in {
def : Pat<(int_nvvm_ff2bf16x2_rs f32:$a, f32:$b, i32:$c),
(CVT_bf16x2_f32_rs $a, $b, $c, CvtRS)>;
def : Pat<(int_nvvm_ff2bf16x2_rs_relu f32:$a, f32:$b, i32:$c),
@@ -2094,7 +2094,7 @@ let Predicates = [hasPTX<81>, hasSM<80>] in {
def : Pat<(int_nvvm_ff2f16x2_rz_relu_satfinite f32:$a, f32:$b), (CVT_f16x2_f32_sf $a, $b, CvtRZ_RELU)>;
}
-let Predicates = [hasPTX<87>, hasSM100aOrSM103a] in {
+let Predicates = [callSubtarget<"hasCVTRs">] in {
def : Pat<(int_nvvm_ff2f16x2_rs f32:$a, f32:$b, i32:$c),
(CVT_f16x2_f32_rs $a, $b, $c, CvtRS)>;
def : Pat<(int_nvvm_ff2f16x2_rs_relu f32:$a, f32:$b, i32:$c),
@@ -2291,7 +2291,7 @@ multiclass CVT_F32X4_TO_FPX4_RS_SF_VEC<string FPName, VTVec RetTy> {
}
// RS rounding mode conversions
-let Predicates = [hasPTX<87>, hasSM100aOrSM103a] in {
+let Predicates = [callSubtarget<"hasCVTRs">] in {
// FP8x4 conversions
defm : CVT_F32X4_TO_FPX4_RS_SF_VEC<"e4m3", v4i8>;
defm : CVT_F32X4_TO_FPX4_RS_SF_VEC<"e5m2", v4i8>;
@@ -4857,12 +4857,16 @@ class WMMA_REGINFO<WMMA_REGS r, string op, string metadata = "", string kind = "
// all fragments of the instruction are viable.
list<Predicate> Predicates = !cond(
!or(!eq(op, "mma.block_scale"),
- !eq(op, "mma.sp.block_scale")) : [hasSM120a, hasPTX<88>],
+ !eq(op, "mma.sp.block_scale")) : [hasSM<120>, hasFamilySpecificFeatures],
+
+ !or(!eq(op, "mma.sp.block_scale"),
+ !eq(kind, "mxf4nvf4"),
+ !eq(kind, "mxf4")) : [callSubtarget<"hasSparseMmaWithBlockScaleF4">],
!or(!eq(ptx_elt_type, "e3m2"),
!eq(ptx_elt_type, "e2m3"),
!eq(ptx_elt_type, "e2m1"),
- !ne(kind, "")) : [hasSM120a, hasPTX<87>],
+ !ne(kind, "")) : [hasSM<120>, hasFamilySpecificFeatures],
!and(!or(!eq(ptx_elt_type,"e4m3"),
!eq(ptx_elt_type,"e5m2")),
@@ -4949,30 +4953,30 @@ class WMMA_REGINFO<WMMA_REGS r, string op, string metadata = "", string kind = "
!and(!eq(op, "ldmatrix"),
!eq(ptx_elt_type, "b8"),
- !eq(geom, "m16n16")) : [hasSM<100>, hasArchAccelFeatures, hasPTX<86>],
+ !eq(geom, "m16n16")) : [callSubtarget<"hasLdStmatrixBlackwellSupport">],
!and(!eq(op, "ldmatrix"),
!eq(ptx_elt_type, "b8x16.b6x16_p32"),
- !eq(geom, "m16n16")) : [hasSM<100>, hasArchAccelFeatures, hasPTX<86>],
+ !eq(geom, "m16n16")) : [callSubtarget<"hasLdStmatrixBlackwellSupport">],
!and(!eq(op, "ldmatrix"),
!eq(ptx_elt_type, "b8x16.b4x16_p64"),
- !eq(geom, "m16n16")) : [hasSM<100>, hasArchAccelFeatures, hasPTX<86>],
+ !eq(geom, "m16n16")) : [callSubtarget<"hasLdStmatrixBlackwellSupport">],
!and(!eq(op, "ldmatrix"),
!eq(ptx_elt_type, "b8x16.b6x16_p32"),
- !eq(geom, "m8n16")) : [hasSM<100>, hasArchAccelFeatures, hasPTX<86>],
+ !eq(geom, "m8n16")) : [callSubtarget<"hasLdStmatrixBlackwellSupport">],
!and(!eq(op, "ldmatrix"),
!eq(ptx_elt_type, "b8x16.b4x16_p64"),
- !eq(geom, "m8n16")) : [hasSM<100>, hasArchAccelFeatures, hasPTX<86>],
+ !eq(geom, "m8n16")) : [callSubtarget<"hasLdStmatrixBlackwellSupport">],
!and(!eq(op, "stmatrix"),!eq(ptx_elt_type, "b16"),
!eq(geom, "m8n8")) : [hasSM<90>, hasPTX<78>],
!and(!eq(op, "stmatrix"),
!eq(ptx_elt_type, "b8"),
- !eq(geom, "m16n8")) : [hasSM<100>, hasArchAccelFeatures, hasPTX<86>]);
+ !eq(geom, "m16n8")) : [callSubtarget<"hasLdStmatrixBlackwellSupport">]);
// template DAGs for instruction inputs/output.
dag Outs = !dag(outs, ptx_regs, reg_names);
@@ -5537,7 +5541,7 @@ multiclass SET_MAXNREG<string Action, Intrinsic Intr> {
def : BasicNVPTXInst<(outs), (ins i32imm:$reg_count),
"setmaxnreg." # Action # ".sync.aligned.u32",
[(Intr timm:$reg_count)]>,
- Requires<[hasArchAccelFeatures, hasSM<90>, hasPTX<80>]>;
+ Requires<[callSubtarget<"hasSetMaxNRegSupport">]>;
}
defm INT_SET_MAXNREG_INC : SET_MAXNREG<"inc", int_nvvm_setmaxnreg_inc_sync_aligned_u32>;
@@ -5742,7 +5746,7 @@ class TCGEN05_ST_INST<string Shape, int Num, bit Unpack> :
# ";";
}
-let isConvergent = true in {
+let isConvergent = true, Predicates = [callSubtarget<"hasTcgen05InstSupport">] in {
foreach shape = ["16x64b", "16x128b", "16x256b", "32x32b", "16x32bx2"] in {
foreach num = !range(0, 8) in {
@@ -5789,7 +5793,7 @@ def CLUSTERLAUNCHCONTRL_TRY_CANCEL_MULTICAST:
"clusterlaunchcontrol.try_cancel.async.shared::cta.mbarrier::complete_tx::bytes" #
".multicast::cluster::all.b128",
[(int_nvvm_clusterlaunchcontrol_try_cancel_async_multicast_shared addr:$addr, addr:$mbar)]>,
- Requires<[hasSM<100>, hasArchAccelFeatures, hasPTX<86>]>;
+ Requires<[hasSM<100>, hasFamilySpecificFeatures, hasPTX<86>]>;
def SDTClusterLaunchControlQueryCancelIsCanceled: SDTypeProfile<1, 2, []>;
def clusterlaunchcontrol_query_cancel_is_canceled:
@@ -5834,10 +5838,14 @@ foreach dim = ["x", "y", "z"] in {
class Tcgen05MMAInst<bit Sp, string KindStr, string ASpace,
int CtaGroup, string CollectorUsage,
- bit ScaleInputD, bit AShift,
- list<Predicate> PTXPredicates> :
- NVPTXInst<(outs), (ins), "?", []>,
- Requires<PTXPredicates> {
+ bit ScaleInputD, bit AShift> :
+ NVPTXInst<(outs), (ins), "?", []>, Requires<[]> {
+
+ let Predicates = !cond(
+ !eq(ScaleInputD, 1) : [callSubtarget<"hasTcgen05MMAScaleInputDImm">],
+ !eq(KindStr, "i8") : [callSubtarget<"hasTcgen05MMAI8Kind">],
+ true : [callSubtarget<"hasTcgen05InstSupport">]
+ );
Intrinsic Intrin = !cast<Intrinsic>(
NVVM_TCGEN05_MMA<Sp, ASpace, AShift, ScaleInputD>.record_name
@@ -5914,10 +5922,7 @@ foreach sp = [0, 1] in {
foreach ashift = !if(!eq(space, "tensor"), [0, 1], [0]) in {
def : Tcgen05MMAInst<sp, kind, space, cta_group, collector_usage,
- scale_input_d, ashift,
- !if(!eq(scale_input_d, 1),
- [hasTcgen05MMAScaleInputDImm],
- [hasTcgen05Instructions])>;
+ scale_input_d, ashift>;
}
}
}
@@ -5957,10 +5962,15 @@ class Tcgen05MMADisableOutputLaneSDNode<bit Sp, string ASpace,
class Tcgen05MMADisableOutputLaneInst<bit Sp, string ASpace,
string Kind, int CtaGroup, string CollectorUsageStr,
- bit ScaleInputD, bit AShift,
- list<Predicate> PTXPredicates> :
+ bit ScaleInputD, bit AShift> :
NVPTXInst<(outs), (ins), "?", []>,
- Requires<PTXPredicates> {
+ Requires<[]> {
+
+ let Predicates = !cond(
+ !eq(ScaleInputD, 1) : [callSubtarget<"hasTcgen05MMAScaleInputDImm">],
+ !eq(Kind, "i8") : [callSubtarget<"hasTcgen05MMAI8Kind">],
+ true : [callSubtarget<"hasTcgen05InstSupport">]
+ );
SDNode Opcode = Tcgen05MMADisableOutputLaneSDNode<Sp, ASpace, CtaGroup,
ScaleInputD, AShift>;
@@ -6063,10 +6073,7 @@ foreach sp = [0, 1] in {
def :
Tcgen05MMADisableOutputLaneInst<sp, space, kind, cta_group,
collector_usage, scale_input_d,
- ashift,
- !if(!eq(scale_input_d, 1),
- [hasTcgen05MMAScaleInputDImm],
- [hasTcgen05Instructions])>;
+ ashift>;
}
}
}
@@ -6076,10 +6083,17 @@ foreach sp = [0, 1] in {
}
class Tcgen05MMABlockScaleInst<bit Sp, string ASpace, string KindStr,
- int CtaGroup, string ScaleVecSize, string CollectorUsageStr,
- Predicate PTXPredicate>:
- NVPTXInst<(outs), (ins), "?", []>,
- Requires<[hasTcgen05Instructions, PTXPredicate]> {
+ int CtaGroup, string ScaleVecSize, string CollectorUsageStr>:
+ NVPTXInst<(outs), (ins), "?", []>, Requires<[]> {
+
+ let Predicates = !cond(
+ !and(!eq(Sp, 1),
+ !eq(KindStr, "mxf4")) : [callSubtarget<"hasTcgen05MMASparseMxf4">],
+ !and(!eq(Sp, 1),
+ !eq(KindStr, "mxf4nvf4")) : [callSubtarget<"hasTcgen05MMASparseMxf4nvf4">],
+ !ne(ScaleVecSize, "") : [callSubtarget<"hasTcgen05InstSupport">, hasPTX<88>],
+ true : [callSubtarget<"hasTcgen05InstSupport">]
+ );
Intrinsic Intrin = !cast<Intrinsic>(
NVVM_TCGEN05_MMA_BLOCKSCALE<Sp, ASpace, KindStr, ScaleVecSize>.record_name);
@@ -6144,9 +6158,7 @@ foreach sp = [0, 1] in {
foreach collector_usage = ["fill", "use", "lastuse", "discard"] in {
if NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED<kind, scale_vec_size>.ret then {
def : Tcgen05MMABlockScaleInst<sp, space, kind, cta_group, scale_vec_size,
- collector_usage,
- !if(!ne(scale_vec_size, ""),
- hasPTX<88>, hasPTX<86>)>;
+ collector_usage>;
}
}
}
@@ -6162,8 +6174,12 @@ foreach sp = [0, 1] in {
class Tcgen05MMAWSInst<bit Sp, string ASpace, string KindStr,
int CollectorBufferB, string CollectorUsageOpStr,
bit HasZeroColMask> :
- NVPTXInst<(outs), (ins), "?", []>,
- Requires<[hasTcgen05Instructions]> {
+ NVPTXInst<(outs), (ins), "?", []>, Requires<[]> {
+
+ let Predicates = !cond(
+ !eq(KindStr, "i8") : [callSubtarget<"hasTcgen05MMAI8Kind">],
+ true : [callSubtarget<"hasTcgen05InstSupport">]
+ );
Intrinsic Intrin = !cast<Intrinsic>(
NVVM_TCGEN05_MMA_WS<Sp, ASpace, HasZeroColMask>.record_name);
diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
index 6f6057b3689e6..55e42a47bd69b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -119,29 +119,6 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
bool hasDotInstructions() const {
return SmVersion >= 61 && PTXVersion >= 50;
}
- // Tcgen05 instructions in Blackwell family
- bool hasTcgen05Instructions() const {
- bool HasTcgen05 = false;
- unsigned MinPTXVersion = 86;
- switch (FullSmVersion) {
- default:
- break;
- case 1003: // sm_100a
- case 1013: // sm_101a
- HasTcgen05 = true;
- break;
- case 1103: // sm_110a
- HasTcgen05 = true;
- MinPTXVersion = 90;
- break;
- case 1033: // sm_103a
- HasTcgen05 = true;
- MinPTXVersion = 88;
- break;
- }
-
- return HasTcgen05 && PTXVersion >= MinPTXVersion;
- }
// Checks following instructions support:
// - tcgen05.ld/st
@@ -149,6 +126,7 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
// - tcgen05.cp
// - tcgen05.fence/wait
// - tcgen05.commit
+ // - tcgen05.mma
bool hasTcgen05InstSupport() const {
// sm_101 renamed to sm_110 in PTX 9.0
return hasPTXWithFamilySMs(90, {100, 110}) ||
@@ -165,8 +143,32 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
}
bool hasTcgen05MMAScaleInputDImm() const {
- return FullSmVersion == 1003 && PTXVersion >= 86;
+ return hasPTXWithFamilySMs(88, {100}) || hasPTXWithAccelSMs(86, {100});
+ }
+
+ bool hasTcgen05MMAI8Kind() const {
+ return hasPTXWithAccelSMs(86, {100, 101}) ||
+ hasPTXWithAccelSMs(90, {100, 110});
+ }
+
+ bool hasTcgen05MMASparseMxf4nvf4() const {
+ return hasPTXWithAccelSMs(87, {100, 101, 103}) ||
+ hasPTXWithAccelSMs(90, {100, 110, 103});
+ }
+
+ bool hasTcgen05MMASparseMxf4() const {
+ return hasPTXWithAccelSMs(86, {100, 101, 103}) ||
+ hasPTXWithAccelSMs(90, {100, 110, 103});
+ }
+
+ bool hasReduxSyncF32() const {
+ return hasPTXWithFamilySMs(88, {100}) || hasPTXWithAccelSMs(86, {100});
+ }
+
+ bool hasSparseMmaWithBlockScaleF4() const {
+ return hasPTXWithAccelSMs(87, {120, 121});
}
+
// f32x2 instructions in Blackwell family
bool hasF32x2Instructions() const;
@@ -202,6 +204,19 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
hasPTXWithAccelSMs(86, {100, 101, 120});
}
+ bool hasSetMaxNRegSupport() const {
+ return hasPTXWithAccelSMs(80, {90}) ||
+ hasPTXWithFamilySMs(90, {100, 110, 120}) ||
+ hasPTXWithFamilySMs(88, {100, 101, 120}) ||
+ hasPTXWithAccelSMs(86, {100, 101, 120});
+ }
+
+ bool hasLdStmatrixBlackwellSupport() const {
+ return hasPTXWithFamilySMs(90, {100, 110, 120}) ||
+ hasPTXWithFamilySMs(88, {100, 101, 120}) ||
+ hasPTXWithAccelSMs(86, {100, 101, 120});
+ }
+
// Prior to CUDA 12.3 ptxas did not recognize that the trap instruction
// terminates a basic block. Instead, it would assume that control flow
// continued to the next instruction. The next instruction could be in the
@@ -211,6 +226,9 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
// present.
bool hasPTXASUnreachableBug() const { return PTXVersion < 83; }
bool hasCvtaParam() const { return SmVersion >= 70 && PTXVersion >= 77; }
+ bool hasCVTRs() const {
+ return hasPTXWithAccelSMs(87, {100, 103});
+ }
unsigned int getFullSmVersion() const { return FullSmVersion; }
unsigned int getSmVersion() const { return getFullSmVersion() / 10; }
unsigned int getSmFamilyVersion() const { return getFullSmVersion() / 100; }
diff --git a/llvm/test/CodeGen/NVPTX/clusterlaunchcontrol-multicast.ll b/llvm/test/CodeGen/NVPTX/clusterlaunchcontrol-multicast.ll
index d930d1842a1d4..9e6beda9b64aa 100644
--- a/llvm/test/CodeGen/NVPTX/clusterlaunchcontrol-multicast.ll
+++ b/llvm/test/CodeGen/NVPTX/clusterlaunchcontrol-multicast.ll
@@ -3,15 +3,32 @@
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_100a %}
+
+; RUN: llc -o - -mcpu=sm_100f -march=nvptx64 -mattr=+ptx88 %s | FileCheck %s --check-prefixes=CHECK,CHECK-PTX-SHARED64
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx88 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s
+; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | %ptxas-verify -arch=sm_100f %}
+; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx88 --nvptx-short-ptr | %ptxas-verify -arch=sm_100f %}
+
; RUN: llc -o - -mcpu=sm_101a -march=nvptx64 -mattr=+ptx86 %s | FileCheck %s --check-prefixes=CHECK,CHECK-PTX-SHARED64
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_101a -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s
; RUN: %if ptxas-sm_101a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_101a -mattr=+ptx86 | %ptxas-verify -arch=sm_101a %}
; RUN: %if ptxas-sm_101a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_101a -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_101a %}
+
+; RUN: llc -o - -mcpu=sm_110f -march=nvptx64 -mattr=+ptx90 %s | FileCheck %s --check-prefixes=CHECK,CHECK-PTX-SHARED64
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx90 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s
+; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}
+; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx90 --nvptx-short-ptr | %ptxas-verify -arch=sm_110f %}
+
; RUN: llc -o - -mcpu=sm_120a -march=nvptx64 -mattr=+ptx86 %s | FileCheck %s --check-prefixes=CHECK,CHECK-PTX-SHARED64
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_120a -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s
; RUN: %if ptxas-sm_120a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120a -mattr=+ptx86 | %ptxas-verify -arch=sm_120a %}
; RUN: %if ptxas-sm_120a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120a -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_120a %}
+; RUN: llc -o - -mcpu=sm_120f -march=nvptx64 -mattr=+ptx88 %s | FileCheck %s --check-prefixes=CHECK,CHECK-PTX-SHARED64
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx88 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK,CHECK-PTX-SHARED32 %s
+; RUN: %if ptxas-sm_120f && ptxas-isa-8.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx88 | %ptxas-verify -arch=sm_120f %}
+; RUN: %if ptxas-sm_120f && ptxas-isa-8.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx88 --nvptx-short-ptr | %ptxas-verify -arch=sm_120f %}
+
define void @nvvm_clusterlaunchcontrol_try_cancel_multicast(
; CHECK-PTX-SHARED64-LABEL: nvvm_clusterlaunchcontrol_try_cancel_multicast(
; CHECK-PTX-SHARED64: {
diff --git a/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll b/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll
index 38c9234c78feb..6522ac5ff4543 100644
--- a/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll
+++ b/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll
@@ -1,6 +1,8 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | FileCheck %s
; RUN: %if ptxas-sm_100...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/173268
More information about the llvm-commits
mailing list