[PATCH] D125060: [amdgpu] Implement lds kernel id intrinsic

[Jon Chesterfield via Phabricator via llvm-commits]

JonChesterfield marked 5 inline comments as done.
JonChesterfield added a comment.

Fixes applied. This whole patch is the machinery to tag kernels with a uint32_t at compile time then access it from a function at runtime, using the existing machinery for implicit SGPR arguments.

> What is this metadata for? Where does it come from?

The following patch (which is functional but not review ready) annotates each kernel that uses LDS with an arbitrary unique integer stored in this metadata. As of this patch, that's only in test cases.

Said patch also creates an array of LDS addresses in __constant memory and initializes it to match the unique integers stored in said metadata.

Finally the intrinsic introduced here is used to look up which set of addresses corresponds to the kernel that is running the current function in order to find where a given variable is allocated.



================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:171
+def int_amdgcn_lds_kernel_id :
+  ClangBuiltin<"__builtin_amdgcn_lds_kernel_id">,
+  Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>;
----------------
arsenm wrote:
> This absolutely should not get a ClangBuiltin. Also probably shouldn't include "lds" in the name llvm.amdgcn.compiler.kernel.id? Probably should add a comment that this is for internal compiler uses and users should not rely on it
It should have lds in the name because it's currently only used by lds. Specifically I don't want to enumerate kernels that don't need this indirection as that would increase the size of a lookup table.

If this machinery for identifying a kernel at runtime proves to be more generally useful, we can rename these variables and move the enumeration pass out of lowermodulelds when that use case arises.


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp:739
+                          2 + // Implicit buffer ptr
+                          1;  // LDSKernelId
+
----------------
arsenm wrote:
> This doesn't apply since it's not a real user SGPR?
It's not a system sgpr either. I'll name it differently but it's still a preloaded sgpr in the sense that it's an sgpr that gets initialised by mostly hidden magic when necessary.

The alternative seems to be duplicating the current machinery for passing magic hidden registers across function calls and I'd really prefer to hook into the existing subsystem.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D125060/new/

https://reviews.llvm.org/D125060