AMDGPU: Remove nocapture attribute from is.shared and is.private intrinsics

[arsenm]

This should be replaced with captures(address), but tablegen currently has
no way to indicate that on an intrinsic. I opened issue #129184 to fix this.

[arsenm]

• AMDGPU: Remove nocapture attribute from is.shared and is.private intrinsics #129238 👈 (View in Graphite)
• main

This stack of pull requests is managed by Graphite. Learn more about stacking.

[llvmbot]

@llvm/pr-subscribers-llvm-ir

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

Changes

This should be replaced with captures(address), but tablegen currently has
no way to indicate that on an intrinsic. I opened issue #129184 to fix this.

---

Full diff: https://github.com/llvm/llvm-project/pull/129238.diff

1 Files Affected:

• (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+2-2)

diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 876a6f816ad3f..abd0afd413316 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2529,13 +2529,13 @@ def int_amdgcn_set_inactive_chain_arg :
 // Return if the given flat pointer points to a local memory address.
 def int_amdgcn_is_shared : ClangBuiltin<"__builtin_amdgcn_is_shared">,
   DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_ptr_ty],
-  [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>]
+  [IntrNoMem, IntrSpeculatable] // FIXME: This should be capture(caddress)
 >;
 
 // Return if the given flat pointer points to a prvate memory address.
 def int_amdgcn_is_private : ClangBuiltin<"__builtin_amdgcn_is_private">,
   DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_ptr_ty],
-  [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>]
+  [IntrNoMem, IntrSpeculatable] // FIXME: This should be capture(caddress)
 >;
 
 // A uniform tail call to a function with the `amdgpu_cs_chain` or

[jayfoad]

Looks reasonable