https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/230677
>From 3c00b39f69f74c2a6f2cf7aed1ac8b17701b886c Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Fri, 9 Oct 2026 16:30:22 -0500 Subject: [PATCH 1/3] [AMDGPU] Allow the AS identifying builtins to take flat pointers Summary: We have `__builtin_amdgcn_is_shared` and `__builtin_amdgcn_is_private` which are used to determine the underlying address space of a flat pointer. There's some magic around AS(0) depending on the language mode, in OpenCL it's `__generic`, in HIP it's the default AS. In the mode without a language it's still a distinct AS. This PR changes the builtin signature to just be a flat pointer. This should be equivalent on the targets I know of, but makes the expectation clearer and fixes the type issues w/ plain C/C++. Test file hopefully verifies that existing OpenCL / HIP usage is unaffected. --- clang/include/clang/Basic/BuiltinsAMDGPU.td | 4 +-- clang/lib/Headers/amdgpuintrin.h | 10 +++---- clang/lib/Sema/SemaExpr.cpp | 2 +- clang/test/Sema/amdgcn-is-shared-private.c | 32 +++++++++++++++++++++ 4 files changed, 39 insertions(+), 9 deletions(-) create mode 100644 clang/test/Sema/amdgcn-is-shared-private.c diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 483df0561fc20..5e7678593e2eb 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -358,8 +358,8 @@ def __builtin_amdgcn_fcmpf : AMDGPUBuiltin<"uint64_t(float, float, _Constant int //===----------------------------------------------------------------------===// // Flat addressing builtins. //===----------------------------------------------------------------------===// -def __builtin_amdgcn_is_shared : AMDGPUBuiltin<"bool(void const address_space<0> *)", [Const]>; -def __builtin_amdgcn_is_private : AMDGPUBuiltin<"bool(void const address_space<0> *)", [Const]>; +def __builtin_amdgcn_is_shared : AMDGPUBuiltin<"bool(void const *)", [Const]>; +def __builtin_amdgcn_is_private : AMDGPUBuiltin<"bool(void const *)", [Const]>; //===----------------------------------------------------------------------===// // GWS builtins. diff --git a/clang/lib/Headers/amdgpuintrin.h b/clang/lib/Headers/amdgpuintrin.h index dc85b713df67e..eb9c95281075c 100644 --- a/clang/lib/Headers/amdgpuintrin.h +++ b/clang/lib/Headers/amdgpuintrin.h @@ -136,16 +136,14 @@ __gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x, return __builtin_amdgcn_ds_bpermute(__lane << 2, __x); } -// Returns true if the flat pointer points to AMDGPU 'shared' memory. +// Returns true if the pointer points to AMDGPU 'shared' memory. _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)); + return __builtin_amdgcn_is_shared(ptr); } -// Returns true if the flat pointer points to AMDGPU 'private' memory. +// Returns true if the pointer points to AMDGPU 'private' memory. _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)); + return __builtin_amdgcn_is_private(ptr); } // Terminates execution of the associated wavefront. diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index ec74c3432b5cf..d152db5882f7f 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -7090,7 +7090,7 @@ ExprResult Sema::BuildCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc, NDecl = FDecl; Fn = DeclRefExpr::Create( Context, DRE->getQualifierLoc(), SourceLocation(), FDecl, false, - SourceLocation(), Fn->getType() /* BuiltinFnTy */, + DRE->getLocation(), Fn->getType() /* BuiltinFnTy */, Fn->getValueKind(), FDecl, nullptr, DRE->isNonOdrUse()); } } diff --git a/clang/test/Sema/amdgcn-is-shared-private.c b/clang/test/Sema/amdgcn-is-shared-private.c new file mode 100644 index 0000000000000..b70e1baaed8e8 --- /dev/null +++ b/clang/test/Sema/amdgcn-is-shared-private.c @@ -0,0 +1,32 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -verify -x c %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -verify -x c++ %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL1.2 -fsyntax-only -verify -x cl %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -verify -x cl %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL3.0 -fsyntax-only -verify -x cl %s + +// expected-no-diagnostics + +#define AS(N) __attribute__((address_space(N))) +#define OCL(Name) __attribute__((opencl_##Name)) +#define CHECK(p) \ + (void)__builtin_amdgcn_is_shared(p); \ + (void)__builtin_amdgcn_is_private(p) + +#if defined(__HIP__) +__attribute__((device)) +#endif +void test(void *flat, void AS(0) *as0, void AS(1) *as1, void AS(3) *as3, + void AS(5) *as5, void OCL(global) *g, void OCL(local) *l, + void OCL(private) *p, void OCL(constant) *c, void OCL(generic) *gen) { + CHECK(flat); + CHECK(as0); + CHECK(as1); + CHECK(as3); + CHECK(as5); + CHECK(g); + CHECK(l); + CHECK(p); + CHECK(c); + CHECK(gen); +} >From c926cd9cc8a9d7b47f651b7258fbb28a1f13c3a5 Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Sun, 11 Oct 2026 15:09:56 -0500 Subject: [PATCH 2/3] Update amdgcn-is-shared-private.c --- clang/test/Sema/amdgcn-is-shared-private.c | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/test/Sema/amdgcn-is-shared-private.c b/clang/test/Sema/amdgcn-is-shared-private.c index b70e1baaed8e8..3db6e3ffd10d0 100644 --- a/clang/test/Sema/amdgcn-is-shared-private.c +++ b/clang/test/Sema/amdgcn-is-shared-private.c @@ -1,9 +1,9 @@ -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -verify -x c %s -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -verify -x c++ %s -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip -fsyntax-only -verify %s -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL1.2 -fsyntax-only -verify -x cl %s -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -verify -x cl %s -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL3.0 -fsyntax-only -verify -x cl %s +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fsyntax-only -verify -x c %s +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fsyntax-only -verify -x c++ %s +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -cl-std=CL1.2 -fsyntax-only -verify -x cl %s +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -verify -x cl %s +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -cl-std=CL3.0 -fsyntax-only -verify -x cl %s // expected-no-diagnostics >From 9a2007f17537e73caf2d31f0561ad033d20f3c8e Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Sun, 11 Oct 2026 17:34:38 -0500 Subject: [PATCH 3/3] tests --- clang/test/Sema/amdgcn-is-shared-private.c | 65 +++++++++++++--------- 1 file changed, 40 insertions(+), 25 deletions(-) diff --git a/clang/test/Sema/amdgcn-is-shared-private.c b/clang/test/Sema/amdgcn-is-shared-private.c index 3db6e3ffd10d0..479731de36d47 100644 --- a/clang/test/Sema/amdgcn-is-shared-private.c +++ b/clang/test/Sema/amdgcn-is-shared-private.c @@ -1,32 +1,47 @@ -// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fsyntax-only -verify -x c %s -// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fsyntax-only -verify -x c++ %s -// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip -fsyntax-only -verify %s -// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -cl-std=CL1.2 -fsyntax-only -verify -x cl %s -// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -verify -x cl %s -// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -cl-std=CL3.0 -fsyntax-only -verify -x cl %s +// RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -fsyntax-only -verify -x c %s +// RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -fsyntax-only -verify -x c++ %s +// RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -fcuda-is-device -x hip \ +// RUN: -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -cl-std=CL1.2 -fsyntax-only \ +// RUN: -verify -x cl %s +// RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -cl-std=CL2.0 -fsyntax-only \ +// RUN: -verify -x cl %s +// RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -cl-std=CL3.0 -fsyntax-only \ +// RUN: -verify -x cl %s // expected-no-diagnostics -#define AS(N) __attribute__((address_space(N))) -#define OCL(Name) __attribute__((opencl_##Name)) -#define CHECK(p) \ - (void)__builtin_amdgcn_is_shared(p); \ - (void)__builtin_amdgcn_is_private(p) - #if defined(__HIP__) __attribute__((device)) #endif -void test(void *flat, void AS(0) *as0, void AS(1) *as1, void AS(3) *as3, - void AS(5) *as5, void OCL(global) *g, void OCL(local) *l, - void OCL(private) *p, void OCL(constant) *c, void OCL(generic) *gen) { - CHECK(flat); - CHECK(as0); - CHECK(as1); - CHECK(as3); - CHECK(as5); - CHECK(g); - CHECK(l); - CHECK(p); - CHECK(c); - CHECK(gen); +void test(void *flat_ptr, + void __attribute__((address_space(0))) *flat_addrspace, + void __attribute__((address_space(1))) *global_addrspace, + void __attribute__((address_space(3))) *local_addrspace, + void __attribute__((address_space(5))) *private_addrspace, + void __attribute__((opencl_global)) *global_ptr, + void __attribute__((opencl_local)) *local_ptr, + void __attribute__((opencl_private)) *private_ptr, + void __attribute__((opencl_constant)) *constant_ptr, + void __attribute__((opencl_generic)) *generic_ptr) { + (void)__builtin_amdgcn_is_shared(flat_ptr); + (void)__builtin_amdgcn_is_private(flat_ptr); + (void)__builtin_amdgcn_is_shared(flat_addrspace); + (void)__builtin_amdgcn_is_private(flat_addrspace); + (void)__builtin_amdgcn_is_shared(global_addrspace); + (void)__builtin_amdgcn_is_private(global_addrspace); + (void)__builtin_amdgcn_is_shared(local_addrspace); + (void)__builtin_amdgcn_is_private(local_addrspace); + (void)__builtin_amdgcn_is_shared(private_addrspace); + (void)__builtin_amdgcn_is_private(private_addrspace); + (void)__builtin_amdgcn_is_shared(global_ptr); + (void)__builtin_amdgcn_is_private(global_ptr); + (void)__builtin_amdgcn_is_shared(local_ptr); + (void)__builtin_amdgcn_is_private(local_ptr); + (void)__builtin_amdgcn_is_shared(private_ptr); + (void)__builtin_amdgcn_is_private(private_ptr); + (void)__builtin_amdgcn_is_shared(constant_ptr); + (void)__builtin_amdgcn_is_private(constant_ptr); + (void)__builtin_amdgcn_is_shared(generic_ptr); + (void)__builtin_amdgcn_is_private(generic_ptr); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
