[clang] ced82d8 - [Clang] Make the pointers to gpuintrin AS query const (#204492)

[via cfe-commits]

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;
 }