Skip to content

Commit c4a9552

Browse files
committed
Add test source kernel
1 parent df3dfab commit c4a9552

File tree

1 file changed

+21
-0
lines changed

1 file changed

+21
-0
lines changed

dpctl/tests/test_work_group_memory.py

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,27 @@ def get_spirv_abspath(fn):
3030
return spirv_file
3131

3232

33+
# The kernel in the SPIR-V file used in this test was generated from the
34+
# following SYCL source code:
35+
# #include <sycl/sycl.hpp>
36+
# using namespace sycl;
37+
# namespace syclexp = sycl::ext::oneapi::experimental;
38+
# namespace syclext = sycl::ext::oneapi;
39+
# using data_t = int32_t;
40+
#
41+
# extern "C" SYCL_EXTERNAL
42+
# SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
43+
# void local_mem_kernel(data_t* in, data_t* out,
44+
# syclexp::work_group_memory<data_t> mem){
45+
# auto* local_mem = &mem;
46+
# auto item = syclext::this_work_item::get_nd_item<1>();
47+
# size_t global_id = item.get_global_linear_id();
48+
# size_t local_id = item.get_local_linear_id();
49+
# local_mem[local_id] = in[global_id];
50+
# out[global_id] = local_mem[local_id];
51+
# }
52+
53+
3354
def test_submit_work_group_memory():
3455
if not dpctl.experimental.WorkGroupMemory.is_available():
3556
pytest.skip("Work group memory extension not supported")

0 commit comments

Comments
 (0)