[llvm] [AMDGPU] SelDAG: fix lowering of undefined workitem intrinsics (PR #126058)

[Robert Imschweiler via llvm-commits]

================
@@ -128,6 +129,30 @@ define amdgpu_kernel void @test_reqd_workgroup_size_z_only(ptr %out) !reqd_work_
   ret void
 }
 
+define amdgpu_kernel void @undefined_workitem_x_only() {
+; UNDEF-LABEL: undefined_workitem_x_only:
+; UNDEF:       ; %bb.0:
+; UNDEF-NEXT:    s_endpgm
+  %id.x = call i32 @llvm.amdgcn.workitem.id.x()
+  ret void
+}
+
+define amdgpu_kernel void @undefined_workitem_y_only() {
+; UNDEF-LABEL: undefined_workitem_y_only:
+; UNDEF:       ; %bb.0:
+; UNDEF-NEXT:    s_endpgm
+  %id.y = call i32 @llvm.amdgcn.workitem.id.y()
+  ret void
+}
+
+define amdgpu_kernel void @undefined_workitem_z_only() {
+; UNDEF-LABEL: undefined_workitem_z_only:
+; UNDEF:       ; %bb.0:
+; UNDEF-NEXT:    s_endpgm
+  %id.z = call i32 @llvm.amdgcn.workitem.id.z()
----------------
ro-i wrote:

I changed the tests and moved them to a separate file because they need to use a different calling convention than `amdgpu_kernel` so that the workitem.id registers are not defined. For the SelDAG case, I need `-O0`, otherwise everything is optimized away (which seems maybe a bit weird).

https://github.com/llvm/llvm-project/pull/126058