AMDGPU/SI: Add implicitarg.ptr intrinsic.

[jvesely]

Summary

Points to the start of implicit arguments (appended after explicit arguments)

[arsenm]

We don't actually want these. We now have the kernarg.segment.ptr intrinsic, so the library should just directly read the offsets from there

[jvesely]

In D20297#431268, @arsenm wrote:

We don't actually want these. We now have the kernarg.segment.ptr intrinsic, so the library should just directly read the offsets from there

I assume it points to the beginning of the kernel args.
Is the idea to have another intrinsic to read from the end (workdim, global offset), or is it OK to have intrinsics for those?

[arsenm]

In D20297#431298, @jvesely wrote:

In D20297#431268, @arsenm wrote:

We don't actually want these. We now have the kernarg.segment.ptr intrinsic, so the library should just directly read the offsets from there

I assume it points to the beginning of the kernel args.
Is the idea to have another intrinsic to read from the end (workdim, global offset), or is it OK to have intrinsics for those?

We should fix clover to not read from the end of the arguments (I thought this is what it already did)?

[jvesely]

In D20297#431315, @arsenm wrote:

In D20297#431298, @jvesely wrote:

In D20297#431268, @arsenm wrote:

We don't actually want these. We now have the kernarg.segment.ptr intrinsic, so the library should just directly read the offsets from there

I assume it points to the beginning of the kernel args.
Is the idea to have another intrinsic to read from the end (workdim, global offset), or is it OK to have intrinsics for those?

We should fix clover to not read from the end of the arguments (I thought this is what it already did)?

some are before the kernel args (global size, local size, ngroups -- this is done by radeonsi/r600 mesa driver). others are after the kernel arguments (work_dim, global-offset -- this is done by clover).
The intention was to move all implicit arguments after the explicit ones, so new implicit arguments can be added without breaking ABI (moving explicit arguments).

[arsenm]

In D20297#431388, @jvesely wrote:

In D20297#431315, @arsenm wrote:

In D20297#431298, @jvesely wrote:

In D20297#431268, @arsenm wrote:

We don't actually want these. We now have the kernarg.segment.ptr intrinsic, so the library should just directly read the offsets from there

I assume it points to the beginning of the kernel args.
Is the idea to have another intrinsic to read from the end (workdim, global offset), or is it OK to have intrinsics for those?

We should fix clover to not read from the end of the arguments (I thought this is what it already did)?

some are before the kernel args (global size, local size, ngroups -- this is done by radeonsi/r600 mesa driver). others are after the kernel arguments (work_dim, global-offset -- this is done by clover).
The intention was to move all implicit arguments after the explicit ones, so new implicit arguments can be added without breaking ABI (moving explicit arguments).

I think clover should move towards matching the HSA ABI closer. Most of the implicit arguments would then be user SGPR inputs like HSA uses, and the number of implicit args would be reduced.

[jvesely]

In D20297#431404, @arsenm wrote:

In D20297#431388, @jvesely wrote:

In D20297#431315, @arsenm wrote:

In D20297#431298, @jvesely wrote:

In D20297#431268, @arsenm wrote:

We don't actually want these. We now have the kernarg.segment.ptr intrinsic, so the library should just directly read the offsets from there

I assume it points to the beginning of the kernel args.
Is the idea to have another intrinsic to read from the end (workdim, global offset), or is it OK to have intrinsics for those?

We should fix clover to not read from the end of the arguments (I thought this is what it already did)?

some are before the kernel args (global size, local size, ngroups -- this is done by radeonsi/r600 mesa driver). others are after the kernel arguments (work_dim, global-offset -- this is done by clover).
The intention was to move all implicit arguments after the explicit ones, so new implicit arguments can be added without breaking ABI (moving explicit arguments).

I think clover should move towards matching the HSA ABI closer. Most of the implicit arguments would then be user SGPR inputs like HSA uses, and the number of implicit args would be reduced.

new implicit arguments were appended to allow newer mesa to work with older llvm (without ifdef hell) any kind of significant abi change would break that (otherwise it'd be in one palce and we'd avoi this problem).
I'm not sure how much of a problem it is and what other (possible) users of clover might want. anyway, it's beyond scope of me trying to make get_global_offset() work. would you be OK if I restricted the changes to r600?

[jvesely]

Drop new amdgcn intrinsics. make kernarg segment mesa compatible instead.

[• tstellarAMD]

