[libclc] e67360e - libclc: Implement address space qualifier functions for amdgpu (#184766)

[via cfe-commits]

Author: Matt Arsenault
Date: 2026-03-05T12:30:09+01:00
New Revision: e67360ec319a642e56abb0c3f0c3e627a98dd225

URL: https://github.com/llvm/llvm-project/commit/e67360ec319a642e56abb0c3f0c3e627a98dd225
DIFF: https://github.com/llvm/llvm-project/commit/e67360ec319a642e56abb0c3f0c3e627a98dd225.diff

LOG: libclc: Implement address space qualifier functions for amdgpu (#184766)

Added: 
    libclc/clc/include/clc/address_space/qualifier.h
    libclc/clc/lib/amdgcn/address_space/qualifier.cl
    libclc/clc/lib/generic/shared/clc_qualifier.cl
    libclc/opencl/lib/generic/address_space/qualifier.cl

Modified: 
    libclc/clc/lib/amdgcn/SOURCES
    libclc/clc/lib/generic/SOURCES
    libclc/opencl/lib/generic/SOURCES

Removed: 
    


################################################################################
diff  --git a/libclc/clc/include/clc/address_space/qualifier.h b/libclc/clc/include/clc/address_space/qualifier.h
new file mode 100644
index 0000000000000..28b8ab33d5df8
--- /dev/null
+++ b/libclc/clc/include/clc/address_space/qualifier.h
@@ -0,0 +1,33 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_ADDRESS_SPACE_CLC_QUALIFIER_H__
+#define __CLC_ADDRESS_SPACE_CLC_QUALIFIER_H__
+
+#include <clc/clcfunc.h>
+
+#if _CLC_GENERIC_AS_SUPPORTED
+
+_CLC_OVERLOAD _CLC_DECL _CLC_CONST cl_mem_fence_flags __clc_get_fence(void *p);
+_CLC_OVERLOAD _CLC_DECL _CLC_CONST cl_mem_fence_flags
+__clc_get_fence(const void *p);
+
+_CLC_OVERLOAD _CLC_DECL _CLC_CONST __global void *__clc_to_global(void *p);
+_CLC_OVERLOAD _CLC_DECL _CLC_CONST const __global void *
+__clc_to_global(const void *p);
+
+_CLC_OVERLOAD _CLC_DECL _CLC_CONST __local void *__clc_to_local(void *p);
+_CLC_OVERLOAD _CLC_DECL _CLC_CONST const __local void *
+__clc_to_local(const void *p);
+
+_CLC_OVERLOAD _CLC_DECL _CLC_CONST __private void *__clc_to_private(void *p);
+_CLC_OVERLOAD _CLC_DECL _CLC_CONST const __private void *
+__clc_to_private(const void *p);
+#endif // _CLC_GENERIC_AS_SUPPORTED
+
+#endif // __CLC_ADDRESS_SPACE_CLC_QUALIFIER_H__

diff  --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES
index 5040bf1158674..b4557b0a26f70 100644
--- a/libclc/clc/lib/amdgcn/SOURCES
+++ b/libclc/clc/lib/amdgcn/SOURCES
@@ -1,3 +1,4 @@
+address_space/qualifier.cl
 math/clc_ldexp.cl
 mem_fence/clc_mem_fence.cl
 synchronization/clc_work_group_barrier.cl

diff  --git a/libclc/clc/lib/amdgcn/address_space/qualifier.cl b/libclc/clc/lib/amdgcn/address_space/qualifier.cl
new file mode 100644
index 0000000000000..82d9f5351c446
--- /dev/null
+++ b/libclc/clc/lib/amdgcn/address_space/qualifier.cl
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/address_space/qualifier.h>
+
+#if _CLC_GENERIC_AS_SUPPORTED
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST cl_mem_fence_flags
+__clc_get_fence(const __generic void *p) {
+  return __builtin_amdgcn_is_shared(p) ? CLK_LOCAL_MEM_FENCE
+                                       : CLK_GLOBAL_MEM_FENCE;
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST const __global void *
+__clc_to_global(const __generic void *p) {
+  return __builtin_amdgcn_is_private(p) || __builtin_amdgcn_is_shared(p)
+             ? NULL
+             : (const __global void *)p;
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST const __local void *
+__clc_to_local(const __generic void *p) {
+  return __builtin_amdgcn_is_shared(p) ? (__local void *)p : NULL;
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST const __private void *
+__clc_to_private(const __generic void *p) {
+  return __builtin_amdgcn_is_private(p) ? (__private void *)p : NULL;
+}
+
+#endif

diff  --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES
index ac820239baa17..421eb638720e3 100644
--- a/libclc/clc/lib/generic/SOURCES
+++ b/libclc/clc/lib/generic/SOURCES
@@ -173,6 +173,7 @@ relational/clc_signbit.cl
 shared/clc_clamp.cl
 shared/clc_max.cl
 shared/clc_min.cl
+shared/clc_qualifier.cl
 shared/clc_vload.cl
 shared/clc_vstore.cl
 workitem/clc_get_local_linear_id.cl

diff  --git a/libclc/clc/lib/generic/shared/clc_qualifier.cl b/libclc/clc/lib/generic/shared/clc_qualifier.cl
new file mode 100644
index 0000000000000..978f6677798cd
--- /dev/null
+++ b/libclc/clc/lib/generic/shared/clc_qualifier.cl
@@ -0,0 +1,33 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/address_space/qualifier.h>
+
+#if _CLC_GENERIC_AS_SUPPORTED
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST cl_mem_fence_flags
+__clc_get_fence(__generic void *p) {
+  return __clc_get_fence((const __generic void *)p);
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST __global void *
+__clc_to_global(__generic void *p) {
+  return (__global void *)__clc_to_global((const __generic void *)p);
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST __local void *
+__clc_to_local(__generic void *p) {
+  return (__local void *)__clc_to_local((const __generic void *)p);
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST __private void *
+__clc_to_private(__generic void *p) {
+  return (__private void *)__clc_to_private((const __generic void *)p);
+}
+
+#endif

diff  --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index c820c6c3c0c06..bb5e8ab08a711 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -1,3 +1,4 @@
+address_space/qualifier.cl
 subnormal_config.cl
 async/async_work_group_copy.cl
 async/async_work_group_strided_copy.cl

diff  --git a/libclc/opencl/lib/generic/address_space/qualifier.cl b/libclc/opencl/lib/generic/address_space/qualifier.cl
new file mode 100644
index 0000000000000..8ad4a373465bd
--- /dev/null
+++ b/libclc/opencl/lib/generic/address_space/qualifier.cl
@@ -0,0 +1,35 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "clc/address_space/qualifier.h"
+
+#if _CLC_GENERIC_AS_SUPPORTED
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST cl_mem_fence_flags
+get_fence(__generic void *p) {
+  return __clc_get_fence(p);
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST cl_mem_fence_flags
+get_fence(const __generic void *p) {
+  return __clc_get_fence(p);
+}
+
+_CLC_DEF _CLC_CONST __global void *__to_global(__generic void *p) {
+  return __clc_to_global(p);
+}
+
+_CLC_DEF _CLC_CONST __local void *__to_local(__generic void *p) {
+  return __clc_to_local(p);
+}
+
+_CLC_DEF _CLC_CONST __private void *__to_private(__generic void *p) {
+  return __clc_to_private(p);
+}
+
+#endif // _CLC_GENERIC_AS_SUPPORTED