Skip to content

Commit 5660a61

Browse files
rajatbajpairajatbajpai
committed
[NVPTX] Update architecture support checks for tcgen05
This change updates architecture support checks for tcgen05 intrinsics (except tcgen05.mma.*). The newer checks will support family-specific architecture variants as well. After this change, the arch checks will be accurate and match with PTX ISA. Intrinsics affected: - tcgen05.ld/st - tcgen05.alloc/dealloc/relinquish - tcgen05.cp - tcgen05.fence/wait - tcgen05.commit - tcgen05.shift
1 parent a33544b commit 5660a61

File tree

12 files changed

+140
-16
lines changed

12 files changed

+140
-16
lines changed

‎llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -280,6 +280,10 @@ static unsigned getTcgen05LdOpcode(unsigned IID, bool enablePack) {
280280
}
281281

282282
void NVPTXDAGToDAGISel::SelectTcgen05Ld(SDNode *N, bool hasOffset) {
283+
if (!Subtarget->hasTcgen05InstSupport())
284+
report_fatal_error(
285+
"tcgen05.ld is not supported on this architecture variant");
286+
283287
SDLoc DL(N);
284288
unsigned IID = cast<ConstantSDNode>(N->getOperand(1))->getZExtValue();
285289

@@ -2136,6 +2140,10 @@ static unsigned getTcgen05StOpcode(unsigned IID, bool enableUnpack) {
21362140
}
21372141

21382142
void NVPTXDAGToDAGISel::SelectTcgen05St(SDNode *N, bool hasOffset) {
2143+
if (!Subtarget->hasTcgen05InstSupport())
2144+
report_fatal_error(
2145+
"tcgen05.st is not supported on this architecture variant");
2146+
21392147
SDLoc DL(N);
21402148
unsigned IID = cast<ConstantSDNode>(N->getOperand(1))->getZExtValue();
21412149

‎llvm/lib/Target/NVPTX/NVPTXInstrInfo.td‎

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -99,6 +99,45 @@ def PrmtMode : Operand<i32> {
9999
// NVPTX Instruction Predicate Definitions
100100
//===----------------------------------------------------------------------===//
101101

102+
// Helper predicate to compose multiple predicates.
103+
class AnyPred<list<Predicate> Preds>
104+
: Predicate<"(" #
105+
!interleave(!foreach(pred, Preds, pred.CondString),
106+
") || (") #
107+
")">;
108+
109+
// Checks PTX version and family-specific and architecture-specific SM versions.
110+
// For example, sm_100{f/a} and any future variants in the same family will match.
111+
class PTXWithFamilySMs<int PTXVersion, list<int> SMVersions> :
112+
Predicate<"Subtarget->getPTXVersion() >= " # PTXVersion #
113+
" && Subtarget->hasFamilySpecificFeatures()" #
114+
" && (" #
115+
!interleave(!foreach(sm, SMVersions,
116+
"(Subtarget->getSmFamilyVersion() == " # !div(sm, 10) #
117+
" && Subtarget->getSmVersion() >= " # sm # ")"),
118+
" || ") #
119+
")">;
120+
121+
// Checks PTX version and architecture-specific SM versions.
122+
// For example, sm_100{a} will match.
123+
class PTXWithAccelSMs<int PTXVersion, list<int> SMVersions> :
124+
Predicate<"Subtarget->getPTXVersion() >= " # PTXVersion #
125+
" && Subtarget->hasArchAccelFeatures()" #
126+
" && (" #
127+
!interleave(!foreach(sm, SMVersions,
128+
"Subtarget->getSmVersion() == " # sm),
129+
" || ") #
130+
")">;
131+
132+
// Helper predicate to call a subtarget method.
133+
class callSubtarget<string SubtargetMethod> : Predicate<"Subtarget->" # SubtargetMethod # "()">;
134+
135+
// Composed predicate to check tcgen05.shift instructions support.
136+
def hasTcgen05ShiftSupport : AnyPred<[
137+
PTXWithAccelSMs<90, [100, 110, 103]>,
138+
PTXWithAccelSMs<88, [100, 101, 103]>,
139+
PTXWithAccelSMs<86, [100, 101]>
140+
]>;
102141

103142
def hasAtomAddF64 : Predicate<"Subtarget->hasAtomAddF64()">;
104143
def hasAtomScope : Predicate<"Subtarget->hasAtomScope()">;

‎llvm/lib/Target/NVPTX/NVPTXIntrinsics.td‎

Lines changed: 18 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -5024,8 +5024,8 @@ let Predicates = [hasSM<90>, hasPTX<78>] in {
50245024
def EXIT : NullaryInst<"exit", int_nvvm_exit>;
50255025

50265026
// Tcgen05 intrinsics
5027-
let isConvergent = true, Predicates = [hasTcgen05Instructions] in {
5028-
5027+
let isConvergent = true in {
5028+
let Predicates = [callSubtarget<"hasTcgen05InstSupport">] in {
50295029
multiclass TCGEN05_ALLOC_INTR<string AS, string num, Intrinsic Intr> {
50305030
def "" : BasicNVPTXInst<(outs),
50315031
(ins ADDR:$dst, B32:$ncols),
@@ -5077,15 +5077,6 @@ defm TCGEN05_COMMIT_CG2 : TCGEN05_COMMIT_INTR<"", "2">;
50775077
defm TCGEN05_COMMIT_S64_CG1 : TCGEN05_COMMIT_INTR<"shared", "1">;
50785078
defm TCGEN05_COMMIT_S64_CG2 : TCGEN05_COMMIT_INTR<"shared", "2">;
50795079

5080-
multiclass TCGEN05_SHIFT_INTR<string num, Intrinsic Intr> {
5081-
def "" : BasicNVPTXInst<(outs),
5082-
(ins ADDR:$tmem_addr),
5083-
"tcgen05.shift.cta_group::" # num # ".down",
5084-
[(Intr addr:$tmem_addr)]>;
5085-
}
5086-
defm TCGEN05_SHIFT_CG1: TCGEN05_SHIFT_INTR<"1", int_nvvm_tcgen05_shift_down_cg1>;
5087-
defm TCGEN05_SHIFT_CG2: TCGEN05_SHIFT_INTR<"2", int_nvvm_tcgen05_shift_down_cg2>;
5088-
50895080
multiclass TCGEN05_CP_INTR<string shape, string src_fmt, string mc = ""> {
50905081
defvar dst_fmt = !if(!eq(src_fmt, ""), "", ".b8x16");
50915082
defvar fmt_asm = StrJoin<".", [dst_fmt, src_fmt]>.ret;
@@ -5116,9 +5107,22 @@ foreach src_fmt = ["", "b6x16_p32", "b4x16_p64"] in {
51165107
defm TCGEN05_CP_64x128_2 # src_fmt : TCGEN05_CP_INTR<"64x128b", src_fmt, "warpx2::01_23">;
51175108
defm TCGEN05_CP_32x128 # src_fmt : TCGEN05_CP_INTR<"32x128b", src_fmt, "warpx4">;
51185109
}
5110+
} // Predicates
5111+
5112+
let Predicates = [hasTcgen05ShiftSupport] in {
5113+
multiclass TCGEN05_SHIFT_INTR<string num, Intrinsic Intr> {
5114+
def "" : BasicNVPTXInst<(outs),
5115+
(ins ADDR:$tmem_addr),
5116+
"tcgen05.shift.cta_group::" # num # ".down",
5117+
[(Intr addr:$tmem_addr)]>;
5118+
}
5119+
defm TCGEN05_SHIFT_CG1: TCGEN05_SHIFT_INTR<"1", int_nvvm_tcgen05_shift_down_cg1>;
5120+
defm TCGEN05_SHIFT_CG2: TCGEN05_SHIFT_INTR<"2", int_nvvm_tcgen05_shift_down_cg2>;
5121+
} // Predicates
5122+
51195123
} // isConvergent
51205124

5121-
let hasSideEffects = 1, Predicates = [hasTcgen05Instructions] in {
5125+
let hasSideEffects = 1, Predicates = [callSubtarget<"hasTcgen05InstSupport">] in {
51225126

51235127
def tcgen05_fence_before_thread_sync: NullaryInst<
51245128
"tcgen05.fence::before_thread_sync", int_nvvm_tcgen05_fence_before_thread_sync>;
@@ -5152,8 +5156,7 @@ class TCGEN05_LDST_REGINFO<int Veclen> {
51525156
//
51535157

51545158
class TCGEN05_LD_INST<string Shape, int Num, bit Pack> :
5155-
NVPTXInst<(outs), (ins), "?", []>,
5156-
Requires<[hasTcgen05Instructions]> {
5159+
NVPTXInst<(outs), (ins), "?", []> {
51575160

51585161
TCGEN05_LDST_REGINFO Info = TCGEN05_LDST_REGINFO<
51595162
NVVM_TCGEN05_LDST_ACCESS_SIZE<Shape, Num>.veclen>;
@@ -5177,8 +5180,7 @@ class TCGEN05_LD_INST<string Shape, int Num, bit Pack> :
51775180
//
51785181

51795182
class TCGEN05_ST_INST<string Shape, int Num, bit Unpack> :
5180-
NVPTXInst<(outs), (ins), "?", []>,
5181-
Requires<[hasTcgen05Instructions]> {
5183+
NVPTXInst<(outs), (ins), "?", []> {
51825184

51835185
TCGEN05_LDST_REGINFO Info = TCGEN05_LDST_REGINFO<
51845186
NVVM_TCGEN05_LDST_ACCESS_SIZE<Shape, Num>.veclen>;

‎llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp‎

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,32 @@ const SelectionDAGTargetInfo *NVPTXSubtarget::getSelectionDAGInfo() const {
7272
return TSInfo.get();
7373
}
7474

75+
bool NVPTXSubtarget::hasPTXWithFamilySMs(unsigned PTXVersion,
76+
ArrayRef<unsigned> SMVersions) const {
77+
if (!hasFamilySpecificFeatures() || getPTXVersion() < PTXVersion)
78+
return false;
79+
80+
for (unsigned SM : SMVersions) {
81+
if (getSmFamilyVersion() == SM / 10 && getSmVersion() >= SM)
82+
return true;
83+
}
84+
85+
return false;
86+
}
87+
88+
bool NVPTXSubtarget::hasPTXWithAccelSMs(unsigned PTXVersion,
89+
ArrayRef<unsigned> SMVersions) const {
90+
if (!hasArchAccelFeatures() || getPTXVersion() < PTXVersion)
91+
return false;
92+
93+
for (unsigned SM : SMVersions) {
94+
if (getSmVersion() == SM)
95+
return true;
96+
}
97+
98+
return false;
99+
}
100+
75101
bool NVPTXSubtarget::allowFP16Math() const {
76102
return hasFP16Math() && NoF16Math == false;
77103
}

‎llvm/lib/Target/NVPTX/NVPTXSubtarget.h‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -73,6 +73,16 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
7373

7474
const SelectionDAGTargetInfo *getSelectionDAGInfo() const override;
7575

76+
// Checks PTX version and family-specific and architecture-specific SM
77+
// versions. For example, sm_100{f/a} and any future variants in the same
78+
// family will match.
79+
bool hasPTXWithFamilySMs(unsigned PTXVersion,
80+
ArrayRef<unsigned> SMVersions) const;
81+
// Checks PTX version and architecture-specific SM versions.
82+
// For example, sm_100{a} will match.
83+
bool hasPTXWithAccelSMs(unsigned PTXVersion,
84+
ArrayRef<unsigned> SMVersions) const;
85+
7686
bool has256BitVectorLoadStore(unsigned AS) const {
7787
return SmVersion >= 100 && PTXVersion >= 88 &&
7888
AS == NVPTXAS::ADDRESS_SPACE_GLOBAL;
@@ -127,6 +137,18 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
127137
return HasTcgen05 && PTXVersion >= MinPTXVersion;
128138
}
129139

140+
// Checks following instructions support:
141+
// - tcgen05.ld/st
142+
// - tcgen05.alloc/dealloc/relinquish
143+
// - tcgen05.cp
144+
// - tcgen05.fence/wait
145+
// - tcgen05.commit
146+
bool hasTcgen05InstSupport() const {
147+
return hasPTXWithFamilySMs(90, {100, 110}) ||
148+
hasPTXWithFamilySMs(88, {100, 101}) ||
149+
hasPTXWithAccelSMs(86, {100, 101});
150+
}
151+
130152
bool hasTcgen05MMAScaleInputDImm() const {
131153
return FullSmVersion == 1003 && PTXVersion >= 86;
132154
}
@@ -158,6 +180,7 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
158180
bool hasCvtaParam() const { return SmVersion >= 70 && PTXVersion >= 77; }
159181
unsigned int getFullSmVersion() const { return FullSmVersion; }
160182
unsigned int getSmVersion() const { return getFullSmVersion() / 10; }
183+
unsigned int getSmFamilyVersion() const { return getFullSmVersion() / 100; }
161184
// GPUs with "a" suffix have architecture-accelerated features that are
162185
// supported on the specified architecture only, hence such targets do not
163186
// follow the onion layer model. hasArchAccelFeatures() allows distinguishing

‎llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2,9 +2,13 @@
22
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck --check-prefixes=CHECK_PTX64 %s
33
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK_PTX64_SHARED32 %s
44
; RUN: llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck --check-prefixes=CHECK_PTX64 %s
5+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | FileCheck --check-prefixes=CHECK_PTX64 %s
6+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | FileCheck --check-prefixes=CHECK_PTX64 %s
57
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}
68
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_100a %}
79
; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}
10+
; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | %ptxas-verify -arch=sm_100f %}
11+
; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}
812

913

1014
declare void @llvm.nvvm.tcgen05.alloc.cg1(ptr %addr, i32 %ncols)

‎llvm/test/CodeGen/NVPTX/tcgen05-commit.ll‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2,9 +2,13 @@
22
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck --check-prefixes=CHECK_PTX64 %s
33
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK_PTX64_SHARED32 %s
44
; RUN: llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck --check-prefixes=CHECK_PTX64 %s
5+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | FileCheck --check-prefixes=CHECK_PTX64 %s
6+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | FileCheck --check-prefixes=CHECK_PTX64 %s
57
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}
68
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_100a %}
79
; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}
10+
; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | %ptxas-verify -arch=sm_100f %}
11+
; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}
812

913
declare void @llvm.nvvm.tcgen05.commit.cg1(ptr %bar_addr)
1014
declare void @llvm.nvvm.tcgen05.commit.cg2(ptr %bar_addr)

‎llvm/test/CodeGen/NVPTX/tcgen05-cp.ll‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,12 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
22
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck --check-prefixes=CHECK %s
33
; RUN: llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck --check-prefixes=CHECK %s
4+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | FileCheck --check-prefixes=CHECK %s
5+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | FileCheck --check-prefixes=CHECK %s
46
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}
57
; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}
8+
; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | %ptxas-verify -arch=sm_100f %}
9+
; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}
610

711
; CHECK-LABEL: test_tcgen05_cp_64x128_v1
812
define void @test_tcgen05_cp_64x128_v1(ptr addrspace(6) %addr, i64 %sdesc) {

‎llvm/test/CodeGen/NVPTX/tcgen05-fence.ll‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,12 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
22
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck --check-prefixes=CHECK %s
33
; RUN: llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck --check-prefixes=CHECK %s
4+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | FileCheck --check-prefixes=CHECK %s
5+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | FileCheck --check-prefixes=CHECK %s
46
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}
57
; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}
8+
; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | %ptxas-verify -arch=sm_100f %}
9+
; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}
610

