Skip to content

Commit 9138b8e

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 a406eb4 commit 9138b8e

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
@@ -101,6 +101,45 @@ def PrmtMode : Operand<i32> {
101101
// NVPTX Instruction Predicate Definitions
102102
//===----------------------------------------------------------------------===//
103103

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

105144
def hasAtomAddF64 : Predicate<"Subtarget->hasAtomAddF64()">;
106145
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
@@ -5091,8 +5091,8 @@ let Predicates = [hasSM<90>, hasPTX<78>] in {
50915091
def EXIT : NullaryInst<"exit", int_nvvm_exit>;
50925092

50935093
// Tcgen05 intrinsics
5094-
let isConvergent = true, Predicates = [hasTcgen05Instructions] in {
5095-
5094+
let isConvergent = true in {
5095+
let Predicates = [callSubtarget<"hasTcgen05InstSupport">] in {
50965096
multiclass TCGEN05_ALLOC_INTR<string AS, string num, Intrinsic Intr> {
50975097
def "" : BasicNVPTXInst<(outs),
50985098
(ins ADDR:$dst, B32:$ncols),
@@ -5144,15 +5144,6 @@ defm TCGEN05_COMMIT_CG2 : TCGEN05_COMMIT_INTR<"", "2">;
51445144
defm TCGEN05_COMMIT_S64_CG1 : TCGEN05_COMMIT_INTR<"shared", "1">;
51455145
defm TCGEN05_COMMIT_S64_CG2 : TCGEN05_COMMIT_INTR<"shared", "2">;
51465146

5147-
multiclass TCGEN05_SHIFT_INTR<string num, Intrinsic Intr> {
5148-
def "" : BasicNVPTXInst<(outs),
5149-
(ins ADDR:$tmem_addr),
5150-
"tcgen05.shift.cta_group::" # num # ".down",
5151-
[(Intr addr:$tmem_addr)]>;
5152-
}
5153-
defm TCGEN05_SHIFT_CG1: TCGEN05_SHIFT_INTR<"1", int_nvvm_tcgen05_shift_down_cg1>;
5154-
defm TCGEN05_SHIFT_CG2: TCGEN05_SHIFT_INTR<"2", int_nvvm_tcgen05_shift_down_cg2>;
5155-
51565147
multiclass TCGEN05_CP_INTR<string shape, string src_fmt, string mc = ""> {
51575148
defvar dst_fmt = !if(!eq(src_fmt, ""), "", ".b8x16");
51585149
defvar fmt_asm = StrJoin<".", [dst_fmt, src_fmt]>.ret;
@@ -5183,9 +5174,22 @@ foreach src_fmt = ["", "b6x16_p32", "b4x16_p64"] in {
51835174
defm TCGEN05_CP_64x128_2 # src_fmt : TCGEN05_CP_INTR<"64x128b", src_fmt, "warpx2::01_23">;
51845175
defm TCGEN05_CP_32x128 # src_fmt : TCGEN05_CP_INTR<"32x128b", src_fmt, "warpx4">;
51855176
}
5177+
} // Predicates
5178+
5179+
let Predicates = [hasTcgen05ShiftSupport] in {
5180+
multiclass TCGEN05_SHIFT_INTR<string num, Intrinsic Intr> {
5181+
def "" : BasicNVPTXInst<(outs),
5182+
(ins ADDR:$tmem_addr),
5183+
"tcgen05.shift.cta_group::" # num # ".down",
5184+
[(Intr addr:$tmem_addr)]>;
5185+
}
5186+
defm TCGEN05_SHIFT_CG1: TCGEN05_SHIFT_INTR<"1", int_nvvm_tcgen05_shift_down_cg1>;
5187+
defm TCGEN05_SHIFT_CG2: TCGEN05_SHIFT_INTR<"2", int_nvvm_tcgen05_shift_down_cg2>;
5188+
} // Predicates
5189+
51865190
} // isConvergent
51875191

5188-
let hasSideEffects = 1, Predicates = [hasTcgen05Instructions] in {
5192+
let hasSideEffects = 1, Predicates = [callSubtarget<"hasTcgen05InstSupport">] in {
51895193

51905194
def tcgen05_fence_before_thread_sync: NullaryInst<
51915195
"tcgen05.fence::before_thread_sync", int_nvvm_tcgen05_fence_before_thread_sync>;
@@ -5219,8 +5223,7 @@ class TCGEN05_LDST_REGINFO<int Veclen> {
52195223
//
52205224

52215225
class TCGEN05_LD_INST<string Shape, int Num, bit Pack> :
5222-
NVPTXInst<(outs), (ins), "?", []>,
5223-
Requires<[hasTcgen05Instructions]> {
5226+
NVPTXInst<(outs), (ins), "?", []> {
52245227

52255228
TCGEN05_LDST_REGINFO Info = TCGEN05_LDST_REGINFO<
52265229
NVVM_TCGEN05_LDST_ACCESS_SIZE<Shape, Num>.veclen>;
@@ -5244,8 +5247,7 @@ class TCGEN05_LD_INST<string Shape, int Num, bit Pack> :
52445247
//
52455248

52465249
class TCGEN05_ST_INST<string Shape, int Num, bit Unpack> :
5247-
NVPTXInst<(outs), (ins), "?", []>,
5248-
Requires<[hasTcgen05Instructions]> {
5250+
NVPTXInst<(outs), (ins), "?", []> {
52495251

52505252
TCGEN05_LDST_REGINFO Info = TCGEN05_LDST_REGINFO<
52515253
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)