Skip to content

Commit b921023

Browse files
ro-iro-i
committed
implement feedback
1 parent 9c8af8a commit b921023

File tree

4 files changed

+58
-56
lines changed

4 files changed

+58
-56
lines changed

‎llvm/lib/Target/AMDGPU/SIISelLowering.cpp‎

Lines changed: 8 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -8619,6 +8619,11 @@ SDValue SITargetLowering::lowerWorkitemID(SelectionDAG &DAG, SDValue Op,
86198619
if (MaxID == 0)
86208620
return DAG.getConstant(0, SL, MVT::i32);
86218621

8622+
// It's undefined behavior if a function marked with the amdgpu-no-*
8623+
// attributes uses the corresponding intrinsic.
8624+
if (!Arg)
8625+
return DAG.getUNDEF(EVT::getIntegerVT(*DAG.getContext(), 32));
8626+
86228627
SDValue Val = loadInputValue(DAG, &AMDGPU::VGPR_32RegClass, MVT::i32,
86238628
SDLoc(DAG.getEntryNode()), Arg);
86248629

@@ -8790,28 +8795,11 @@ SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
87908795
AMDGPUFunctionArgInfo::LDS_KERNEL_ID);
87918796
}
87928797
case Intrinsic::amdgcn_workitem_id_x:
8793-
if (!MFI->getArgInfo().WorkItemIDX) {
8794-
// It's undefined behavior if a function marked with the amdgpu-no-*
8795-
// attributes uses the corresponding intrinsic.
8796-
return DAG.getConstant(0, SDLoc(Op),
8797-
EVT::getIntegerVT(*DAG.getContext(), 32));
8798-
} else {
8799-
return lowerWorkitemID(DAG, Op, 0, MFI->getArgInfo().WorkItemIDX);
8800-
}
8798+
return lowerWorkitemID(DAG, Op, 0, MFI->getArgInfo().WorkItemIDX);
88018799
case Intrinsic::amdgcn_workitem_id_y:
8802-
if (!MFI->getArgInfo().WorkItemIDY) {
8803-
return DAG.getConstant(0, SDLoc(Op),
8804-
EVT::getIntegerVT(*DAG.getContext(), 32));
8805-
} else {
8806-
return lowerWorkitemID(DAG, Op, 1, MFI->getArgInfo().WorkItemIDY);
8807-
}
8800+
return lowerWorkitemID(DAG, Op, 1, MFI->getArgInfo().WorkItemIDY);
88088801
case Intrinsic::amdgcn_workitem_id_z:
8809-
if (!MFI->getArgInfo().WorkItemIDZ) {
8810-
return DAG.getConstant(0, SDLoc(Op),
8811-
EVT::getIntegerVT(*DAG.getContext(), 32));
8812-
} else {
8813-
return lowerWorkitemID(DAG, Op, 2, MFI->getArgInfo().WorkItemIDZ);
8814-
}
8802+
return lowerWorkitemID(DAG, Op, 2, MFI->getArgInfo().WorkItemIDZ);
88158803
case Intrinsic::amdgcn_wavefrontsize:
88168804
return DAG.getConstant(MF.getSubtarget<GCNSubtarget>().getWavefrontSize(),
88178805
SDLoc(Op), MVT::i32);

‎llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.workitem.id.ll‎

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@
99
; RUN: llc -global-isel -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx90a -verify-machineinstrs < %t.v4.ll | FileCheck -check-prefixes=ALL,PACKED-TID %s
1010
; RUN: llc -global-isel -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx1100 -verify-machineinstrs -amdgpu-enable-vopd=0 < %t.v4.ll | FileCheck -check-prefixes=ALL,PACKED-TID %s
1111
; RUN: llc -global-isel -mtriple=amdgcn-unknown-amdhsa --amdhsa-code-object-version=6 -mcpu=gfx11-generic -verify-machineinstrs -amdgpu-enable-vopd=0 < %t.v6.ll | FileCheck -check-prefixes=ALL,PACKED-TID %s
12+
; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx942 -o - %t.v4.ll | FileCheck --check-prefixes=UNDEF %s
1213

