Skip to content

[SYCL] Add SPIR-V atomic builtins declarations to clang/lib/Sema/SPIRVBuiltins.td #17471

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 11 commits into from
Apr 21, 2025
Merged

[SYCL] Add SPIR-V atomic builtins declarations to clang/lib/Sema/SPIRVBuiltins.td #17471

Changes from 1 commit
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
8088b0a
[SYCL] Move SPIR-V atomic builtins declarations into clang/lib/Sema/S…
wenju-he Mar 14, 2025
0e45a71
remove const from pointer argument, add missing ##OP_MANGLED
wenju-he Mar 17, 2025
6f69afe
add back atomic declarations in spirv_ops.hpp
wenju-he Mar 17, 2025
ad33aa8
update test amdgpu_unsafe_atomics.cpp: remove volatile
wenju-he Mar 17, 2025
e630132
add clang test
wenju-he Mar 20, 2025
105d495
Merge branch 'sycl' into move-atomic-builtin-to-SPIRVBuiltins.td
wenju-he Mar 24, 2025
764b3f3
add noexcept to declarations in header
wenju-he Mar 24, 2025
a4c12e5
Merge branch 'sycl' into move-atomic-builtin-to-SPIRVBuiltins.td
wenju-he Apr 2, 2025
bbafeca
Merge branch 'sycl' into move-atomic-builtin-to-SPIRVBuiltins.td
wenju-he Apr 10, 2025
3b087df
clang-format
wenju-he Apr 10, 2025
7dab59d
Merge branch 'sycl' into move-atomic-builtin-to-SPIRVBuiltins.td
wenju-he Apr 21, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
98 changes: 89 additions & 9 deletions sycl/include/sycl/__spirv/spirv_ops.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -232,37 +232,117 @@ extern __DPCPP_SYCL_EXTERNAL
#define __SYCL_OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy
#define __SYCL_OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy

// Atomic SPIR-V builtins
#define __SPIRV_ATOMIC_LOAD(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicLoad(AS Type *P, int S, \
int O);
#define __SPIRV_ATOMIC_STORE(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL void __spirv_AtomicStore(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_EXCHANGE(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicExchange(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_CMP_EXCHANGE(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicCompareExchange( \
AS Type *P, int S, int E, int U, Type V, Type C);
#define __SPIRV_ATOMIC_IADD(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicIAdd(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_ISUB(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicISub(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_FADD(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicFAddEXT(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_SMIN(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicSMin(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_UMIN(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicUMin(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_FMIN(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicFMinEXT(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_SMAX(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicSMax(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_UMAX(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicUMax(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_FMAX(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicFMaxEXT(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_AND(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicAnd(AS Type *P, int S, \
int O, Type V);
#define __SPIRV_ATOMIC_OR(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicOr(AS Type *P, int S, int O, \
Type V);
#define __SPIRV_ATOMIC_XOR(AS, Type) \
extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicXor(AS Type *P, int S, \
Copy link
Contributor

@aelovikov-intel aelovikov-intel Apr 21, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

For my education, why do we still need these?

Copy link
Contributor Author

@wenju-he wenju-he Apr 22, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@aelovikov-intel it is related to long long type, please check comment: #17471 (comment)

int O, Type V);

#define __SPIRV_ATOMIC_FLOAT(AS, Type) \
__SPIRV_ATOMIC_FADD(AS, Type) \
__SPIRV_ATOMIC_FMIN(AS, Type) \
__SPIRV_ATOMIC_FMAX(AS, Type) \
__SPIRV_ATOMIC_LOAD(AS, Type) \
__SPIRV_ATOMIC_STORE(AS, Type) \
__SPIRV_ATOMIC_EXCHANGE(AS, Type)

#define __SPIRV_ATOMIC_BASE(AS, Type) \
__SPIRV_ATOMIC_FLOAT(AS, Type) \
__SPIRV_ATOMIC_CMP_EXCHANGE(AS, Type) \
__SPIRV_ATOMIC_IADD(AS, Type) \
__SPIRV_ATOMIC_ISUB(AS, Type) \
__SPIRV_ATOMIC_AND(AS, Type) \
__SPIRV_ATOMIC_OR(AS, Type) \
__SPIRV_ATOMIC_XOR(AS, Type)

#define __SPIRV_ATOMIC_SIGNED(AS, Type) \
__SPIRV_ATOMIC_BASE(AS, Type) \
__SPIRV_ATOMIC_SMIN(AS, Type) \
__SPIRV_ATOMIC_SMAX(AS, Type)

#define __SPIRV_ATOMIC_UNSIGNED(AS, Type) \
__SPIRV_ATOMIC_BASE(AS, Type) \
__SPIRV_ATOMIC_UMIN(AS, Type) \
__SPIRV_ATOMIC_UMAX(AS, Type)

