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/2] [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/2] 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
