Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

[SYCL] Add support for the generic address space in tests for sycl::a… #619

Merged
merged 8 commits into from Dec 23, 2021
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
12 changes: 8 additions & 4 deletions SYCL/AtomicRef/add.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -11,10 +11,14 @@ using namespace sycl;

// Floating-point types do not support pre- or post-increment
template <> void add_test<float>(queue q, size_t N) {
add_fetch_test<::sycl::ext::oneapi::atomic_ref, float>(q, N);
add_fetch_test<::sycl::atomic_ref, float>(q, N);
add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, float>(q, N);
add_plus_equal_test<::sycl::atomic_ref, float>(q, N);
add_fetch_test<::sycl::ext::oneapi::atomic_ref,
access::address_space::global_space, float>(q, N);
add_fetch_test<::sycl::atomic_ref, access::address_space::global_space,
float>(q, N);
add_plus_equal_test<::sycl::ext::oneapi::atomic_ref,
access::address_space::global_space, float>(q, N);
add_plus_equal_test<::sycl::atomic_ref, access::address_space::global_space,
float>(q, N);
}

int main() {
Expand Down
58 changes: 42 additions & 16 deletions SYCL/AtomicRef/add.h
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,8 @@ using namespace sycl::ext::oneapi;

template <template <typename, memory_order, memory_scope, access::address_space>
class AtomicRef,
typename T, typename Difference = T>
access::address_space address_space, typename T,
typename Difference = T>
void add_fetch_test(queue q, size_t N) {
T sum = 0;
std::vector<T> output(N);
Expand All @@ -27,7 +28,7 @@ void add_fetch_test(queue q, size_t N) {
cgh.parallel_for(range<1>(N), [=](item<1> it) {
int gid = it.get_id(0);
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
access::address_space::global_space>(sum[0]);
address_space>(sum[0]);
out[gid] = atm.fetch_add(Difference(1));
});
});
Expand All @@ -48,7 +49,8 @@ void add_fetch_test(queue q, size_t N) {

template <template <typename, memory_order, memory_scope, access::address_space>
class AtomicRef,
typename T, typename Difference = T>
access::address_space address_space, typename T,
typename Difference = T>
void add_plus_equal_test(queue q, size_t N) {
T sum = 0;
std::vector<T> output(N);
Expand All @@ -64,7 +66,7 @@ void add_plus_equal_test(queue q, size_t N) {
cgh.parallel_for(range<1>(N), [=](item<1> it) {
int gid = it.get_id(0);
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
access::address_space::global_space>(sum[0]);
address_space>(sum[0]);
out[gid] = atm += Difference(1);
});
});
Expand All @@ -85,7 +87,8 @@ void add_plus_equal_test(queue q, size_t N) {

template <template <typename, memory_order, memory_scope, access::address_space>
class AtomicRef,
typename T, typename Difference = T>
access::address_space address_space, typename T,
typename Difference = T>
void add_pre_inc_test(queue q, size_t N) {
T sum = 0;
std::vector<T> output(N);
Expand All @@ -101,7 +104,7 @@ void add_pre_inc_test(queue q, size_t N) {
cgh.parallel_for(range<1>(N), [=](item<1> it) {
int gid = it.get_id(0);
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
access::address_space::global_space>(sum[0]);
address_space>(sum[0]);
out[gid] = ++atm;
});
});
Expand All @@ -122,7 +125,8 @@ void add_pre_inc_test(queue q, size_t N) {

template <template <typename, memory_order, memory_scope, access::address_space>
class AtomicRef,
typename T, typename Difference = T>
access::address_space address_space, typename T,
typename Difference = T>
void add_post_inc_test(queue q, size_t N) {
T sum = 0;
std::vector<T> output(N);
Expand All @@ -138,7 +142,7 @@ void add_post_inc_test(queue q, size_t N) {
cgh.parallel_for(range<1>(N), [=](item<1> it) {
int gid = it.get_id(0);
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
access::address_space::global_space>(sum[0]);
address_space>(sum[0]);
out[gid] = atm++;
});
});
Expand All @@ -159,12 +163,34 @@ void add_post_inc_test(queue q, size_t N) {

template <typename T, typename Difference = T>
void add_test(queue q, size_t N) {
add_fetch_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
add_fetch_test<::sycl::atomic_ref, T, Difference>(q, N);
add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
add_plus_equal_test<::sycl::atomic_ref, T, Difference>(q, N);
add_pre_inc_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
add_pre_inc_test<::sycl::atomic_ref, T, Difference>(q, N);
add_post_inc_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
add_post_inc_test<::sycl::atomic_ref, T, Difference>(q, N);
add_fetch_test<::sycl::ext::oneapi::atomic_ref,
access::address_space::global_space, T, Difference>(q, N);
add_fetch_test<::sycl::atomic_ref, access::address_space::global_space, T,
Difference>(q, N);
add_plus_equal_test<::sycl::ext::oneapi::atomic_ref,
access::address_space::global_space, T, Difference>(q, N);
add_plus_equal_test<::sycl::atomic_ref, access::address_space::global_space,
T, Difference>(q, N);
add_pre_inc_test<::sycl::ext::oneapi::atomic_ref,
access::address_space::global_space, T, Difference>(q, N);
add_pre_inc_test<::sycl::atomic_ref, access::address_space::global_space, T,
Difference>(q, N);
add_post_inc_test<::sycl::ext::oneapi::atomic_ref,
access::address_space::global_space, T, Difference>(q, N);
add_post_inc_test<::sycl::atomic_ref, access::address_space::global_space, T,
Difference>(q, N);
}

template <typename T, typename Difference = T>
void add_generic_test(queue q, size_t N) {
add_fetch_test<::sycl::atomic_ref, access::address_space::generic_space, T,
Difference>(q, N);
add_plus_equal_test<::sycl::atomic_ref, access::address_space::generic_space,
T, Difference>(q, N);
add_pre_inc_test<::sycl::atomic_ref, access::address_space::generic_space, T,
Difference>(q, N);
add_post_inc_test<::sycl::atomic_ref, access::address_space::generic_space, T,
Difference>(q, N);
add_post_inc_test<::sycl::atomic_ref, access::address_space::global_space, T,
Difference>(q, N);
}
12 changes: 8 additions & 4 deletions SYCL/AtomicRef/add_atomic64.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -11,10 +11,14 @@ using namespace sycl;

// Floating-point types do not support pre- or post-increment
template <> void add_test<double>(queue q, size_t N) {
add_fetch_test<::sycl::ext::oneapi::atomic_ref, double>(q, N);
add_fetch_test<::sycl::atomic_ref, double>(q, N);
add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, double>(q, N);
add_plus_equal_test<::sycl::atomic_ref, double>(q, N);
add_fetch_test<::sycl::ext::oneapi::atomic_ref,
access::address_space::global_space, double>(q, N);
add_fetch_test<::sycl::atomic_ref, access::address_space::global_space,
double>(q, N);
add_plus_equal_test<::sycl::ext::oneapi::atomic_ref,
access::address_space::global_space, double>(q, N);
add_plus_equal_test<::sycl::atomic_ref, access::address_space::global_space,
double>(q, N);
}

int main() {
Expand Down
52 changes: 52 additions & 0 deletions SYCL/AtomicRef/add_atomic64_generic.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,52 @@
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out \
// RUN: -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_60
// RUN: %HOST_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

// CUDA backend has had no support for the generic address space yet
// XFAIL: cuda

#include "add.h"
#include <iostream>
using namespace sycl;

// Floating-point types do not support pre- or post-increment
template <> void add_generic_test<double>(queue q, size_t N) {
add_fetch_test<::sycl::atomic_ref, access::address_space::generic_space,
double>(q, N);
add_plus_equal_test<::sycl::atomic_ref, access::address_space::generic_space,
double>(q, N);
}

int main() {
queue q;

if (!q.get_device().has(aspect::atomic64)) {
std::cout << "Skipping test\n";
return 0;
}

constexpr int N = 32;
add_generic_test<double>(q, N);

// Include long tests if they are 64 bits wide
if constexpr (sizeof(long) == 8) {
add_generic_test<long>(q, N);
add_generic_test<unsigned long>(q, N);
}

// Include long long tests if they are 64 bits wide
if constexpr (sizeof(long long) == 8) {
add_generic_test<long long>(q, N);
add_generic_test<unsigned long long>(q, N);
}

// Include pointer tests if they are 64 bits wide
if constexpr (sizeof(char *) == 8) {
add_generic_test<char *, ptrdiff_t>(q, N);
}

std::cout << "Test passed." << std::endl;
}
43 changes: 43 additions & 0 deletions SYCL/AtomicRef/add_generic.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,43 @@
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out \
// RUN: -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_60
// RUN: %HOST_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

// CUDA backend has had no support for the generic address space yet
// XFAIL: cuda

#include "add.h"
#include <iostream>
using namespace sycl;

// Floating-point types do not support pre- or post-increment
template <> void add_generic_test<float>(queue q, size_t N) {
add_fetch_test<::sycl::atomic_ref, access::address_space::generic_space,
float>(q, N);
add_plus_equal_test<::sycl::atomic_ref, access::address_space::generic_space,
float>(q, N);
}

int main() {
queue q;

constexpr int N = 32;
add_generic_test<int>(q, N);
add_generic_test<unsigned int>(q, N);
add_generic_test<float>(q, N);

// Include long tests if they are 32 bits wide
if constexpr (sizeof(long) == 4) {
add_generic_test<long>(q, N);
add_generic_test<unsigned long>(q, N);
}

// Include pointer tests if they are 32 bits wide
if constexpr (sizeof(char *) == 4) {
add_generic_test<char *, ptrdiff_t>(q, N);
}

std::cout << "Test passed." << std::endl;
}
19 changes: 13 additions & 6 deletions SYCL/AtomicRef/assignment.h
Original file line number Diff line number Diff line change
Expand Up @@ -11,12 +11,12 @@ using namespace sycl::ext::oneapi;

template <template <typename, memory_order, memory_scope, access::address_space>
class AtomicRef,
typename T>
access::address_space address_space, typename T>
class assignment_kernel;

template <template <typename, memory_order, memory_scope, access::address_space>
class AtomicRef,
typename T>
access::address_space address_space, typename T>
void assignment_test(queue q, size_t N) {
T initial = T(N);
T assignment = initial;
Expand All @@ -25,11 +25,11 @@ void assignment_test(queue q, size_t N) {
q.submit([&](handler &cgh) {
auto st =
assignment_buf.template get_access<access::mode::read_write>(cgh);
cgh.parallel_for<assignment_kernel<AtomicRef, T>>(
cgh.parallel_for<assignment_kernel<AtomicRef, address_space, T>>(
range<1>(N), [=](item<1> it) {
size_t gid = it.get_id(0);
auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
access::address_space::global_space>(st[0]);
address_space>(st[0]);
atm = T(gid);
});
});
Expand All @@ -42,6 +42,13 @@ void assignment_test(queue q, size_t N) {
}

template <typename T> void assignment_test(queue q, size_t N) {
assignment_test<::sycl::ext::oneapi::atomic_ref, T>(q, N);
assignment_test<::sycl::atomic_ref, T>(q, N);
assignment_test<::sycl::ext::oneapi::atomic_ref,
access::address_space::global_space, T>(q, N);
assignment_test<::sycl::atomic_ref, access::address_space::global_space, T>(
q, N);
}

template <typename T> void assignment_generic_test(queue q, size_t N) {
assignment_test<::sycl::atomic_ref, access::address_space::generic_space, T>(
q, N);
}
43 changes: 43 additions & 0 deletions SYCL/AtomicRef/assignment_atomic64_generic.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,43 @@
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
// RUN: %HOST_RUN_PLACEHOLDER %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

// CUDA backend has had no support for the generic address space yet
// XFAIL: cuda

#include "assignment.h"
#include <iostream>
using namespace sycl;

int main() {
queue q;

if (!q.get_device().has(aspect::atomic64)) {
std::cout << "Skipping test\n";
return 0;
}

constexpr int N = 32;
assignment_generic_test<double>(q, N);

// Include long tests if they are 64 bits wide
if constexpr (sizeof(long) == 8) {
assignment_generic_test<long>(q, N);
assignment_generic_test<unsigned long>(q, N);
}

// Include long long tests if they are 64 bits wide
if constexpr (sizeof(long long) == 8) {
assignment_generic_test<long long>(q, N);
assignment_generic_test<unsigned long long>(q, N);
}

// Include pointer tests if they are 64 bits wide
if constexpr (sizeof(char *) == 8) {
assignment_generic_test<char *>(q, N);
}

std::cout << "Test passed." << std::endl;
}
34 changes: 34 additions & 0 deletions SYCL/AtomicRef/assignment_generic.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,34 @@
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
// RUN: %HOST_RUN_PLACEHOLDER %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

// CUDA backend has had no support for the generic address space yet
// XFAIL: cuda

#include "assignment.h"
#include <iostream>
using namespace sycl;

int main() {
queue q;

constexpr int N = 32;
assignment_generic_test<int>(q, N);
assignment_generic_test<unsigned int>(q, N);
assignment_generic_test<float>(q, N);

// Include long tests if they are 32 bits wide
if constexpr (sizeof(long) == 4) {
assignment_generic_test<long>(q, N);
assignment_generic_test<unsigned long>(q, N);
}

// Include pointer tests if they are 32 bits wide
if constexpr (sizeof(char *) == 4) {
assignment_generic_test<char *>(q, N);
}

std::cout << "Test passed." << std::endl;
}
Loading