Skip to content

Commit

Permalink
Changed try_compile with try_run in cmake SYCL feature check
Browse files Browse the repository at this point in the history
  • Loading branch information
Luigi-Crisci committed Jun 3, 2024
1 parent 83712f6 commit 993728e
Show file tree
Hide file tree
Showing 5 changed files with 47 additions and 15 deletions.
13 changes: 7 additions & 6 deletions cmake/HasFeatures.cmake
Original file line number Diff line number Diff line change
@@ -1,12 +1,13 @@
macro(check_feature VAR FILENAME)
if(NOT DEFINED SYCL_BENCH_HAS_${VAR})
try_compile(SYCL_BENCH_HAS_${VAR} ${CMAKE_BINARY_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/cmake/has-features/src/${FILENAME}
CMAKE_FLAGS ${CMAKE_CXX_FLAGS}
OUTPUT_VARIABLE OUTPUT_VAR
)
if(NOT DEFINED RUN_RES_${VAR})
try_run(RUN_RES_${VAR} COMPILE_RES_${VAR} ${CMAKE_BINARY_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/cmake/has-features/src/${FILENAME}
CMAKE_FLAGS ${CMAKE_CXX_FLAGS}
COMPILE_OUTPUT_VARIABLE OUTPUT_VAR
RUN_OUTPUT_VARIABLE RUN_VAR
)
endif()

if (SYCL_BENCH_HAS_${VAR})
if (COMPILE_RES_${VAR} AND RUN_RES_${VAR} EQUAL 0)
set(RES ON)
else()
set(RES OFF)
Expand Down
4 changes: 4 additions & 0 deletions cmake/has-features/src/fp64_support_dummy.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -8,4 +8,8 @@ int main() {
sycl::accessor a(x, cgh, sycl::read_write);
cgh.parallel_for<class dummy>(sycl::range<1>(1), [=](sycl::id<1> idx) { a[idx] = 0; });
});

sycl::host_accessor host{x};
assert(host[0] == 0);

}
13 changes: 9 additions & 4 deletions cmake/has-features/src/group_algorithms_dummy.cpp
Original file line number Diff line number Diff line change
@@ -1,12 +1,17 @@
#include <sycl/sycl.hpp>
#include <iostream>


int main() {
sycl::queue q;

int* i = sycl::malloc_shared<int>(1, q);
q.submit([&](sycl::handler& cgh) {
q.parallel_for(sycl::nd_range<1>{{1}, {1}}, [=](sycl::nd_item<1> item) {
cgh.parallel_for(sycl::nd_range<1>{{1}, {1}}, [=](sycl::nd_item<1> item) {
// call only the group algorithms used in SYCL-Bench
sycl::reduce_over_group(item.get_group(), 0, sycl::plus<int>{});
*i = sycl::reduce_over_group(item.get_group(), 1, sycl::plus<int>{});
});
});
}).wait();

assert(*i == 1);
sycl::free(i, q);
}
7 changes: 5 additions & 2 deletions cmake/has-features/src/kernel_reduction_dummy.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,9 @@ int main() {
auto r = sycl::reduction(x, cgh, sycl::plus<int>{});
#endif

cgh.parallel_for(sycl::range<1>{1}, r, [=](sycl::id<1> idx, auto& op) { op.combine(1); });
});
cgh.parallel_for(sycl::range<1>{5}, r, [=](sycl::id<1> idx, auto& op) { op.combine(1); });
}).wait();

sycl::host_accessor host{x};
assert(host[0] == 5);
}
25 changes: 22 additions & 3 deletions cmake/has-features/src/spec_constants_dummy.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -6,14 +6,33 @@ static constexpr sycl::specialization_id<int> x;

int main() {
sycl::queue q;

q.submit([&](sycl::handler& cgh) { cgh.set_specialization_constant<x>(5); });
int* i = sycl::malloc_shared<int>(1, q);
q.submit([&](sycl::handler& cgh) {
cgh.set_specialization_constant<x>(5);
cgh.parallel_for(sycl::range(1), [=](sycl::item<1> item, sycl::kernel_handler h) {
*i = h.get_specialization_constant<x>();
});
}).wait();

assert(*i == 5);
sycl::free(i, q);
}

#else

// AdaptiveCpp implements sycl::specialized instead of spec constants

int main() { sycl::specialized<int> x(5); }
int main() {
sycl::queue q;
sycl::specialized<int> x;
x = 5; //Requires copy assignment operator
int* i = sycl::malloc_shared<int>(1, q);
q.parallel_for(sycl::range(1), [=](sycl::id<1> idx) {
*i = x;
}).wait();

assert(*i == 5);
sycl::free(i, q);
}

#endif

0 comments on commit 993728e

Please sign in to comment.