In D20297#431404, @arsenm wrote:

In D20297#431388, @jvesely wrote:

In D20297#431315, @arsenm wrote:

In D20297#431298, @jvesely wrote:

In D20297#431268, @arsenm wrote:

We don't actually want these. We now have the kernarg.segment.ptr intrinsic, so the library should just directly read the offsets from there

I assume it points to the beginning of the kernel args.
Is the idea to have another intrinsic to read from the end (workdim, global offset), or is it OK to have intrinsics for those?

We should fix clover to not read from the end of the arguments (I thought this is what it already did)?

some are before the kernel args (global size, local size, ngroups -- this is done by radeonsi/r600 mesa driver). others are after the kernel arguments (work_dim, global-offset -- this is done by clover).
The intention was to move all implicit arguments after the explicit ones, so new implicit arguments can be added without breaking ABI (moving explicit arguments).

I think clover should move towards matching the HSA ABI closer. Most of the implicit arguments would then be user SGPR inputs like HSA uses, and the number of implicit args would be reduced.

I agree. I actually started working on this last week. I think all implicit args that aren't passed in SGPRs should be at the end of the kernarg segment.

lib/Target/AMDGPU/SIISelLowering.cpp

1574–1579

There should be a separate intrinsic which points to the start of implicit args. As the ABI is currently, the library could use kernarg_segment_ptr to load the workgroup size information.

[jvesely]

In D20297#435687, @tstellarAMD wrote:

In D20297#431404, @arsenm wrote:

In D20297#431388, @jvesely wrote:

In D20297#431315, @arsenm wrote:

In D20297#431298, @jvesely wrote:

In D20297#431268, @arsenm wrote:

We don't actually want these. We now have the kernarg.segment.ptr intrinsic, so the library should just directly read the offsets from there

I assume it points to the beginning of the kernel args.
Is the idea to have another intrinsic to read from the end (workdim, global offset), or is it OK to have intrinsics for those?

We should fix clover to not read from the end of the arguments (I thought this is what it already did)?

some are before the kernel args (global size, local size, ngroups -- this is done by radeonsi/r600 mesa driver). others are after the kernel arguments (work_dim, global-offset -- this is done by clover).
The intention was to move all implicit arguments after the explicit ones, so new implicit arguments can be added without breaking ABI (moving explicit arguments).

I think clover should move towards matching the HSA ABI closer. Most of the implicit arguments would then be user SGPR inputs like HSA uses, and the number of implicit args would be reduced.

This is a bit confusing. If the implicit args should be in SGPRs, wouldn't we need one intrinsic per implicit arg? and how would the number of implicit args be reduced? the only redundant one is global size (libclc computes it as num_groups * local_size).

I agree. I actually started working on this last week. I think all implicit args that aren't passed in SGPRs should be at the end of the kernarg segment.

My idea was to switch work_dim and newly implemented global_offset, as those are already appended (the rest can be switched by patches to libclc and clover). However, doesn't this contradict Matt's suggestion to pass implicit arguments in SGPRs?

[• tstellarAMD]

I think clover should move towards matching the HSA ABI closer. Most of the implicit arguments would then be user SGPR inputs like HSA uses, and the number of implicit args would be reduced.

This is a bit confusing. If the implicit args should be in SGPRs, wouldn't we need one intrinsic per implicit arg? and how would the number of implicit args be reduced? the only redundant one is global size (libclc computes it as num_groups * local_size).

I agree. I actually started working on this last week. I think all implicit args that aren't passed in SGPRs should be at the end of the kernarg segment.

My idea was to switch work_dim and newly implemented global_offset, as those are already appended (the rest can be switched by patches to libclc and clover). However, doesn't this contradict Matt's suggestion to pass implicit arguments in SGPRs?

SGPR space is limited, so we won't be able to pass all implicit arguments this way, so some will need to be added to the kernarg buffer. The types of values that should be passed in SGPRs are things that tend to be common across all runtimes, like work-group/work-item size, scratch buffer pointers, etc.

[• tstellarAMD]

This comment was removed by • tstellarAMD.

[arsenm]

We could do what we are planning for OpenCL and add a struct pointer for future expansion at the end of some decided N bytes to reserve.

If 256 bytes were reserved at the beginning, that would be more than enough for any implicit needs and would keep the user arg base alignment the same (not sure that really matters).

[jvesely]

In D20297#436251, @tstellarAMD wrote:

