Skip to content

[SYCL][libclc][E2E] atomic work_item scope fallback #16172

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 2 commits into from
Dec 2, 2024
Merged

[SYCL][libclc][E2E] atomic work_item scope fallback #16172

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
1e09ffc
[SYCL][libclc][E2E] atomic work_item scope fallback
Seanst98 Nov 20, 2024
eb65fce
Add missing free to test
Seanst98 Nov 25, 2024
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
1 change: 1 addition & 0 deletions libclc/amdgcn-amdhsa/libspirv/atomic/atomic_helpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,7 @@ extern int __oclc_amdgpu_reflect(__constant char *);
OUT_ORDER) \
{ \
switch (IN_SCOPE) { \
case Invocation: \
case Subgroup: \
OUT_SCOPE = __HIP_MEMORY_SCOPE_WAVEFRONT; \
break; \
Expand Down
2 changes: 2 additions & 0 deletions libclc/ptx-nvidiacl/libspirv/atomic/atomic_cmpxchg.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,7 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
#define __CLC_NVVM_ATOMIC_CAS_IMPL_ORDER(TYPE, TYPE_NV, TYPE_MANGLED_NV, OP, \
ADDR_SPACE, ADDR_SPACE_NV, ORDER) \
switch (scope) { \
case Invocation: \
case Subgroup: \
case Workgroup: { \
if (__clc_nvvm_reflect_arch() >= 600) { \
Expand Down Expand Up @@ -44,6 +45,7 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
#define __CLC_NVVM_ATOMIC_CAS_IMPL_ACQUIRE_FENCE( \
TYPE, TYPE_NV, TYPE_MANGLED_NV, OP, ADDR_SPACE, ADDR_SPACE_NV) \
switch (scope) { \
case Invocation: \
case Subgroup: \
case Workgroup: { \
if (__clc_nvvm_reflect_arch() >= 600) { \
Expand Down
2 changes: 2 additions & 0 deletions libclc/ptx-nvidiacl/libspirv/atomic/atomic_helpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,7 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
#define __CLC_NVVM_ATOMIC_IMPL_ORDER(TYPE, TYPE_NV, TYPE_MANGLED_NV, OP, \
ADDR_SPACE, ADDR_SPACE_NV, ORDER) \
switch (scope) { \
case Invocation: \
case Subgroup: \
case Workgroup: { \
if (__clc_nvvm_reflect_arch() >= 600) { \
Expand Down Expand Up @@ -46,6 +47,7 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
#define __CLC_NVVM_ATOMIC_IMPL_ACQUIRE_FENCE(TYPE, TYPE_NV, TYPE_MANGLED_NV, \
OP, ADDR_SPACE, ADDR_SPACE_NV) \
switch (scope) { \
case Invocation: \
case Subgroup: \
case Workgroup: { \
if (__clc_nvvm_reflect_arch() >= 600) { \
Expand Down
1 change: 1 addition & 0 deletions libclc/ptx-nvidiacl/libspirv/atomic/atomic_load.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,7 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
#define __CLC_NVVM_ATOMIC_LOAD_IMPL_ORDER(TYPE, TYPE_NV, TYPE_MANGLED_NV, \
ADDR_SPACE, ADDR_SPACE_NV, ORDER) \
switch (scope) { \
case Invocation: \
case Subgroup: \
case Workgroup: { \
TYPE_NV res = __nvvm##ORDER##_cta_ld##ADDR_SPACE_NV##TYPE_MANGLED_NV( \
Expand Down
1 change: 1 addition & 0 deletions libclc/ptx-nvidiacl/libspirv/atomic/atomic_store.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,7 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int, unsigned int);
#define __CLC_NVVM_ATOMIC_STORE_IMPL_ORDER(TYPE, TYPE_NV, TYPE_MANGLED_NV, \
ADDR_SPACE, ADDR_SPACE_NV, ORDER) \
switch (scope) { \
case Invocation: \
case Subgroup: \
case Workgroup: { \
__nvvm##ORDER##_cta_st##ADDR_SPACE_NV##TYPE_MANGLED_NV( \
Expand Down
65 changes: 65 additions & 0 deletions sycl/test-e2e/AtomicRef/work_item_scope.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,65 @@
// RUN: %{build} -o %t.out
// RUN: %{run} %t.out

#include <iostream>
#include <sycl/atomic_ref.hpp>
#include <sycl/detail/core.hpp>
#include <sycl/usm.hpp>

int main() {

sycl::device dev;
sycl::queue q(dev);
auto ctxt = q.get_context();

// This test does not validate any output
// Only that the work_item scope does not error
try {

// Allocate device memory
int *data = sycl::malloc_device<int>(1, q);
Copy link
Contributor

@aelovikov-intel aelovikov-intel Nov 25, 2024

Choose a reason for hiding this comment

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

This needs to be free'ed properly.

Copy link
Contributor Author

@Seanst98 Seanst98 Nov 25, 2024

Choose a reason for hiding this comment

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

Good catch. Fixed.

Copy link
Contributor

@aelovikov-intel aelovikov-intel Nov 25, 2024

Choose a reason for hiding this comment

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

Fixed.

That's why I dislike this part of the SYCL - no RAII for USM :(

Is that common for other e2e tests to "leak" USM on exceptions?


q.submit([&](sycl::handler &cgh) {
cgh.parallel_for(10, [=](sycl::id<> id) {
data[0] = 0;

// Check atomic_ref functionality
sycl::atomic_ref<int, sycl::memory_order::relaxed,
sycl::memory_scope::work_item,
sycl::access::address_space::generic_space>
at(data[0]);

auto lock = at.is_lock_free();
at.store(1);
auto load = at.load();
auto xch = at.exchange(2);
auto weak =
at.compare_exchange_weak(data[0], 3, sycl::memory_order::relaxed,
sycl::memory_order::relaxed);
auto strong =
at.compare_exchange_strong(data[0], 4, sycl::memory_order::relaxed,
sycl::memory_order::relaxed);
auto fetch_add = at.fetch_add(5);
auto fetch_sub = at.fetch_sub(6);
auto fetch_and = at.fetch_and(7);
auto fetch_or = at.fetch_or(8);
auto fetch_xor = at.fetch_xor(9);
auto fetch_min = at.fetch_min(10);
auto fetch_max = at.fetch_max(11);
});
});
q.wait_and_throw();

sycl::free(data, q);

} catch (sycl::exception e) {
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
return 1;
} catch (...) {
std::cerr << "Unknown exception caught!\n";
return 2;
}

std::cout << "Test passed!" << std::endl;
return 0;
}