[all-commits] [llvm/llvm-project] dd5895: AMDGPU: Use the implicit kernargs for code object ...

[Changpeng Fang via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: dd5895cc39864393f8ca357bc4e23e8d7b5b9723
      https://github.com/llvm/llvm-project/commit/dd5895cc39864393f8ca357bc4e23e8d7b5b9723
  Author: Changpeng Fang <Changpeng.Fang at amd.com>
  Date:   2022-03-17 (Thu, 17 Mar 2022)

  Changed paths:
    M clang/lib/CodeGen/CGBuiltin.cpp
    M clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
    M llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
    M llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
    M llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h
    M llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
    M llvm/lib/Target/AMDGPU/SIDefines.h
    M llvm/lib/Target/AMDGPU/SIISelLowering.cpp
    M llvm/lib/Target/AMDGPU/SIISelLowering.h
    M llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
    M llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
    A llvm/test/CodeGen/AMDGPU/GlobalISel/implicit-kernarg-backend-usage-global-isel.ll
    A llvm/test/CodeGen/AMDGPU/implicit-kernarg-backend-usage.ll

  Log Message:
  -----------
  AMDGPU: Use the implicit kernargs for code object version 5

Summary:
  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}.

Reviewers: arsenm, sameerds, yaxunl

Differential Revision: https://reviews.llvm.org/D120265