[PATCH] D120265: AMDGPU: Use the implicit kernargs for code object version 5

[Changpeng Fang via Phabricator via llvm-commits]

cfang created this revision.
cfang added reviewers: b-sumner, tony-tye, arsenm, sameerds, bcahoon, yaxunl, kzhuravl, AMDGPU.
Herald added subscribers: foad, okura, kuter, kerbowa, hiraditya, t-tye, tpr, dstuttard, nhaehnle, jvesely.
cfang requested review of this revision.
Herald added a subscriber: wdng.
Herald added a reviewer: jdoerfert.
Herald added a reviewer: sstefan1.
Herald added a reviewer: baziotis.
Herald added a project: LLVM.

In this work, we implement the use of the implicit kernargs in the backend
and clang under code object version 5.

      

Specifically, for trap handling, for targets that do not support getDoorbellID,
we load the queue_ptr from the implicit kernarg, and move queue_ptr to s[0:1].
To get aperture bases when targets do not have aperture registers, we load
private_base or shared_base directly from the implicit kernarg. In clang, we use
implicitarg_ptr + offsets to implement __builtin_amdgcn_workgroup_size_{xyz}.

Fixes:

  SWDEV-307189


https://reviews.llvm.org/D120265

Files:
  clang/lib/CodeGen/CGBuiltin.cpp
  clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
  clang/test/CodeGenOpenCL/builtins-amdgcn-workgroup-size-v5.cl
  llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
  llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
  llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h
  llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
  llvm/lib/Target/AMDGPU/SIISelLowering.cpp
  llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
  llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
  llvm/test/CodeGen/AMDGPU/GlobalISel/implicit-kernarg-backend-usage-global-isel.ll
  llvm/test/CodeGen/AMDGPU/implicit-kernarg-backend-usage.ll

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D120265.410341.patch
Type: text/x-patch
Size: 60137 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20220221/853a8c12/attachment.bin>