@@ -23,15 +23,15 @@ namespace syclexp = sycl::ext::oneapi::experimental;
2323constexpr int SIZE = 16 ;
2424
2525SYCL_EXT_ONEAPI_FUNCTION_PROPERTY ((syclexp::nd_range_kernel<1 >))
26- void double_kernel (float *src, float *dst) {
26+ void scratch_kernel (float *src, float *dst) {
2727 size_t lid = syclext::this_work_item::get_nd_item<1 >().get_local_linear_id ();
2828 float *local_mem = (float *)syclexp::get_work_group_scratch_memory ();
2929 local_mem[lid] = 2 * src[lid];
3030 dst[lid] = local_mem[lid];
3131}
3232
3333SYCL_EXT_ONEAPI_FUNCTION_PROPERTY ((syclexp::nd_range_kernel<1 >))
34- void square_kernel (float *src, float *dst) {
34+ void static_kernel (float *src, float *dst) {
3535 sycl::nd_item<1 > item = syclext::this_work_item::get_nd_item<1 >();
3636 size_t lid = item.get_local_linear_id ();
3737 syclexp::work_group_static<float [SIZE]> local_mem;
@@ -54,27 +54,31 @@ int main() {
5454 src[i] = i;
5555 }
5656
57- auto kbndl =
58- syclexp::get_kernel_bundle<double_kernel, sycl::bundle_state::executable>(
57+ auto scratchbndl = syclexp::get_kernel_bundle<scratch_kernel,
58+ sycl::bundle_state::executable>(
59+ q.get_context ());
60+ auto staticbndl =
61+ syclexp::get_kernel_bundle<static_kernel, sycl::bundle_state::executable>(
5962 q.get_context ());
60- sycl::kernel DoubleKernel =
61- kbndl.template ext_oneapi_get_kernel <double_kernel>();
62- sycl::kernel SquareKernel =
63- kbndl.template ext_oneapi_get_kernel <square_kernel>();
64- syclexp::launch_config DoubleKernelcfg{
63+
64+ sycl::kernel ScratchKernel =
65+ scratchbndl.template ext_oneapi_get_kernel <scratch_kernel>();
66+ sycl::kernel StaticKernel =
67+ staticbndl.template ext_oneapi_get_kernel <static_kernel>();
68+ syclexp::launch_config ScratchKernelcfg{
6569 ::sycl::nd_range<1 >(::sycl::range<1 >(SIZE), ::sycl::range<1 >(SIZE)),
6670 syclexp::properties{
6771 syclexp::work_group_scratch_size (SIZE * sizeof (float ))}};
68- syclexp::launch_config SquareKernelcfg {
72+ syclexp::launch_config StaticKernelcfg {
6973 ::sycl::nd_range<1 >(::sycl::range<1 >(SIZE), ::sycl::range<1 >(SIZE))};
7074
71- syclexp::nd_launch (q, DoubleKernelcfg, DoubleKernel , src, dst);
75+ syclexp::nd_launch (q, ScratchKernelcfg, ScratchKernel , src, dst);
7276 q.wait ();
7377 for (int i = 0 ; i < SIZE; i++) {
7478 assert (dst[i] == 2 * src[i]);
7579 }
7680
77- syclexp::nd_launch (q, SquareKernelcfg, SquareKernel , src, dst);
81+ syclexp::nd_launch (q, StaticKernelcfg, StaticKernel , src, dst);
7882 q.wait ();
7983 for (int i = 0 ; i < SIZE; i++) {
8084 assert (dst[i] == src[i] * src[i]);
0 commit comments