diff --git a/libclc/generic/include/lp64_types.h b/libclc/generic/include/lp64_types.h index f5609d4d3c047..2cf6a948a5d4b 100644 --- a/libclc/generic/include/lp64_types.h +++ b/libclc/generic/include/lp64_types.h @@ -120,7 +120,7 @@ typedef _Float16 __clc_vec8_float16_t __attribute__((ext_vector_type(8))); typedef _Float16 __clc_vec16_float16_t __attribute__((ext_vector_type(16))); #endif -typedef __clc_int64_t __clc_size_t; +typedef __clc_uint64_t __clc_size_t; typedef event_t __clc_event_t; diff --git a/sycl/test/regression/async_work_group_copy.cpp b/sycl/test/regression/async_work_group_copy.cpp new file mode 100644 index 0000000000000..e42a644a0fae6 --- /dev/null +++ b/sycl/test/regression/async_work_group_copy.cpp @@ -0,0 +1,63 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o - + +// Test checks for that no compile errors occur for +// builtin async_work_group_copy +#include + +using namespace cl::sycl; + +// Define the number of work items to enqueue. +const size_t NElems = 32; +const size_t WorkGroupSize = 8; +const size_t NWorkGroups = NElems / WorkGroupSize; + +template void async_work_group_test() { + queue Q; + + buffer InBuf(NElems); + buffer OutBuf(NElems); + + Q.submit([&](handler &CGH) { + auto In = InBuf.template get_access(CGH); + auto Out = OutBuf.template get_access(CGH); + accessor Local( + range<1>{WorkGroupSize}, CGH); + + nd_range<1> NDR{range<1>(NElems), range<1>(WorkGroupSize)}; + CGH.parallel_for(NDR, [=](nd_item<1> NDId) { + auto GrId = NDId.get_group_linear_id(); + auto Group = NDId.get_group(); + size_t Offset = GrId * WorkGroupSize; + auto E = NDId.async_work_group_copy( + Local.get_pointer(), In.get_pointer() + Offset, WorkGroupSize); + E = NDId.async_work_group_copy(Out.get_pointer() + Offset, + Local.get_pointer(), WorkGroupSize); + }); + }).wait(); +} + +template void test() { + async_work_group_test(); + async_work_group_test>(); + async_work_group_test>(); + async_work_group_test>(); + async_work_group_test>(); + async_work_group_test>(); + async_work_group_test>(); + async_work_group_test, 2>>(); + async_work_group_test, 3>>(); + async_work_group_test, 4>>(); + async_work_group_test, 8>>(); + async_work_group_test, 16>>(); +} + +int main() { + test(); + test(); + test(); + test(); + test(); + test(); + test(); + return 1; +}