Skip to content

Commit 56b6209

Browse files
modikingmodiking
committed
review feedback
1 parent f7a17c4 commit 56b6209

File tree

7 files changed

+171
-85
lines changed

7 files changed

+171
-85
lines changed

‎llvm/docs/NVPTXUsage.rst

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -108,6 +108,7 @@ The NVPTX back-end uses the following address space mapping:
108108
3 Shared
109109
4 Constant
110110
5 Local
111+
7 Shared Cluster
111112
============= ======================
112113

113114
Every global variable and pointer type is assigned to one of these address
@@ -306,6 +307,32 @@ If the given pointer in the generic address space refers to memory which falls
306307
within the state space of the intrinsic (and therefore could be safely address
307308
space casted to this space), 1 is returned, otherwise 0 is returned.
308309

310+
'``llvm.nvvm.mapa.*``' Intrinsics
311+
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
312+
313+
Syntax:
314+
"""""""
315+
316+
.. code-block:: llvm
317+
318+
declare ptr @llvm.nvvm.mapa(ptr %p, i32 %rank)
319+
declare ptr addrspace(7) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %p, i32 %rank)
320+
321+
Overview:
322+
"""""""""
323+
324+
The '``llvm.nvvm.mapa.*``' intrinsics map a shared memory pointer ``p`` of another CTA with ``%rank`` to the current CTA.
325+
The ``llvm.nvvm.mapa`` form expects a generic pointer to shared memory and returns a generic pointer to shared cluster memory.
326+
The ``llvm.nvvm.mapa.shared.cluster`` form expects a pointer to shared memory and returns a pointer to shared cluster memory.
327+
They corresponds directly to the ``mapa`` and ``mapa.shared.cluster`` PTX instructions.
328+
329+
Semantics:
330+
""""""""""
331+
332+
If the given pointer in the generic address space refers to memory which falls
333+
within the state space of the intrinsic (and therefore could be safely address
334+
space casted to this space), 1 is returned, otherwise 0 is returned.
335+
309336
Arithmetic Intrinsics
310337
---------------------
311338

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

Lines changed: 8 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -3046,36 +3046,20 @@ SDValue NVPTXTargetLowering::LowerADDRSPACECAST(SDValue Op,
30463046
DestAS != llvm::ADDRESS_SPACE_GENERIC) {
30473047
// Shared and SharedCluster can be converted to each other through generic
30483048
// space
3049-
if (SrcAS == llvm::ADDRESS_SPACE_SHARED &&
3050-
DestAS == llvm::ADDRESS_SPACE_SHARED_CLUSTER) {
3049+
if ((SrcAS == llvm::ADDRESS_SPACE_SHARED &&
3050+
DestAS == llvm::ADDRESS_SPACE_SHARED_CLUSTER) ||
3051+
(SrcAS == llvm::ADDRESS_SPACE_SHARED_CLUSTER &&
3052+
DestAS == llvm::ADDRESS_SPACE_SHARED)) {
30513053
const MVT GenerictVT =
30523054
getPointerTy(DAG.getDataLayout(), ADDRESS_SPACE_GENERIC);
3053-
const MVT SharedClusterVT =
3054-
getPointerTy(DAG.getDataLayout(), ADDRESS_SPACE_SHARED_CLUSTER);
3055-
SDValue GenericConversion =
3056-
DAG.getAddrSpaceCast(SDLoc(), GenerictVT, Op.getOperand(0),
3057-
ADDRESS_SPACE_SHARED, ADDRESS_SPACE_GENERIC);
3055+
const MVT OutputVT = getPointerTy(DAG.getDataLayout(), DestAS);
3056+
SDValue GenericConversion = DAG.getAddrSpaceCast(
3057+
SDLoc(), GenerictVT, Op.getOperand(0), SrcAS, ADDRESS_SPACE_GENERIC);
30583058
SDValue SharedClusterConversion = DAG.getAddrSpaceCast(
3059-
SDLoc(), SharedClusterVT, GenericConversion, ADDRESS_SPACE_GENERIC,
3060-
ADDRESS_SPACE_SHARED_CLUSTER);
3059+
SDLoc(), OutputVT, GenericConversion, ADDRESS_SPACE_GENERIC, DestAS);
30613060
return SharedClusterConversion;
30623061
}
30633062

3064-
if (SrcAS == llvm::ADDRESS_SPACE_SHARED_CLUSTER &&
3065-
DestAS == llvm::ADDRESS_SPACE_SHARED) {
3066-
const MVT GenerictVT =
3067-
getPointerTy(DAG.getDataLayout(), ADDRESS_SPACE_GENERIC);
3068-
const MVT SharedVT =
3069-
getPointerTy(DAG.getDataLayout(), ADDRESS_SPACE_SHARED);
3070-
SDValue GenericConversion = DAG.getAddrSpaceCast(
3071-
SDLoc(), GenerictVT, Op.getOperand(0), ADDRESS_SPACE_SHARED_CLUSTER,
3072-
ADDRESS_SPACE_GENERIC);
3073-
SDValue SharedConversion =
3074-
DAG.getAddrSpaceCast(SDLoc(), SharedVT, GenericConversion,
3075-
ADDRESS_SPACE_GENERIC, ADDRESS_SPACE_SHARED);
3076-
return SharedConversion;
3077-
}
3078-
30793063
return DAG.getUNDEF(Op.getValueType());
30803064
}
30813065

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