1314
declare i32 @llvm.amdgcn.workitem.id.x() #0
1415
declare i32 @llvm.amdgcn.workitem.id.y() #0
@@ -195,6 +196,30 @@ define amdgpu_kernel void @test_reqd_workgroup_size_z_only(ptr %out) !reqd_work_
195196
ret void
196197
}
197198

199+
define amdgpu_kernel void @undefined_workitem_x_only() {
200+
; UNDEF-LABEL: undefined_workitem_x_only:
201+
; UNDEF: ; %bb.0:
202+
; UNDEF-NEXT: s_endpgm
203+
%id.x = call i32 @llvm.amdgcn.workitem.id.x()
204+
ret void
205+
}
206+
207+
define amdgpu_kernel void @undefined_workitem_y_only() {
208+
; UNDEF-LABEL: undefined_workitem_y_only:
209+
; UNDEF: ; %bb.0:
210+
; UNDEF-NEXT: s_endpgm
211+
%id.y = call i32 @llvm.amdgcn.workitem.id.y()
212+
ret void
213+
}
214+
215+
define amdgpu_kernel void @undefined_workitem_z_only() {
216+
; UNDEF-LABEL: undefined_workitem_z_only:
217+
; UNDEF: ; %bb.0:
218+
; UNDEF-NEXT: s_endpgm
219+
%id.z = call i32 @llvm.amdgcn.workitem.id.z()
220+
ret void
221+
}
222+
198223
attributes #0 = { nounwind readnone }
199224
attributes #1 = { nounwind }
200225

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.workitem.id-undefined-attr.ll‎

Lines changed: 0 additions & 36 deletions
This file was deleted.

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.workitem.id.ll‎

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
; RUN: llc -mtriple=amdgcn-unknown-mesa3d -mcpu=tonga -mattr=-flat-for-global < %t.bc | FileCheck -check-prefixes=ALL,MESA3D,UNPACKED %s
66
; RUN: llc -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx90a < %t.bc | FileCheck -check-prefixes=ALL,PACKED-TID %s
77
; RUN: llc -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx1100 -amdgpu-enable-vopd=0 < %t.bc | FileCheck -check-prefixes=ALL,PACKED-TID %s
8+
; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -o - %s | FileCheck --check-prefix=UNDEF %s
89

910
declare i32 @llvm.amdgcn.workitem.id.x() #0
1011
declare i32 @llvm.amdgcn.workitem.id.y() #0
@@ -128,6 +129,30 @@ define amdgpu_kernel void @test_reqd_workgroup_size_z_only(ptr %out) !reqd_work_
128129
ret void
129130
}
130131

132+
define amdgpu_kernel void @undefined_workitem_x_only() {
133+
; UNDEF-LABEL: undefined_workitem_x_only:
134+
; UNDEF: ; %bb.0:
135+
; UNDEF-NEXT: s_endpgm
136+
%id.x = call i32 @llvm.amdgcn.workitem.id.x()
137+
ret void
138+
}
139+
140+
define amdgpu_kernel void @undefined_workitem_y_only() {
141+
; UNDEF-LABEL: undefined_workitem_y_only:
142+
; UNDEF: ; %bb.0:
143+
; UNDEF-NEXT: s_endpgm
144+
%id.y = call i32 @llvm.amdgcn.workitem.id.y()
145+
ret void
146+
}
147+
148+
define amdgpu_kernel void @undefined_workitem_z_only() {
149+
; UNDEF-LABEL: undefined_workitem_z_only:
150+
; UNDEF: ; %bb.0:
151+
; UNDEF-NEXT: s_endpgm
152+
%id.z = call i32 @llvm.amdgcn.workitem.id.z()
153+
ret void
154+
}
155+
131156
attributes #0 = { nounwind readnone }
132157
attributes #1 = { nounwind }
133158

0 commit comments

Comments
 (0)