Skip to content

Commit 00f5763

Browse files
authored
AMDGPU: Remove nocapture attribute from is.shared and is.private intrinsics (#129238)
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.
1 parent 2477f82 commit 00f5763

File tree

1 file changed

+2
-2
lines changed

1 file changed

+2
-2
lines changed

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2529,13 +2529,13 @@ def int_amdgcn_set_inactive_chain_arg :
25292529
// Return if the given flat pointer points to a local memory address.
25302530
def int_amdgcn_is_shared : ClangBuiltin<"__builtin_amdgcn_is_shared">,
25312531
DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_ptr_ty],
2532-
[IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>]
2532+
[IntrNoMem, IntrSpeculatable] // FIXME: This should be captures(ret: address)
25332533
>;
25342534

25352535
// Return if the given flat pointer points to a prvate memory address.
25362536
def int_amdgcn_is_private : ClangBuiltin<"__builtin_amdgcn_is_private">,
25372537
DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_ptr_ty],
2538-
[IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>]
2538+
[IntrNoMem, IntrSpeculatable] // FIXME: This should be captures(ret: address)
25392539
>;
25402540

25412541
// A uniform tail call to a function with the `amdgpu_cs_chain` or

0 commit comments

Comments
 (0)