AMDGPU: Add new amdgcn workitem intrinsics

[arsenm]

Summary

These use the correct prefix and follow the HSA naming convention
rather than the config register option names.

[arsenm]

Stop running HSA tests in existing work-item-intrinsics tests. Some of these do not work correctly with HSA, and will be errors in a future commit.

[• tstellarAMD]

LGTM.

[arsenm]

r259293

include/llvm/IR/IntrinsicsAMDGPU.td

 //===----------------------------------------------------------------------===//
 
 class AMDGPUReadPreloadRegisterIntrinsic<string name>
   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
     GCCBuiltin<name>;
 
 let TargetPrefix = "r600" in {
 
-class R600ReadPreloadRegisterIntrinsic<string name>
-  : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
-    GCCBuiltin<name>;
-
-multiclass R600ReadPreloadRegisterIntrinsic_xyz<string prefix> {
-  def _x : R600ReadPreloadRegisterIntrinsic<!strconcat(prefix, "_x")>;
-  def _y : R600ReadPreloadRegisterIntrinsic<!strconcat(prefix, "_y")>;
-  def _z : R600ReadPreloadRegisterIntrinsic<!strconcat(prefix, "_z")>;
+multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz<string prefix> {
+  def _x : AMDGPUReadPreloadRegisterIntrinsic<!strconcat(prefix, "_x")>;
+  def _y : AMDGPUReadPreloadRegisterIntrinsic<!strconcat(prefix, "_y")>;
+  def _z : AMDGPUReadPreloadRegisterIntrinsic<!strconcat(prefix, "_z")>;
 }
 
-defm int_r600_read_global_size : R600ReadPreloadRegisterIntrinsic_xyz <
+defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz <
                                        "__builtin_r600_read_global_size">;
-defm int_r600_read_local_size : R600ReadPreloadRegisterIntrinsic_xyz <
+defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz <
                                        "__builtin_r600_read_local_size">;
-defm int_r600_read_ngroups : R600ReadPreloadRegisterIntrinsic_xyz <
+defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz <
                                        "__builtin_r600_read_ngroups">;
-defm int_r600_read_tgid : R600ReadPreloadRegisterIntrinsic_xyz <
+defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz <
                                        "__builtin_r600_read_tgid">;
-defm int_r600_read_tidig : R600ReadPreloadRegisterIntrinsic_xyz <
+defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz <
                                        "__builtin_r600_read_tidig">;
 
 def int_r600_rat_store_typed :
   // 1st parameter: Data
   // 2nd parameter: Index
   // 3rd parameter: Constant RAT ID
   Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>,
   GCCBuiltin<"__builtin_r600_rat_store_typed">;

 let TargetPrefix = "AMDGPU" in {
 def int_AMDGPU_ldexp : Intrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
 >;
 }
 
 let TargetPrefix = "amdgcn" in {
 
+defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz <
+  "__builtin_amdgcn_workitem_id">;
+defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz <
+  "__builtin_amdgcn_workgroup_id">;
+
 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
   Intrinsic<[], [], [IntrConvergent]>;
 
 def int_amdgcn_umul24 : GCCBuiltin<"__builtin_amdgcn_umul24">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]
 >;
 
 def int_amdgcn_umulhi24 : GCCBuiltin<"__builtin_amdgcn_umulhi24">,

lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp

bool AMDGPUAnnotateKernelFeatures::addAttrsForIntrinsics(
   return Changed;
 }
 
 bool AMDGPUAnnotateKernelFeatures::runOnModule(Module &M) {
   Triple TT(M.getTargetTriple());
 
   static const StringRef IntrinsicToAttr[][2] = {
     // .x omitted
+    { "llvm.amdgcn.workitem.id.y", "amdgpu-work-item-id-y" },
+    { "llvm.amdgcn.workitem.id.z", "amdgpu-work-item-id-z" },
+
+    { "llvm.amdgcn.workgroup.id.y", "amdgpu-work-group-id-y" },
+    { "llvm.amdgcn.workgroup.id.z", "amdgpu-work-group-id-z" },
+
     { "llvm.r600.read.tgid.y", "amdgpu-work-group-id-y" },
     { "llvm.r600.read.tgid.z", "amdgpu-work-group-id-z" },
 
     // .x omitted
     { "llvm.r600.read.tidig.y", "amdgpu-work-item-id-y" },
     { "llvm.r600.read.tidig.z", "amdgpu-work-item-id-z" }
 
   };

lib/Target/AMDGPU/SIISelLowering.cpp

SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
   case Intrinsic::r600_read_local_size_z:
     return lowerImplicitZextParam(DAG, Op, MVT::i16,
                                   SI::KernelInputOffsets::LOCAL_SIZE_Z);
   case Intrinsic::amdgcn_read_workdim:
   case AMDGPUIntrinsic::AMDGPU_read_workdim: // Legacy name.
     // Really only 2 bits.
     return lowerImplicitZextParam(DAG, Op, MVT::i8,
                                   getImplicitParameterOffset(MFI, GRID_DIM));
+  case Intrinsic::amdgcn_workgroup_id_x:
   case Intrinsic::r600_read_tgid_x:
     return CreateLiveInRegister(DAG, &AMDGPU::SReg_32RegClass,
       TRI->getPreloadedValue(MF, SIRegisterInfo::WORKGROUP_ID_X), VT);
+  case Intrinsic::amdgcn_workgroup_id_y:
   case Intrinsic::r600_read_tgid_y:
     return CreateLiveInRegister(DAG, &AMDGPU::SReg_32RegClass,
       TRI->getPreloadedValue(MF, SIRegisterInfo::WORKGROUP_ID_Y), VT);
+  case Intrinsic::amdgcn_workgroup_id_z:
   case Intrinsic::r600_read_tgid_z:
     return CreateLiveInRegister(DAG, &AMDGPU::SReg_32RegClass,
       TRI->getPreloadedValue(MF, SIRegisterInfo::WORKGROUP_ID_Z), VT);
+  case Intrinsic::amdgcn_workitem_id_x:
   case Intrinsic::r600_read_tidig_x:
     return CreateLiveInRegister(DAG, &AMDGPU::VGPR_32RegClass,
       TRI->getPreloadedValue(MF, SIRegisterInfo::WORKITEM_ID_X), VT);
+  case Intrinsic::amdgcn_workitem_id_y:
   case Intrinsic::r600_read_tidig_y:
     return CreateLiveInRegister(DAG, &AMDGPU::VGPR_32RegClass,
       TRI->getPreloadedValue(MF, SIRegisterInfo::WORKITEM_ID_Y), VT);
+  case Intrinsic::amdgcn_workitem_id_z:
   case Intrinsic::r600_read_tidig_z:
     return CreateLiveInRegister(DAG, &AMDGPU::VGPR_32RegClass,
       TRI->getPreloadedValue(MF, SIRegisterInfo::WORKITEM_ID_Z), VT);
   case AMDGPUIntrinsic::SI_load_const: {
     SDValue Ops[] = {
       Op.getOperand(1),
       Op.getOperand(2)
     };

test/CodeGen/AMDGPU/llvm.amdgcn.workgroup.id.ll

+; RUN: llc -march=amdgcn -mtriple=amdgcn-unknown-amdhsa -verify-machineinstrs < %s | FileCheck -check-prefix=ALL -check-prefix=HSA %s
+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=ALL -check-prefix=MESA %s
+
+declare i32 @llvm.amdgcn.workgroup.id.x() #0
+declare i32 @llvm.amdgcn.workgroup.id.y() #0
+declare i32 @llvm.amdgcn.workgroup.id.z() #0
+
+; ALL-LABEL {{^}}test_workgroup_id_x:
+
+; HSA: compute_pgm_rsrc2_tgid_x_en = 1
+; HSA: compute_pgm_rsrc2_tgid_y_en = 0
+; HSA: compute_pgm_rsrc2_tgid_z_en = 0
+
+; MESA: v_mov_b32_e32 [[VCOPY:v[0-9]+]], s2{{$}}
+; HSA: v_mov_b32_e32 [[VCOPY:v[0-9]+]], s6{{$}}
+
+; ALL-NOT: [[VCOPY]]
+; ALL: {{buffer|flat}}_store_dword [[VCOPY]]
+define void @test_workgroup_id_x(i32 addrspace(1)* %out) #1 {
+  %id = call i32 @llvm.amdgcn.workgroup.id.x()
+  store i32 %id, i32 addrspace(1)* %out
+  ret void
+}
+
+; ALL-LABEL {{^}}test_workgroup_id_y:
+; HSA: compute_pgm_rsrc2_tgid_x_en = 1
+; HSA: compute_pgm_rsrc2_tgid_y_en = 1
+; HSA: compute_pgm_rsrc2_tgid_z_en = 0
+
+; MESA: v_mov_b32_e32 [[VCOPY:v[0-9]+]], s3{{$}}
+; HSA: v_mov_b32_e32 [[VCOPY:v[0-9]+]], s7{{$}}
+
+; ALL-NOT: [[VCOPY]]
+; ALL: {{buffer|flat}}_store_dword [[VCOPY]]
+define void @test_workgroup_id_y(i32 addrspace(1)* %out) #1 {
+  %id = call i32 @llvm.amdgcn.workgroup.id.y()
+  store i32 %id, i32 addrspace(1)* %out
+  ret void
+}
+
+; ALL-LABEL {{^}}test_workgroup_id_z:
+; HSA: compute_pgm_rsrc2_tgid_x_en = 1
+; HSA: compute_pgm_rsrc2_tgid_y_en = 1
+; HSA: compute_pgm_rsrc2_tgid_z_en = 1
+
+; MESA: v_mov_b32_e32 [[VCOPY:v[0-9]+]], s4{{$}}
+; HSA: v_mov_b32_e32 [[VCOPY:v[0-9]+]], s8{{$}}
+
+; ALL-NOT: [[VCOPY]]
+; ALL: {{buffer|flat}}_store_dword [[VCOPY]]
+define void @test_workgroup_id_z(i32 addrspace(1)* %out) #1 {
+  %id = call i32 @llvm.amdgcn.workgroup.id.z()
+  store i32 %id, i32 addrspace(1)* %out
+  ret void
+}
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { nounwind }

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

+; RUN: llc -march=amdgcn -mtriple=amdgcn-unknown-amdhsa -verify-machineinstrs < %s | FileCheck -check-prefix=ALL -check-prefix=HSA %s
+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=ALL -check-prefix=MESA %s
+
+declare i32 @llvm.amdgcn.workitem.id.x() #0
+declare i32 @llvm.amdgcn.workitem.id.y() #0
+declare i32 @llvm.amdgcn.workitem.id.z() #0
+
+; ALL-LABEL {{^}}test_workitem_id_x:
+; HSA: compute_pgm_rsrc2_tidig_comp_cnt = 0
+
+; ALL-NOT: v0
+; ALL: {{buffer|flat}}_store_dword v0
+define void @test_workitem_id_x(i32 addrspace(1)* %out) #1 {
+  %id = call i32 @llvm.amdgcn.workitem.id.x()
+  store i32 %id, i32 addrspace(1)* %out
+  ret void
+}
+
+; ALL-LABEL {{^}}test_workitem_id_y:
+; HSA: compute_pgm_rsrc2_tidig_comp_cnt = 1
+
+; ALL-NOT: v1
+; ALL: {{buffer|flat}}_store_dword v1
+define void @test_workitem_id_y(i32 addrspace(1)* %out) #1 {
+  %id = call i32 @llvm.amdgcn.workitem.id.y()
+  store i32 %id, i32 addrspace(1)* %out
+  ret void
+}
+
+; ALL-LABEL {{^}}test_workitem_id_z:
+; HSA: compute_pgm_rsrc2_tidig_comp_cnt = 2
+
+; ALL-NOT: v2
+; ALL: {{buffer|flat}}_store_dword v2
+define void @test_workitem_id_z(i32 addrspace(1)* %out) #1 {
+  %id = call i32 @llvm.amdgcn.workitem.id.z()
+  store i32 %id, i32 addrspace(1)* %out
+  ret void
+}
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { nounwind }