I think clover should move towards matching the HSA ABI closer. Most of the implicit arguments would then be user SGPR inputs like HSA uses, and the number of implicit args would be reduced.

This is a bit confusing. If the implicit args should be in SGPRs, wouldn't we need one intrinsic per implicit arg? and how would the number of implicit args be reduced? the only redundant one is global size (libclc computes it as num_groups * local_size).

I agree. I actually started working on this last week. I think all implicit args that aren't passed in SGPRs should be at the end of the kernarg segment.

My idea was to switch work_dim and newly implemented global_offset, as those are already appended (the rest can be switched by patches to libclc and clover). However, doesn't this contradict Matt's suggestion to pass implicit arguments in SGPRs?

SGPR space is limited, so we won't be able to pass all implicit arguments this way, so some will need to be added to the kernarg buffer. The types of values that should be passed in SGPRs are things that tend to be common across all runtimes, like work-group/work-item size, scratch buffer pointers, etc.

OK, so to be specific about currently passed information.
workdim, wg_size, num_group should be eventually passed in SGPRs and therefore should have their own intrinsic, correct?
should those values also be duplicated in the kernarg segment?

the rest (global_size, global_offset) are loaded via kernargs segment ptr

do you still want two pointers (beginning of kernarg and beginning of implicit args) for mesa, or is it ok to have one if all the information is present at the appended location?

In D20297#436254, @arsenm wrote:

We could do what we are planning for OpenCL and add a struct pointer for future expansion at the end of some decided N bytes to reserve.

If 256 bytes were reserved at the beginning, that would be more than enough for any implicit needs and would keep the user arg base alignment the same (not sure that really matters).

reserving 256 bytes at the beginning would waste half a KC block for r600 (and might run into limitations for other hw), so clover is probably best of appending the information to minimize space (it should work until there are kernels with variable parameters). radeonsi driver is free to change this for GCN hw.

[• tstellarAMD]

In D20297#436378, @jvesely wrote:

In D20297#436251, @tstellarAMD wrote:

I think clover should move towards matching the HSA ABI closer. Most of the implicit arguments would then be user SGPR inputs like HSA uses, and the number of implicit args would be reduced.

This is a bit confusing. If the implicit args should be in SGPRs, wouldn't we need one intrinsic per implicit arg? and how would the number of implicit args be reduced? the only redundant one is global size (libclc computes it as num_groups * local_size).

I agree. I actually started working on this last week. I think all implicit args that aren't passed in SGPRs should be at the end of the kernarg segment.

My idea was to switch work_dim and newly implemented global_offset, as those are already appended (the rest can be switched by patches to libclc and clover). However, doesn't this contradict Matt's suggestion to pass implicit arguments in SGPRs?

SGPR space is limited, so we won't be able to pass all implicit arguments this way, so some will need to be added to the kernarg buffer. The types of values that should be passed in SGPRs are things that tend to be common across all runtimes, like work-group/work-item size, scratch buffer pointers, etc.

OK, so to be specific about currently passed information.
workdim, wg_size, num_group should be eventually passed in SGPRs and therefore should have their own intrinsic, correct?

If you look at enum PreloadedValue in SIRegisterInfo.h (this may have moved to SIMachineFunctionInfo.h by now), those are the values preloaded into SGPRs by the HSA runtime. If we are going to be changing the ABI for radeonsi clover, I would really like it to match what we use for HSA. For r600, it doesn't really matter to me what the ABI ends up being.

should those values also be duplicated in the kernarg segment?

the rest (global_size, global_offset) are loaded via kernargs segment ptr

do you still want two pointers (beginning of kernarg and beginning of implicit args) for mesa, or is it ok to have one if all the information is present at the appended location?

I think there needs to be both. The intrinsics names should correctly describe what they do, so I don't want to have kernarg.segment.ptr mean different things for different arches.

In D20297#436254, @arsenm wrote:

We could do what we are planning for OpenCL and add a struct pointer for future expansion at the end of some decided N bytes to reserve.

If 256 bytes were reserved at the beginning, that would be more than enough for any implicit needs and would keep the user arg base alignment the same (not sure that really matters).

reserving 256 bytes at the beginning would waste half a KC block for r600 (and might run into limitations for other hw), so clover is probably best of appending the information to minimize space (it should work until there are kernels with variable parameters). radeonsi driver is free to change this for GCN hw.

[jvesely]

Create new intrinsic for implicit args

