[libclc] libclc: Implement address space qualifier functions for amdgpu (PR #184766)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 5 02:31:11 PST 2026
https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/184766
None
>From 7d6936192280d4596de2f0640368453730ba3139 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 5 Mar 2026 09:48:18 +0100
Subject: [PATCH] libclc: Implement address space qualifier functions for
amdgpu
---
.../clc/include/clc/address_space/qualifier.h | 33 +++++++++++++++++
libclc/clc/lib/amdgcn/SOURCES | 1 +
.../clc/lib/amdgcn/address_space/qualifier.cl | 36 +++++++++++++++++++
libclc/clc/lib/generic/SOURCES | 1 +
.../clc/lib/generic/shared/clc_qualifier.cl | 33 +++++++++++++++++
libclc/opencl/lib/generic/SOURCES | 1 +
.../lib/generic/address_space/qualifier.cl | 35 ++++++++++++++++++
7 files changed, 140 insertions(+)
create mode 100644 libclc/clc/include/clc/address_space/qualifier.h
create mode 100644 libclc/clc/lib/amdgcn/address_space/qualifier.cl
create mode 100644 libclc/clc/lib/generic/shared/clc_qualifier.cl
create mode 100644 libclc/opencl/lib/generic/address_space/qualifier.cl
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
More information about the cfe-commits
mailing list