[clang] [Clang] Simplify target specific implementations in gpuintrin.h (PR #194669)
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Tue Apr 28 09:24:24 PDT 2026
https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/194669
Summary:
Previously we had a three-way dance where the top-level defined an impl
then the lower level defined potentially an override and then we
stitched those together.
This simplifies that, instead just defining a macro that says whether
or not we need it. This works because the target-specific portions are
always included first, so this basically says "Do we need to make a
default version".
>From 86b1f6d114b0a6965aab61644ea5b137148d92d4 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Tue, 28 Apr 2026 11:00:32 -0500
Subject: [PATCH] [Clang] Simplify target specific implementations in
gpuintrin.h
Summary:
Previously we had a three-way dance where the top-level defined an impl
then the lower level defined potentially an override and then we
stitched those together.
This simplifies that, instead just defining a macro that says whether
or not we need it. This works because the target-specific portions are
always included first, so this basically says "Do we need to make a
default version".
---
clang/lib/Headers/amdgpuintrin.h | 24 -----------------
clang/lib/Headers/gpuintrin.h | 44 ++++++++++++--------------------
clang/lib/Headers/nvptxintrin.h | 32 +++++++++--------------
clang/lib/Headers/spirvintrin.h | 24 -----------------
4 files changed, 28 insertions(+), 96 deletions(-)
diff --git a/clang/lib/Headers/amdgpuintrin.h b/clang/lib/Headers/amdgpuintrin.h
index e0989e0a2d097..3f4bba0a03930 100644
--- a/clang/lib/Headers/amdgpuintrin.h
+++ b/clang/lib/Headers/amdgpuintrin.h
@@ -136,30 +136,6 @@ __gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x,
return __builtin_amdgcn_ds_bpermute(__lane << 2, __x);
}
-// Returns a bitmask marking all lanes that have the same value of __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_any_u32(uint64_t __lane_mask, uint32_t __x) {
- return __gpu_match_any_u32_impl(__lane_mask, __x);
-}
-
-// Returns a bitmask marking all lanes that have the same value of __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_any_u64(uint64_t __lane_mask, uint64_t __x) {
- return __gpu_match_any_u64_impl(__lane_mask, __x);
-}
-
-// Returns the current lane mask if every lane contains __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_all_u32(uint64_t __lane_mask, uint32_t __x) {
- return __gpu_match_all_u32_impl(__lane_mask, __x);
-}
-
-// Returns the current lane mask if every lane contains __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_all_u64(uint64_t __lane_mask, uint64_t __x) {
- return __gpu_match_all_u64_impl(__lane_mask, __x);
-}
-
// Returns true if the flat pointer points to AMDGPU 'shared' memory.
_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(void *ptr) {
return __builtin_amdgcn_is_shared((void [[clang::address_space(0)]] *)((
diff --git a/clang/lib/Headers/gpuintrin.h b/clang/lib/Headers/gpuintrin.h
index 253a426730791..9716541dc4b5d 100644
--- a/clang/lib/Headers/gpuintrin.h
+++ b/clang/lib/Headers/gpuintrin.h
@@ -32,30 +32,6 @@ _Pragma("push_macro(\"bool\")");
#define bool _Bool
#endif
-_Pragma("omp begin declare target device_type(nohost)");
-_Pragma("omp begin declare variant match(device = {kind(gpu)})");
-
-// Forward declare a few functions for the implementation header.
-
-// Returns a bitmask marking all lanes that have the same value of __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_any_u32_impl(uint64_t __lane_mask, uint32_t __x);
-
-// Returns a bitmask marking all lanes that have the same value of __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_any_u64_impl(uint64_t __lane_mask, uint64_t __x);
-
-// Returns the current lane mask if every lane contains __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_all_u32_impl(uint64_t __lane_mask, uint32_t __x);
-
-// Returns the current lane mask if every lane contains __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_all_u64_impl(uint64_t __lane_mask, uint64_t __x);
-
-_Pragma("omp end declare variant");
-_Pragma("omp end declare target");
-
#if defined(__NVPTX__)
#include <nvptxintrin.h>
#elif defined(__AMDGPU__)
@@ -289,8 +265,9 @@ __DO_LANE_OPS(double, __GPU_OP, -__builtin_inf(), maxnum, f64);
#undef __DO_LANE_OPS
// Returns a bitmask marking all lanes that have the same value of __x.
+#ifndef __gpu_match_any_u32_impl
_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_any_u32_impl(uint64_t __lane_mask, uint32_t __x) {
+__gpu_match_any_u32(uint64_t __lane_mask, uint32_t __x) {
uint64_t __match_mask = 0;
bool __done = 0;
@@ -308,10 +285,13 @@ __gpu_match_any_u32_impl(uint64_t __lane_mask, uint32_t __x) {
}
return __match_mask;
}
+#endif
+#undef __gpu_match_any_u32_impl
// Returns a bitmask marking all lanes that have the same value of __x.
+#ifndef __gpu_match_any_u64_impl
_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_any_u64_impl(uint64_t __lane_mask, uint64_t __x) {
+__gpu_match_any_u64(uint64_t __lane_mask, uint64_t __x) {
uint64_t __match_mask = 0;
bool __done = 0;
@@ -329,24 +309,32 @@ __gpu_match_any_u64_impl(uint64_t __lane_mask, uint64_t __x) {
}
return __match_mask;
}
+#endif
+#undef __gpu_match_any_u64_impl
// Returns the current lane mask if every lane contains __x.
+#ifndef __gpu_match_all_u32_impl
_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_all_u32_impl(uint64_t __lane_mask, uint32_t __x) {
+__gpu_match_all_u32(uint64_t __lane_mask, uint32_t __x) {
uint32_t __first = __gpu_shuffle_idx_u32(
__lane_mask, __builtin_ctzg(__lane_mask), __x, __gpu_num_lanes());
uint64_t __ballot = __gpu_ballot(__lane_mask, __x == __first);
return __ballot == __lane_mask ? __lane_mask : UINT64_C(0);
}
+#endif
+#undef __gpu_match_all_u32_impl
// Returns the current lane mask if every lane contains __x.
+#ifndef __gpu_match_all_u64_impl
_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_all_u64_impl(uint64_t __lane_mask, uint64_t __x) {
+__gpu_match_all_u64(uint64_t __lane_mask, uint64_t __x) {
uint64_t __first = __gpu_shuffle_idx_u64(
__lane_mask, __builtin_ctzg(__lane_mask), __x, __gpu_num_lanes());
uint64_t __ballot = __gpu_ballot(__lane_mask, __x == __first);
return __ballot == __lane_mask ? __lane_mask : UINT64_C(0);
}
+#endif
+#undef __gpu_match_all_u64_impl
_Pragma("omp end declare variant");
_Pragma("omp end declare target");
diff --git a/clang/lib/Headers/nvptxintrin.h b/clang/lib/Headers/nvptxintrin.h
index 57a6a2cd08633..df87cf4eaeaaa 100644
--- a/clang/lib/Headers/nvptxintrin.h
+++ b/clang/lib/Headers/nvptxintrin.h
@@ -144,50 +144,42 @@ __gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x,
}
// Returns a bitmask marking all lanes that have the same value of __x.
+#if __CUDA_ARCH__ >= 700
+#define __gpu_match_any_u32_impl
_DEFAULT_FN_ATTRS static __inline__ uint64_t
__gpu_match_any_u32(uint64_t __lane_mask, uint32_t __x) {
- // Newer targets can use the dedicated CUDA support.
-#if __CUDA_ARCH__ >= 700
return __nvvm_match_any_sync_i32(__lane_mask, __x);
-#else
- return __gpu_match_any_u32_impl(__lane_mask, __x);
-#endif
}
+#endif
// Returns a bitmask marking all lanes that have the same value of __x.
+#if __CUDA_ARCH__ >= 700
+#define __gpu_match_any_u64_impl
_DEFAULT_FN_ATTRS static __inline__ uint64_t
__gpu_match_any_u64(uint64_t __lane_mask, uint64_t __x) {
- // Newer targets can use the dedicated CUDA support.
-#if __CUDA_ARCH__ >= 700
return __nvvm_match_any_sync_i64(__lane_mask, __x);
-#else
- return __gpu_match_any_u64_impl(__lane_mask, __x);
-#endif
}
+#endif
// Returns the current lane mask if every lane contains __x.
+#if __CUDA_ARCH__ >= 700
+#define __gpu_match_all_u32_impl
_DEFAULT_FN_ATTRS static __inline__ uint64_t
__gpu_match_all_u32(uint64_t __lane_mask, uint32_t __x) {
- // Newer targets can use the dedicated CUDA support.
-#if __CUDA_ARCH__ >= 700
int predicate;
return __nvvm_match_all_sync_i32p(__lane_mask, __x, &predicate);
-#else
- return __gpu_match_all_u32_impl(__lane_mask, __x);
-#endif
}
+#endif
// Returns the current lane mask if every lane contains __x.
+#if __CUDA_ARCH__ >= 700
+#define __gpu_match_all_u64_impl
_DEFAULT_FN_ATTRS static __inline__ uint64_t
__gpu_match_all_u64(uint64_t __lane_mask, uint64_t __x) {
- // Newer targets can use the dedicated CUDA support.
-#if __CUDA_ARCH__ >= 700
int predicate;
return __nvvm_match_all_sync_i64p(__lane_mask, __x, &predicate);
-#else
- return __gpu_match_all_u64_impl(__lane_mask, __x);
-#endif
}
+#endif
// Returns true if the flat pointer points to CUDA 'shared' memory.
_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(void *ptr) {
diff --git a/clang/lib/Headers/spirvintrin.h b/clang/lib/Headers/spirvintrin.h
index 9658f280b247d..a172bf8ae3ac8 100644
--- a/clang/lib/Headers/spirvintrin.h
+++ b/clang/lib/Headers/spirvintrin.h
@@ -143,30 +143,6 @@ __gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x,
return __builtin_spirv_subgroup_shuffle(__x, __lane);
}
-// Returns a bitmask marking all lanes that have the same value of __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_any_u32(uint64_t __lane_mask, uint32_t __x) {
- return __gpu_match_any_u32_impl(__lane_mask, __x);
-}
-
-// Returns a bitmask marking all lanes that have the same value of __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_any_u64(uint64_t __lane_mask, uint64_t __x) {
- return __gpu_match_any_u64_impl(__lane_mask, __x);
-}
-
-// Returns the current lane mask if every lane contains __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_all_u32(uint64_t __lane_mask, uint32_t __x) {
- return __gpu_match_all_u32_impl(__lane_mask, __x);
-}
-
-// Returns the current lane mask if every lane contains __x.
-_DEFAULT_FN_ATTRS static __inline__ uint64_t
-__gpu_match_all_u64(uint64_t __lane_mask, uint64_t __x) {
- return __gpu_match_all_u64_impl(__lane_mask, __x);
-}
-
// SPIR-V does not expose this, always return false.
_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(void *ptr) {
return 0;
More information about the cfe-commits
mailing list