// Helper atomic operations which select correct signed/unsigned version
// of atomic min/max based on the type
#define __SPIRV_ATOMIC_MINMAX(AS, Op) \
template <typename T> \
typename std::enable_if_t< \
std::is_integral<T>::value && std::is_signed<T>::value, T> \
__spirv_Atomic##Op(AS T *Ptr, __spv::Scope::Flag Memory, \
__spv::MemorySemanticsMask::Flag Semantics, \
T Value) { \
__spirv_Atomic##Op(AS T *Ptr, int Memory, int Semantics, T Value) { \
return __spirv_AtomicS##Op(Ptr, Memory, Semantics, Value); \
} \
template <typename T> \
typename std::enable_if_t< \
std::is_integral<T>::value && !std::is_signed<T>::value, T> \
__spirv_Atomic##Op(AS T *Ptr, __spv::Scope::Flag Memory, \
__spv::MemorySemanticsMask::Flag Semantics, \
T Value) { \
__spirv_Atomic##Op(AS T *Ptr, int Memory, int Semantics, T Value) { \
return __spirv_AtomicU##Op(Ptr, Memory, Semantics, Value); \
} \
template <typename T> \
typename std::enable_if_t<std::is_floating_point<T>::value, T> \
__spirv_Atomic##Op(AS T *Ptr, __spv::Scope::Flag Memory, \
__spv::MemorySemanticsMask::Flag Semantics, \
T Value) { \
__spirv_Atomic##Op(AS T *Ptr, int Memory, int Semantics, T Value) { \
return __spirv_AtomicF##Op##EXT(Ptr, Memory, Semantics, Value); \
}

#define __SPIRV_ATOMICS(macro, Arg) \
macro(__attribute__((opencl_global)), Arg) \
macro(__attribute__((opencl_local)), Arg) macro(, Arg)

__SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, _Float16)
__SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, float)
__SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, double)
__SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED, int)
__SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED, long)
Copy link
Contributor Author

@wenju-he wenju-he Mar 18, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I can't remove this part and below unsigned int/long variants because when a test uses both long long type from this header and non-long long type from clang SPIRVBuiltins.td, there is build error

using AtomicRefT =
    atomic_ref<long long, memory_order::relaxed, memory_scope::device>;

using AtomicRefT2 =
    atomic_ref<long, memory_order::relaxed, memory_scope::device>;

In file included from Regression/pf-wg-atomic64.cpp:9:
In file included from intel-llvm/build/bin/../include/sycl/atomic_ref.hpp:13:
In file included from intel-llvm/build/bin/../include/sycl/ext/oneapi/experimental/address_cast.hpp:11:
intel-llvm/build/bin/../include/sycl/detail/spirv.hpp:654:10: error: no matching function for call to '__spirv_AtomicIAdd'
  654 |   return __spirv_AtomicIAdd(Ptr, SPIRVScope, SPIRVOrder, Value);
      |          ^~~~~~~~~~~~~~~~~~
intel-llvm/build/bin/../include/sycl/atomic_ref.hpp:311:27: note: in instantiation of function template specialization 'sycl::detail::spirv::AtomicIAdd<long, sycl::access::address_space::generic_space, sycl::access::decorated::no>' requested here
  311 |     return detail::spirv::AtomicIAdd(ptr, scope, order, operand);
      |                           ^
intel-llvm/build/bin/../include/sycl/atomic_ref.hpp:319:12: note: in instantiation of member function 'sycl::detail::atomic_ref_impl<long, 4, sycl::memory_order::relaxed, sycl::memory_scope::device, sycl::access::address_space::generic_space>::fetch_add' requested here
  319 |     return fetch_add(operand) + operand;
      |            ^
Regression/pf-wg-atomic64.cpp:45:19: note: in instantiation of member function 'sycl::detail::atomic_ref_impl<long, 4, sycl::memory_order::relaxed, sycl::memory_scope::device, sycl::access::address_space::generic_space>::operator+=' requested here
   45 |           feature += 42;
      |                   ^
intel-llvm/build/bin/../include/sycl/__spirv/spirv_ops.hpp:337:17: note: candidate function not viable: no known conversion from 'long *' to '__global long long *' for 1st argument
  337 | __SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED, long long)
      | ~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
intel-llvm/build/bin/../include/sycl/__spirv/spirv_ops.hpp:334:3: note: expanded from macro '__SPIRV_ATOMICS'
  334 |   macro(__attribute__((opencl_global)), Arg)                                   \
      |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
intel-llvm/build/bin/../include/sycl/__spirv/spirv_ops.hpp:303:3: note: expanded from macro '__SPIRV_ATOMIC_SIGNED'
  303 |   __SPIRV_ATOMIC_BASE(AS, Type)                                                \
      |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~
intel-llvm/build/bin/../include/sycl/__spirv/spirv_ops.hpp:296:3: note: expanded from macro '__SPIRV_ATOMIC_BASE'
  296 |   __SPIRV_ATOMIC_IADD(AS, Type)                                                \
      |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~
intel-llvm/build/bin/../include/sycl/__spirv/spirv_ops.hpp:249:37: note: expanded from macro '__SPIRV_ATOMIC_IADD'
  249 |   extern __DPCPP_SYCL_EXTERNAL Type __spirv_AtomicIAdd(                        \
      |                                     ^

Please advise how to proceed.

Copy link

@Naghasan Naghasan Mar 19, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That's linked to this I think https://github.com/intel/llvm/pull/1384/files#diff-73b1c844b6929ccc37836285e2d87f89ead7074c6693a80e712f06bb54fa5cb2, I'll need to refresh my memory and come back to you as it is a bit old now.

Copy link
Contributor Author

@wenju-he wenju-he Apr 2, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'll keep the declarations in sycl/include/sycl/__spirv/spirv_ops.hpp until there is a stable solution. I'll check it later.

__SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED, long long)
__SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED, unsigned int)
__SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED, unsigned long)
__SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED, unsigned long long)
__SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Min)
__SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max)

Expand Down
Loading