@@ -119,21 +119,30 @@ class atomic_ref_base {
119119 }
120120};
121121
122+ // The spec requires that atomic_ref<T> is only valid for types T that are 4 or 8 bytes in size
123+ // (8 is only valid for devices with aspect::atomic64, but we can't check that at compile time)
124+ template <typename T>
125+ concept ValidAtomicSize = (sizeof (int ) == 4 || sizeof (int ) == 8 );
126+
122127// The spec provides an explicit list of which types are valid for atomic operations
123128template <typename T>
124129concept ValidAtomicIntType
125130 = (std::same_as<T, int > || std::same_as<T, unsigned int > || std::same_as<T, long > || std::same_as<T, unsigned long >
126131 || std::same_as<T, long long > || std::same_as<T, unsigned long long >)
127- && ( sizeof (T) == 4 || sizeof (T) == 8 ) ;
132+ && ValidAtomicSize<T> ;
128133
129134template <typename T>
130- concept ValidAtomicFloatType = std::same_as<T, float > || std::same_as<T, double >;
135+ concept ValidAtomicFloatType = ( std::same_as<T, float > || std::same_as<T, double >) && ValidAtomicSize<T >;
131136// TODO: for double we would also need to check if
132137// "the code containing this sycl::atomic_ref was submitted to a device that has sycl::aspect::atomic64"
133138// that's quite challenging to implement as-is, since we would need to do this at runtime when operations are performed
134139
135140template <typename T>
136- concept ValidAtomicType = ValidAtomicIntType<T> || ValidAtomicFloatType<T>;
141+ concept ValidAtomicPointerType
142+ = std::is_pointer_v<T> && std::is_object_v<std::remove_pointer_t <T>> && ValidAtomicSize<T>;
143+
144+ template <typename T>
145+ concept ValidAtomicType = ValidAtomicIntType<T> || ValidAtomicFloatType<T> || ValidAtomicPointerType<T>;
137146
138147} // namespace simsycl::detail
139148
0 commit comments