Lines changed: 22 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -2387,28 +2387,28 @@ def INT_PTX_LDG_G_v4i32_ELE : VLDG_G_ELE_V4<"u32", Int32Regs>;
23872387
def INT_PTX_LDG_G_v4f32_ELE : VLDG_G_ELE_V4<"f32", Float32Regs>;
23882388

23892389

2390-
multiclass NG_TO_G<string Str, list<Predicate> Preds = []> {
2391-
def "" : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
2392-
"cvta." # Str # ".u32 \t$result, $src;", []>, Requires<Preds>;
2393-
def _64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
2394-
"cvta." # Str # ".u64 \t$result, $src;", []>, Requires<Preds>;
2395-
}
2396-
2397-
multiclass NG_TO_G_64<string Str, list<Predicate> Preds = []> {
2398-
def _64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
2399-
"cvta." # Str # ".u64 \t$result, $src;", []>, Requires<Preds>;
2400-
}
2401-
2402-
multiclass G_TO_NG<string Str, list<Predicate> Preds = []> {
2403-
def "" : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
2404-
"cvta.to." # Str # ".u32 \t$result, $src;", []>, Requires<Preds>;
2405-
def _64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
2406-
"cvta.to." # Str # ".u64 \t$result, $src;", []>, Requires<Preds>;
2390+
multiclass NG_TO_G<string Str, bit Supports32 = 1, list<Predicate> Preds = []> {
2391+
foreach bitwidth = !if(Supports32, ["32", "64"], ["64"]) in {
2392+
if !eq(bitwidth, "32") then {
2393+
def "" : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
2394+
"cvta." # Str # ".u32 \t$result, $src;", []>, Requires<Preds>;
2395+
} else if !eq(bitwidth, "64") then {
2396+
def _64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
2397+
"cvta." # Str # ".u64 \t$result, $src;", []>, Requires<Preds>;
2398+
}
2399+
}
24072400
}
24082401

2409-
multiclass G_TO_NG_64<string Str, list<Predicate> Preds = []> {
2410-
def _64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
2411-
"cvta.to." # Str # ".u64 \t$result, $src;", []>, Requires<Preds>;
2402+
multiclass G_TO_NG<string Str, bit Supports32 = 1, list<Predicate> Preds = []> {
2403+
foreach bitwidth = !if(Supports32, ["32", "64"], ["64"]) in {
2404+
if !eq(bitwidth, "32") then {
2405+
def "" : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
2406+
"cvta.to." # Str # ".u32 \t$result, $src;", []>, Requires<Preds>;
2407+
} else if !eq(bitwidth, "64") then {
2408+
def _64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
2409+
"cvta.to." # Str # ".u64 \t$result, $src;", []>, Requires<Preds>;
2410+
}
2411+
}
24122412
}
24132413

24142414
foreach space = ["local", "shared", "global", "const", "param"] in {
@@ -2417,8 +2417,8 @@ foreach space = ["local", "shared", "global", "const", "param"] in {
24172417
}
24182418

24192419
foreach space = ["shared::cluster"] in {
2420-
defm cvta_#space : NG_TO_G<space>;
2421-
defm cvta_to_#space : G_TO_NG<space>;
2420+
defm cvta_#space : NG_TO_G<space, false, [hasclusters]>;
2421+
defm cvta_to_#space : G_TO_NG<space, false, [hasclusters]>;
24222422
}
24232423

24242424
def : Pat<(int_nvvm_ptr_param_to_gen i32:$src),

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

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -426,10 +426,7 @@ static std::optional<bool> evaluateIsSpace(Intrinsic::ID IID, unsigned AS) {
426426
case Intrinsic::nvvm_isspacep_shared:
427427
return AS == NVPTXAS::ADDRESS_SPACE_SHARED;
428428
case Intrinsic::nvvm_isspacep_shared_cluster:
429-
// We can't tell shared from shared_cluster at compile time from AS alone,
430-
// but it can't be either is AS is not shared.
431-
return AS == NVPTXAS::ADDRESS_SPACE_SHARED ? std::nullopt
432-
: std::optional{false};
429+
return AS == NVPTXAS::ADDRESS_SPACE_SHARED_CLUSTER;
433430
case Intrinsic::nvvm_isspacep_const:
434431
return AS == NVPTXAS::ADDRESS_SPACE_CONST;
435432
default:

‎llvm/test/CodeGen/NVPTX/addrspacecast-ptx64.ll

Lines changed: 108 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -1,15 +1,36 @@
1-
; RUN: llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_80 | FileCheck %s -check-prefixes=ALL,NOPTRCONV,CLS64
2-
; RUN: llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_80 --nvptx-short-ptr | FileCheck %s -check-prefixes=ALL,PTRCONV,CLS64
3-
; RUN: %if ptxas-12.8 %{ llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_80 | %ptxas-verify %}
4-
; RUN: %if ptxas-12.8 %{ llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_80 --nvptx-short-ptr | %ptxas-verify %}
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | FileCheck %s -check-prefixes=NOPTRCONV
3+
; RUN: llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 --nvptx-short-ptr | FileCheck %s -check-prefixes=PTRCONV
4+
; RUN: %if ptxas-12.8 %{ llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify %}
5+
; RUN: %if ptxas-12.8 %{ llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 --nvptx-short-ptr | %ptxas-verify %}
56

67
; ALL-LABEL: conv_shared_cluster_to_generic
78
define i32 @conv_shared_cluster_to_generic(ptr addrspace(7) %ptr) {
89
; CLS32: cvta.shared::cluster.u32
9-
; PTRCONV: cvt.u64.u32
10-
; NOPTRCONV-NOT: cvt.u64.u32
11-
; CLS64: cvta.shared::cluster.u64
12-
; ALL: ld.u32
10+
; NOPTRCONV-LABEL: conv_shared_cluster_to_generic(
11+
; NOPTRCONV: {
12+
; NOPTRCONV-NEXT: .reg .b32 %r<2>;
13+
; NOPTRCONV-NEXT: .reg .b64 %rd<3>;
14+
; NOPTRCONV-EMPTY:
15+
; NOPTRCONV-NEXT: // %bb.0:
16+
; NOPTRCONV-NEXT: ld.param.u64 %rd1, [conv_shared_cluster_to_generic_param_0];
17+
; NOPTRCONV-NEXT: cvta.shared::cluster.u64 %rd2, %rd1;
18+
; NOPTRCONV-NEXT: ld.u32 %r1, [%rd2];
19+
; NOPTRCONV-NEXT: st.param.b32 [func_retval0], %r1;
20+
; NOPTRCONV-NEXT: ret;
21+
;
22+
; PTRCONV-LABEL: conv_shared_cluster_to_generic(
23+
; PTRCONV: {
24+
; PTRCONV-NEXT: .reg .b32 %r<3>;
25+
; PTRCONV-NEXT: .reg .b64 %rd<3>;
26+
; PTRCONV-EMPTY:
27+
; PTRCONV-NEXT: // %bb.0:
28+
; PTRCONV-NEXT: ld.param.u32 %r1, [conv_shared_cluster_to_generic_param_0];
29+
; PTRCONV-NEXT: cvt.u64.u32 %rd1, %r1;
30+
; PTRCONV-NEXT: cvta.shared::cluster.u64 %rd2, %rd1;
31+
; PTRCONV-NEXT: ld.u32 %r2, [%rd2];
32+
; PTRCONV-NEXT: st.param.b32 [func_retval0], %r2;
33+
; PTRCONV-NEXT: ret;
1334
%genptr = addrspacecast ptr addrspace(7) %ptr to ptr
1435
%val = load i32, ptr %genptr
1536
ret i32 %val
@@ -18,31 +39,99 @@ define i32 @conv_shared_cluster_to_generic(ptr addrspace(7) %ptr) {
1839
; ALL-LABEL: conv_generic_to_shared_cluster
1940
define i32 @conv_generic_to_shared_cluster(ptr %ptr) {
2041
; CLS32: cvta.to.shared::cluster.u32
21-
; CLS64: cvta.to.shared::cluster.u64
22-
; PTRCONV: cvt.u32.u64
23-
; NOPTRCONV-NOT: cvt.u32.u64
24-
; ALL: ld.shared::cluster.u32
42+
; NOPTRCONV-LABEL: conv_generic_to_shared_cluster(
43+
; NOPTRCONV: {
44+
; NOPTRCONV-NEXT: .reg .b32 %r<2>;
45+
; NOPTRCONV-NEXT: .reg .b64 %rd<3>;
46+
; NOPTRCONV-EMPTY:
47+
; NOPTRCONV-NEXT: // %bb.0:
48+
; NOPTRCONV-NEXT: ld.param.u64 %rd1, [conv_generic_to_shared_cluster_param_0];
49+
; NOPTRCONV-NEXT: cvta.to.shared::cluster.u64 %rd2, %rd1;
50+
; NOPTRCONV-NEXT: ld.shared::cluster.u32 %r1, [%rd2];
51+
; NOPTRCONV-NEXT: st.param.b32 [func_retval0], %r1;
52+
; NOPTRCONV-NEXT: ret;
53+
;
54+
; PTRCONV-LABEL: conv_generic_to_shared_cluster(
55+
; PTRCONV: {
56+
; PTRCONV-NEXT: .reg .b32 %r<3>;
57+
; PTRCONV-NEXT: .reg .b64 %rd<3>;
58+
; PTRCONV-EMPTY:
59+
; PTRCONV-NEXT: // %bb.0:
60+
; PTRCONV-NEXT: ld.param.u64 %rd1, [conv_generic_to_shared_cluster_param_0];
61+
; PTRCONV-NEXT: cvta.to.shared::cluster.u64 %rd2, %rd1;
62+
; PTRCONV-NEXT: cvt.u32.u64 %r1, %rd2;
63+
; PTRCONV-NEXT: ld.shared::cluster.u32 %r2, [%r1];
64+
; PTRCONV-NEXT: st.param.b32 [func_retval0], %r2;
65+
; PTRCONV-NEXT: ret;
2566
%specptr = addrspacecast ptr %ptr to ptr addrspace(7)
2667
%val = load i32, ptr addrspace(7) %specptr
2768
ret i32 %val
2869
}
2970

3071
; ALL-LABEL: conv_shared_to_shared_cluster
3172
define i32 @conv_shared_to_shared_cluster(ptr addrspace(3) %ptr) {
32-
; CLS64: cvta.shared.u64
33-
; CLS64: cvta.to.shared::cluster.u64
34-
; ALL: ld.shared::cluster.u32
73+
; NOPTRCONV-LABEL: conv_shared_to_shared_cluster(
74+
; NOPTRCONV: {
75+
; NOPTRCONV-NEXT: .reg .b32 %r<2>;
76+
; NOPTRCONV-NEXT: .reg .b64 %rd<4>;
77+
; NOPTRCONV-EMPTY:
78+
; NOPTRCONV-NEXT: // %bb.0:
79+
; NOPTRCONV-NEXT: ld.param.u64 %rd1, [conv_shared_to_shared_cluster_param_0];
80+
; NOPTRCONV-NEXT: cvta.shared.u64 %rd2, %rd1;
81+
; NOPTRCONV-NEXT: cvta.to.shared::cluster.u64 %rd3, %rd2;
82+
; NOPTRCONV-NEXT: ld.shared::cluster.u32 %r1, [%rd3];
83+
; NOPTRCONV-NEXT: st.param.b32 [func_retval0], %r1;
84+
; NOPTRCONV-NEXT: ret;
85+
;
86+
; PTRCONV-LABEL: conv_shared_to_shared_cluster(
87+
; PTRCONV: {
88+
; PTRCONV-NEXT: .reg .b32 %r<4>;
89+
; PTRCONV-NEXT: .reg .b64 %rd<4>;
90+
; PTRCONV-EMPTY:
91+
; PTRCONV-NEXT: // %bb.0:
92+
; PTRCONV-NEXT: ld.param.u32 %r1, [conv_shared_to_shared_cluster_param_0];
93+
; PTRCONV-NEXT: cvt.u64.u32 %rd1, %r1;
94+
; PTRCONV-NEXT: cvta.shared.u64 %rd2, %rd1;
95+
; PTRCONV-NEXT: cvta.to.shared::cluster.u64 %rd3, %rd2;
96+
; PTRCONV-NEXT: cvt.u32.u64 %r2, %rd3;
97+
; PTRCONV-NEXT: ld.shared::cluster.u32 %r3, [%r2];
98+
; PTRCONV-NEXT: st.param.b32 [func_retval0], %r3;
99+
; PTRCONV-NEXT: ret;
35100
%specptr = addrspacecast ptr addrspace(3) %ptr to ptr addrspace(7)
36101
%val = load i32, ptr addrspace(7) %specptr
37102
ret i32 %val
38103
}
39104

40105
; ALL-LABEL: conv_shared_cluster_to_shared
41106
define i32 @conv_shared_cluster_to_shared(ptr addrspace(7) %ptr) {
42-
; CLS64: cvta.shared::cluster.u64
43-
; CLS64: cvta.to.shared.u64
44-
; ALL: ld.shared.u32
107+
; NOPTRCONV-LABEL: conv_shared_cluster_to_shared(
108+
; NOPTRCONV: {
109+
; NOPTRCONV-NEXT: .reg .b32 %r<2>;
110+
; NOPTRCONV-NEXT: .reg .b64 %rd<4>;
111+
; NOPTRCONV-EMPTY:
112+
; NOPTRCONV-NEXT: // %bb.0:
113+
; NOPTRCONV-NEXT: ld.param.u64 %rd1, [conv_shared_cluster_to_shared_param_0];
114+
; NOPTRCONV-NEXT: cvta.shared::cluster.u64 %rd2, %rd1;
115+
; NOPTRCONV-NEXT: cvta.to.shared.u64 %rd3, %rd2;
116+
; NOPTRCONV-NEXT: ld.shared.u32 %r1, [%rd3];
117+
; NOPTRCONV-NEXT: st.param.b32 [func_retval0], %r1;
118+
; NOPTRCONV-NEXT: ret;
119+
;
120+
; PTRCONV-LABEL: conv_shared_cluster_to_shared(
121+
; PTRCONV: {
122+
; PTRCONV-NEXT: .reg .b32 %r<4>;
123+
; PTRCONV-NEXT: .reg .b64 %rd<4>;
124+
; PTRCONV-EMPTY:
125+
; PTRCONV-NEXT: // %bb.0:
126+
; PTRCONV-NEXT: ld.param.u32 %r1, [conv_shared_cluster_to_shared_param_0];
127+
; PTRCONV-NEXT: cvt.u64.u32 %rd1, %r1;
128+
; PTRCONV-NEXT: cvta.shared::cluster.u64 %rd2, %rd1;
129+
; PTRCONV-NEXT: cvta.to.shared.u64 %rd3, %rd2;
130+
; PTRCONV-NEXT: cvt.u32.u64 %r2, %rd3;
131+
; PTRCONV-NEXT: ld.shared.u32 %r3, [%r2];
132+
; PTRCONV-NEXT: st.param.b32 [func_retval0], %r3;
133+
; PTRCONV-NEXT: ret;
45134
%specptr = addrspacecast ptr addrspace(7) %ptr to ptr addrspace(3)
46135
%val = load i32, ptr addrspace(3) %specptr
47136
ret i32 %val
48-
}
137+
}

‎llvm/test/CodeGen/NVPTX/distributed-shared-cluster.ll

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -4,14 +4,6 @@
44

55
target triple = "nvptx64-nvidia-cuda"
66

7-
@llvm.used = appending global [5 x ptr] [
8-
ptr @test_distributed_shared_cluster_common,
9-
ptr @test_distributed_shared_cluster_float_atomic,
10-
ptr @test_distributed_shared_cluster_int_atomic,
11-
ptr @test_distributed_shared_cluster_bitwise_atomic,
12-
ptr @test_distributed_shared_cluster_cmpxchg
13-
], section "llvm.metadata"
14-
157
declare ptr addrspace(7) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3), i32)
168
declare i1 @llvm.nvvm.isspacep.shared.cluster(ptr)
179
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()

‎llvm/test/Transforms/InferAddressSpaces/NVPTX/isspacep.ll

Lines changed: 5 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -114,17 +114,14 @@ entry:
114114
ret i1 %val
115115
}
116116

117-
define i1 @test_isspacep_cluster_shared_unsure(ptr addrspace(3) %addr) {
118-
; CHECK-LABEL: define i1 @test_isspacep_cluster_shared_unsure(
119-
; CHECK-SAME: ptr addrspace(3) [[ADDR:%.*]]) {
117+
define i1 @test_isspacep_cluster_shared_true(ptr addrspace(7) %addr) {
118+
; CHECK-LABEL: define i1 @test_isspacep_cluster_shared_true(
119+
; CHECK-SAME: ptr addrspace(7) [[ADDR:%.*]]) {
120120
; CHECK-NEXT: [[ENTRY:.*:]]
121-
; CHECK-NEXT: [[ADDR1:%.*]] = getelementptr i8, ptr addrspace(3) [[ADDR]], i32 10
122-
; CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(3) [[ADDR1]] to ptr
123-
; CHECK-NEXT: [[VAL:%.*]] = call i1 @llvm.nvvm.isspacep.shared.cluster(ptr [[TMP0]])
124-
; CHECK-NEXT: ret i1 [[VAL]]
121+
; CHECK-NEXT: ret i1 true
125122
;
126123
entry:
127-
%addr0 = addrspacecast ptr addrspace(3) %addr to ptr
124+
%addr0 = addrspacecast ptr addrspace(7) %addr to ptr
128125
%addr1 = getelementptr i8, ptr %addr0, i32 10
129126
%val = call i1 @llvm.nvvm.isspacep.shared.cluster(ptr %addr1)
130127
ret i1 %val

0 commit comments

Comments
 (0)