Author: Joseph Huber Date: 2026-06-17T19:57:54-05:00 New Revision: ced82d8d8b7924077692a46f30fab67be7268073
URL: https://github.com/llvm/llvm-project/commit/ced82d8d8b7924077692a46f30fab67be7268073 DIFF: https://github.com/llvm/llvm-project/commit/ced82d8d8b7924077692a46f30fab67be7268073.diff LOG: [Clang] Make the pointers to gpuintrin AS query const (#204492) Summary: Right now these force a const cast if the user is checking a read-only pointer, not great. Added: Modified: clang/lib/Headers/amdgpuintrin.h clang/lib/Headers/nvptxintrin.h clang/lib/Headers/spirvintrin.h Removed: ################################################################################ diff --git a/clang/lib/Headers/amdgpuintrin.h b/clang/lib/Headers/amdgpuintrin.h index 3f4bba0a03930..dc85b713df67e 100644 --- a/clang/lib/Headers/amdgpuintrin.h +++ b/clang/lib/Headers/amdgpuintrin.h @@ -137,13 +137,13 @@ __gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x, } // Returns true if the flat pointer points to AMDGPU 'shared' memory. -_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(void *ptr) { +_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(const void *ptr) { return __builtin_amdgcn_is_shared((void [[clang::address_space(0)]] *)(( void [[clang::opencl_generic]] *)ptr)); } // Returns true if the flat pointer points to AMDGPU 'private' memory. -_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_private(void *ptr) { +_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_private(const void *ptr) { return __builtin_amdgcn_is_private((void [[clang::address_space(0)]] *)(( void [[clang::opencl_generic]] *)ptr)); } diff --git a/clang/lib/Headers/nvptxintrin.h b/clang/lib/Headers/nvptxintrin.h index df87cf4eaeaaa..40704949d8e4b 100644 --- a/clang/lib/Headers/nvptxintrin.h +++ b/clang/lib/Headers/nvptxintrin.h @@ -182,12 +182,12 @@ __gpu_match_all_u64(uint64_t __lane_mask, uint64_t __x) { #endif // Returns true if the flat pointer points to CUDA 'shared' memory. -_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(void *ptr) { +_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(const void *ptr) { return __nvvm_isspacep_shared(ptr); } // Returns true if the flat pointer points to CUDA 'local' memory. -_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_private(void *ptr) { +_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_private(const void *ptr) { return __nvvm_isspacep_local(ptr); } diff --git a/clang/lib/Headers/spirvintrin.h b/clang/lib/Headers/spirvintrin.h index a172bf8ae3ac8..0c61a3f9cc8e3 100644 --- a/clang/lib/Headers/spirvintrin.h +++ b/clang/lib/Headers/spirvintrin.h @@ -144,12 +144,12 @@ __gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x, } // SPIR-V does not expose this, always return false. -_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(void *ptr) { +_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(const void *ptr) { return 0; } // SPIR-V does not expose this, always return false. -_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_private(void *ptr) { +_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_private(const void *ptr) { return 0; } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