711
declare void @llvm.nvvm.tcgen05.fence.before.thread.sync()
812
declare void @llvm.nvvm.tcgen05.fence.after.thread.sync()

‎llvm/test/CodeGen/NVPTX/tcgen05-ld.ll‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2,9 +2,13 @@
22
; RUN: llc < %s -o - -mcpu=sm_100a -march=nvptx64 -mattr=+ptx86 | FileCheck %s
33
; RUN: llc < %s -o - -mcpu=sm_101a -march=nvptx64 -mattr=+ptx86 | FileCheck %s
44
; RUN: llc < %s -o - -mcpu=sm_103a -march=nvptx64 -mattr=+ptx88 | FileCheck %s
5+
; RUN: llc < %s -o - -mcpu=sm_100f -march=nvptx64 -mattr=+ptx88 | FileCheck %s
6+
; RUN: llc < %s -o - -mcpu=sm_110f -march=nvptx64 -mattr=+ptx90 | FileCheck %s
57
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mattr=+ptx86 -mcpu=sm_100a | %ptxas-verify -arch=sm_100a %}
68
; RUN: %if ptxas-sm_101a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mattr=+ptx86 -mcpu=sm_101a | %ptxas-verify -arch=sm_101a %}
79
; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mattr=+ptx88 -mcpu=sm_103a | %ptxas-verify -arch=sm_103a %}
10+
; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mattr=+ptx88 -mcpu=sm_100f | %ptxas-verify -arch=sm_100f %}
11+
; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mattr=+ptx90 -mcpu=sm_110f | %ptxas-verify -arch=sm_110f %}
812

913
; CHECK-LABEL: nvvm_tcgen05_ld_16x64b
1014
define void @nvvm_tcgen05_ld_16x64b(ptr addrspace(6) %taddr) {

0 commit comments

Comments
 (0)