[arsenm]

Why not the offset from the base pointer intrinsic rather than an intrinsic to the offset?

[jvesely]

In D20297#441156, @arsenm wrote:

Why not the offset from the base pointer intrinsic rather than an intrinsic to the offset?

Does not really matter, but pointer appeared more flexible. if you decide to move implicit args elsewhere (that is not offset to kernarg segment) you only need to reimplement the intrinsic, without updating library.

[• tstellarAMD]

In D20297#441156, @arsenm wrote:

Why not the offset from the base pointer intrinsic rather than an intrinsic to the offset?

Because if the implict args are stored after the explict kernel args, then the offset from kernarg.base.ptr to the start of implicit args is not known until compile time.

[• tstellarAMD]

This patch LGTM, but I think we should drop segment from the intrinsic name. In HSA 'segment' means address space, and we don't have a separate address space for implicit args.

[jvesely]

drop segment from the intrinsic name

[arsenm]

In D20297#442193, @tstellarAMD wrote:

In D20297#441156, @arsenm wrote:

Why not the offset from the base pointer intrinsic rather than an intrinsic to the offset?

Because if the implict args are stored after the explict kernel args, then the offset from kernarg.base.ptr to the start of implicit args is not known until compile time.

Why is this an issue? This won't be known anyway

[jvesely]

In D20297#444575, @arsenm wrote:

In D20297#442193, @tstellarAMD wrote:

In D20297#441156, @arsenm wrote:

Why not the offset from the base pointer intrinsic rather than an intrinsic to the offset?

Because if the implict args are stored after the explict kernel args, then the offset from kernarg.base.ptr to the start of implicit args is not known until compile time.

Why is this an issue? This won't be known anyway

just to be sure I understand correctly. your suggestion is to have "implicitarg.offset" so the libclc(or any other user) then uses "_builtin_kernarg_segment_ptr() + _builtin_implict_arg_offset()" to read the implicit arguments?

[jvesely]

rebase after r272512

[• tstellarAMD]

In D20297#449024, @jvesely wrote:

In D20297#444575, @arsenm wrote:

In D20297#442193, @tstellarAMD wrote:

In D20297#441156, @arsenm wrote:

Why not the offset from the base pointer intrinsic rather than an intrinsic to the offset?

Because if the implict args are stored after the explict kernel args, then the offset from kernarg.base.ptr to the start of implicit args is not known until compile time.

Why is this an issue? This won't be known anyway

just to be sure I understand correctly. your suggestion is to have "implicitarg.offset" so the libclc(or any other user) then uses "_builtin_kernarg_segment_ptr() + _builtin_implict_arg_offset()" to read the implicit arguments?

I have no preference either way on this. Since this patch has been outstanding for a while, I say just commit it. We can always add _builtin_implict_arg_offset() later if we decide it is a better solution.

[jvesely]

fix build without D20298 (GRID_OFFSET is still needed)

include/llvm/IR/IntrinsicsAMDGPU.td

   GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,
   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
 
+def int_amdgcn_implicitarg_ptr :
+  GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">,
+  Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
+
 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
 def int_amdgcn_interp_p1 :
   GCCBuiltin<"__builtin_amdgcn_interp_p1">,

lib/Target/AMDGPU/AMDGPUISelLowering.h

                                        unsigned Reg, EVT VT) const;
 
   enum ImplicitParameter {
-    GRID_DIM,
-    GRID_OFFSET
+    FIRST_IMPLICIT,
+    GRID_DIM = FIRST_IMPLICIT,
   };
 
   /// \brief Helper function that returns the byte offset of the given

