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

[Matt Arsenault via Phabricator via llvm-commits]

arsenm added inline comments.


================
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]>;
----------------
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


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp:4207
+    B.buildCopy(DstReg,
+                B.buildConstant(LLT::scalar(32), KnownSize->getZExtValue()));
+  }
----------------
Can build the constant directly into DstReg


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp:106
+ConstantInt *AMDGPUMachineFunction::getLDSKernelIdMetadata(const Function &F) {
+  auto MD = F.getMetadata("llvm.amdgcn.lds_kernel_id");
+  if (MD && MD->getNumOperands() == 1) {
----------------
Why '.' and '_'?


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp:739
+                          2 + // Implicit buffer ptr
+                          1;  // LDSKernelId
+
----------------
This doesn't apply since it's not a real user SGPR?


================
Comment at: llvm/lib/Target/AMDGPU/SIISelLowering.cpp:1674
+    return DAG.getConstant(KnownSize->getZExtValue(), SL, MVT::i32);
+  } else {
+    return SDValue();
----------------
No else before return


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D125060