lib/Target/AMDGPU/SIISelLowering.h

 namespace llvm {
 
 class SITargetLowering final : public AMDGPUTargetLowering {
+  SDValue LowerParameterPtr(SelectionDAG &DAG, SDLoc SL, SDValue Chain,
+                            unsigned Offset) const;
   SDValue LowerParameter(SelectionDAG &DAG, EVT VT, EVT MemVT, SDLoc DL,
                          SDValue Chain, unsigned Offset, bool Signed) const;
 

lib/Target/AMDGPU/SIISelLowering.cpp

   return TargetLowering::isTypeDesirableForOp(Op, VT);
 }
 
-SDValue SITargetLowering::LowerParameter(SelectionDAG &DAG, EVT VT, EVT MemVT,
+SDValue SITargetLowering::LowerParameterPtr(SelectionDAG &DAG,
                                             SDLoc SL, SDValue Chain,
-                                         unsigned Offset, bool Signed) const {
+                                            unsigned Offset) const {
   const DataLayout &DL = DAG.getDataLayout();
   MachineFunction &MF = DAG.getMachineFunction();
   const SIRegisterInfo *TRI =
       static_cast<const SIRegisterInfo*>(Subtarget->getRegisterInfo());
   unsigned InputPtrReg = TRI->getPreloadedValue(MF, SIRegisterInfo::KERNARG_SEGMENT_PTR);
 
-  Type *Ty = VT.getTypeForEVT(*DAG.getContext());
-
   MachineRegisterInfo &MRI = DAG.getMachineFunction().getRegInfo();
   MVT PtrVT = getPointerTy(DL, AMDGPUAS::CONSTANT_ADDRESS);
-  PointerType *PtrTy = PointerType::get(Ty, AMDGPUAS::CONSTANT_ADDRESS);
   SDValue BasePtr = DAG.getCopyFromReg(Chain, SL,
                                        MRI.getLiveInVirtReg(InputPtrReg), PtrVT);
-  SDValue Ptr = DAG.getNode(ISD::ADD, SL, PtrVT, BasePtr,
-                            DAG.getConstant(Offset, SL, PtrVT));
+  return DAG.getNode(ISD::ADD, SL, PtrVT, BasePtr,
+}
+SDValue SITargetLowering::LowerParameter(SelectionDAG &DAG, EVT VT, EVT MemVT,
+                                         SDLoc SL, SDValue Chain,
+                                         unsigned Offset, bool Signed) const {
+  const DataLayout &DL = DAG.getDataLayout();
+  Type *Ty = VT.getTypeForEVT(*DAG.getContext());
+  MVT PtrVT = getPointerTy(DL, AMDGPUAS::CONSTANT_ADDRESS);
+  PointerType *PtrTy = PointerType::get(Ty, AMDGPUAS::CONSTANT_ADDRESS);
   SDValue PtrOffset = DAG.getUNDEF(PtrVT);
   MachinePointerInfo PtrInfo(UndefValue::get(PtrTy));
 

   if (MemVT.isFloatingPoint())
     ExtTy = ISD::EXTLOAD;
 
+  SDValue Ptr = LowerParameterPtr(DAG, SL, Chain, Offset);
   return DAG.getLoad(ISD::UNINDEXED, ExtTy,
                      VT, SL, Chain, Ptr, PtrOffset, PtrInfo, MemVT,
                      false, // isVolatile

     return CreateLiveInRegister(DAG, &AMDGPU::SReg_64RegClass,
                                 TRI->getPreloadedValue(MF, Reg), VT);
   }
+  case Intrinsic::amdgcn_implicitarg_ptr: {
+    unsigned offset = getImplicitParameterOffset(MFI, FIRST_IMPLICIT);
+    return LowerParameterPtr(DAG, DL, DAG.getEntryNode(), offset);
+  }
   case Intrinsic::amdgcn_kernarg_segment_ptr: {
     unsigned Reg
       = TRI->getPreloadedValue(MF, SIRegisterInfo::KERNARG_SEGMENT_PTR);

test/CodeGen/AMDGPU/llvm.amdgcn.kernarg.segment.ptr.ll

   ret void
 }
 
+; ALL-LABEL: {{^}}test_implicit:
+; 10 + 9 (36 prepended implicit bytes) + 2(out pointer) = 21 = 0x15
+; MESA: s_load_dword s{{[0-9]+}}, s[0:1], 0x15
+define void @test_implicit(i32 addrspace(1)* %out) #1 {
+  %implicitarg.ptr = call noalias i8 addrspace(2)* @llvm.amdgcn.implicitarg.ptr()
+  %header.ptr = bitcast i8 addrspace(2)* %implicitarg.ptr to i32 addrspace(2)*
+  %gep = getelementptr i32, i32 addrspace(2)* %header.ptr, i64 10
+  %value = load i32, i32 addrspace(2)* %gep
+  store i32 %value, i32 addrspace(1)* %out
+  ret void
+}
+
 declare i8 addrspace(2)* @llvm.amdgcn.kernarg.segment.ptr() #0
+declare i8 addrspace(2)* @llvm.amdgcn.implicitarg.ptr() #0
 
 attributes #0 = { nounwind readnone }
 attributes #1 = { nounwind }