diff --git a/tests/vector_load_store/generate_vector_load_store.py b/tests/vector_load_store/generate_vector_load_store.py index 411a8571b..9e94d8cf4 100755 --- a/tests/vector_load_store/generate_vector_load_store.py +++ b/tests/vector_load_store/generate_vector_load_store.py @@ -32,7 +32,7 @@ TEST_NAME = 'LOAD_STORE' -load_store_test_template = Template( +global_multi_ptr_load_store_test_template = Template( """ { ${type} inputData${type_as_str}${size}[${size}] = {${in_order_vals}}; @@ -56,7 +56,7 @@ auto testVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); auto multiPtrIn${type_as_str}${size} = inPtr${type_as_str}${size}.get_multi_ptr<${decorated}>(); - sycl::global_ptr constMultiPtrIn${type_as_str}${size} = multiPtrIn${type_as_str}${size}; + sycl::global_ptr constMultiPtrIn${type_as_str}${size} = multiPtrIn${type_as_str}${size}; auto multiPtrOut${type_as_str}${size} = outPtr${type_as_str}${size}.get_multi_ptr<${decorated}>(); testVec${type_as_str}${size}.load(0, constMultiPtrIn${type_as_str}${size}); testVec${type_as_str}${size}.store(0, multiPtrOut${type_as_str}${size}); @@ -65,7 +65,7 @@ sycl::vec<${type}, ${size}> swizzledVec {cleanVec${type_as_str}${size}.template swizzle<${swizVals}>()}; auto multiPtrInSwizzle${type_as_str}${size} = swizzleInPtr${type_as_str}${size}.get_multi_ptr<${decorated}>(); - sycl::global_ptr constMultiPtrInSwizzle${type_as_str}${size} = multiPtrInSwizzle${type_as_str}${size}; + sycl::global_ptr constMultiPtrInSwizzle${type_as_str}${size} = multiPtrInSwizzle${type_as_str}${size}; auto multiPtrOutSwizzle${type_as_str}${size} = swizzleOutPtr${type_as_str}${size}.get_multi_ptr<${decorated}>(); swizzledVec.load(0, constMultiPtrInSwizzle${type_as_str}${size}); swizzledVec.store(0, multiPtrOutSwizzle${type_as_str}${size}); @@ -80,6 +80,316 @@ } """) +local_multi_ptr_load_store_test_template = Template( + """ + { + ${type} inputData${type_as_str}${size}[${size}] = {${in_order_vals}}; + ${type} outputData${type_as_str}${size}[${size}] = {${val}}; + ${type} swizzleInputData${type_as_str}${size}[${size}] = {${reverse_order_vals}}; + ${type} swizzleOutputData${type_as_str}${size}[${size}] = {${val}}; + { + sycl::buffer<${type}, 1> inBuffer${type_as_str}${size}(inputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> outBuffer${type_as_str}${size}(outputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleInBuffer${type_as_str}${size}(swizzleInputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleOutBuffer${type_as_str}${size}(swizzleOutputData${type_as_str}${size}, sycl::range<1>(${size})); + + testQueue.submit([&](sycl::handler &cgh) { + auto inPtr${type_as_str}${size} = inBuffer${type_as_str}${size}.get_access(cgh); + auto outPtr${type_as_str}${size} = outBuffer${type_as_str}${size}.get_access(cgh); + + auto swizzleInPtr${type_as_str}${size} = swizzleInBuffer${type_as_str}${size}.get_access(cgh); + auto swizzleOutPtr${type_as_str}${size} = swizzleOutBuffer${type_as_str}${size}.get_access(cgh); + + sycl::local_accessor<${type}> inLocalPtr${type_as_str}${size}(${size}, cgh); + sycl::local_accessor<${type}> outLocalPtr${type_as_str}${size}(${size}, cgh); + + sycl::local_accessor<${type}> swizzleInLocalPtr${type_as_str}${size}(${size}, cgh); + sycl::local_accessor<${type}> swizzleOutLocalPtr${type_as_str}${size}(${size}, cgh); + + cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(1), sycl::range<1>(1)), [=](sycl::nd_item<1>) { + for (unsigned i = 0; i < ${size}; ++i) { + inLocalPtr${type_as_str}${size}[i] = inPtr${type_as_str}${size}[i]; + swizzleInLocalPtr${type_as_str}${size}[i] = swizzleInPtr${type_as_str}${size}[i]; + } + + auto testVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + + auto rawPtrIn${type_as_str}${size} = inLocalPtr${type_as_str}${size}.get_multi_ptr<${decorated}>(); + sycl::local_ptr constRawPtrIn${type_as_str}${size} = rawPtrIn${type_as_str}${size}; + auto rawPtrOut${type_as_str}${size} = outLocalPtr${type_as_str}${size}.get_multi_ptr<${decorated}>(); + testVec${type_as_str}${size}.load(0, constRawPtrIn${type_as_str}${size}); + testVec${type_as_str}${size}.store(0, rawPtrOut${type_as_str}${size}); + + auto cleanVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + sycl::vec<${type}, ${size}> swizzledVec {cleanVec${type_as_str}${size}.template swizzle<${swizVals}>()}; + + auto rawPtrInSwizzle${type_as_str}${size} = swizzleInLocalPtr${type_as_str}${size}.get_multi_ptr<${decorated}>(); + sycl::local_ptr constRawPtrInSwizzle${type_as_str}${size} = rawPtrInSwizzle${type_as_str}${size}; + auto rawPtrOutSwizzle${type_as_str}${size} = swizzleOutLocalPtr${type_as_str}${size}.get_multi_ptr<${decorated}>(); + swizzledVec.load(0, constRawPtrInSwizzle${type_as_str}${size}); + swizzledVec.store(0, rawPtrOutSwizzle${type_as_str}${size}); + + for (unsigned i = 0; i < ${size}; ++i) { + outPtr${type_as_str}${size}[i] = outLocalPtr${type_as_str}${size}[i]; + swizzleOutPtr${type_as_str}${size}[i] = swizzleOutLocalPtr${type_as_str}${size}[i]; + } + }); + }); + + } + check_array_equality<${type}, ${size}>(log, inputData${type_as_str}${size}, outputData${type_as_str}${size}); + check_array_equality<${type}, ${size}>(log, swizzleInputData${type_as_str}${size}, swizzleOutputData${type_as_str}${size}); + + testQueue.wait_and_throw(); + } + """) + +private_multi_ptr_load_store_test_template = Template( + """ + { + ${type} inputData${type_as_str}${size}[${size}] = {${in_order_vals}}; + ${type} outputData${type_as_str}${size}[${size}] = {${val}}; + ${type} swizzleInputData${type_as_str}${size}[${size}] = {${reverse_order_vals}}; + ${type} swizzleOutputData${type_as_str}${size}[${size}] = {${val}}; + { + sycl::buffer<${type}, 1> inBuffer${type_as_str}${size}(inputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> outBuffer${type_as_str}${size}(outputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleInBuffer${type_as_str}${size}(swizzleInputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleOutBuffer${type_as_str}${size}(swizzleOutputData${type_as_str}${size}, sycl::range<1>(${size})); + + testQueue.submit([&](sycl::handler &cgh) { + auto inPtr${type_as_str}${size} = inBuffer${type_as_str}${size}.get_access(cgh); + auto outPtr${type_as_str}${size} = outBuffer${type_as_str}${size}.get_access(cgh); + + auto swizzleInPtr${type_as_str}${size} = swizzleInBuffer${type_as_str}${size}.get_access(cgh); + auto swizzleOutPtr${type_as_str}${size} = swizzleOutBuffer${type_as_str}${size}.get_access(cgh); + + sycl::local_accessor<${type}> inLocalPtr${type_as_str}${size}(${size}, cgh); + sycl::local_accessor<${type}> outLocalPtr${type_as_str}${size}(${size}, cgh); + + sycl::local_accessor<${type}> swizzleInLocalPtr${type_as_str}${size}(${size}, cgh); + sycl::local_accessor<${type}> swizzleOutLocalPtr${type_as_str}${size}(${size}, cgh); + + cgh.single_task([=]() { + ${type} inPrivatePtr${type_as_str}${size}[${size}]; + ${type} outPrivatePtr${type_as_str}${size}[${size}]; + + ${type} swizzleInPrivatePtr${type_as_str}${size}[${size}]; + ${type} swizzleOutPrivatePtr${type_as_str}${size}[${size}]; + + for (unsigned i = 0; i < ${size}; ++i) { + inPrivatePtr${type_as_str}${size}[i] = inPtr${type_as_str}${size}[i]; + swizzleInPrivatePtr${type_as_str}${size}[i] = swizzleInPtr${type_as_str}${size}[i]; + } + + auto testVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + + auto rawPtrIn${type_as_str}${size} = sycl::address_space_cast(inPrivatePtr${type_as_str}${size}); + sycl::private_ptr constRawPtrIn${type_as_str}${size} = rawPtrIn${type_as_str}${size}; + auto rawPtrOut${type_as_str}${size} = sycl::address_space_cast(outPrivatePtr${type_as_str}${size}); + testVec${type_as_str}${size}.load(0, constRawPtrIn${type_as_str}${size}); + testVec${type_as_str}${size}.store(0, rawPtrOut${type_as_str}${size}); + + auto cleanVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + sycl::vec<${type}, ${size}> swizzledVec {cleanVec${type_as_str}${size}.template swizzle<${swizVals}>()}; + + auto rawPtrInSwizzle${type_as_str}${size} = sycl::address_space_cast(swizzleInPrivatePtr${type_as_str}${size}); + sycl::private_ptr constRawPtrInSwizzle${type_as_str}${size} = rawPtrInSwizzle${type_as_str}${size}; + auto rawPtrOutSwizzle${type_as_str}${size} = sycl::address_space_cast(swizzleOutPrivatePtr${type_as_str}${size}); + swizzledVec.load(0, constRawPtrInSwizzle${type_as_str}${size}); + swizzledVec.store(0, rawPtrOutSwizzle${type_as_str}${size}); + + for (unsigned i = 0; i < ${size}; ++i) { + outPtr${type_as_str}${size}[i] = outPrivatePtr${type_as_str}${size}[i]; + swizzleOutPtr${type_as_str}${size}[i] = swizzleOutPrivatePtr${type_as_str}${size}[i]; + } + }); + }); + + } + check_array_equality<${type}, ${size}>(log, inputData${type_as_str}${size}, outputData${type_as_str}${size}); + check_array_equality<${type}, ${size}>(log, swizzleInputData${type_as_str}${size}, swizzleOutputData${type_as_str}${size}); + + testQueue.wait_and_throw(); + } + """) + +global_raw_ptr_load_store_test_template = Template( + """ + { + ${type} inputData${type_as_str}${size}[${size}] = {${in_order_vals}}; + ${type} outputData${type_as_str}${size}[${size}] = {${val}}; + ${type} swizzleInputData${type_as_str}${size}[${size}] = {${reverse_order_vals}}; + ${type} swizzleOutputData${type_as_str}${size}[${size}] = {${val}}; + { + sycl::buffer<${type}, 1> inBuffer${type_as_str}${size}(inputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> outBuffer${type_as_str}${size}(outputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleInBuffer${type_as_str}${size}(swizzleInputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleOutBuffer${type_as_str}${size}(swizzleOutputData${type_as_str}${size}, sycl::range<1>(${size})); + + testQueue.submit([&](sycl::handler &cgh) { + auto inPtr${type_as_str}${size} = inBuffer${type_as_str}${size}.get_access(cgh); + auto outPtr${type_as_str}${size} = outBuffer${type_as_str}${size}.get_access(cgh); + + auto swizzleInPtr${type_as_str}${size} = swizzleInBuffer${type_as_str}${size}.get_access(cgh); + auto swizzleOutPtr${type_as_str}${size} = swizzleOutBuffer${type_as_str}${size}.get_access(cgh); + + cgh.single_task([=]() { + auto testVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + + std::add_pointer_t<${type}> rawPtrIn${type_as_str}${size} = inPtr${type_as_str}${size}.get_multi_ptr().get_raw(); + const std::add_pointer_t<${type}> constRawPtrIn${type_as_str}${size} = rawPtrIn${type_as_str}${size}; + std::add_pointer_t<${type}> rawPtrOut${type_as_str}${size} = outPtr${type_as_str}${size}.get_multi_ptr().get_raw(); + testVec${type_as_str}${size}.load(0, constRawPtrIn${type_as_str}${size}); + testVec${type_as_str}${size}.store(0, rawPtrOut${type_as_str}${size}); + + auto cleanVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + sycl::vec<${type}, ${size}> swizzledVec {cleanVec${type_as_str}${size}.template swizzle<${swizVals}>()}; + + std::add_pointer_t<${type}> rawPtrInSwizzle${type_as_str}${size} = swizzleInPtr${type_as_str}${size}.get_multi_ptr().get_raw(); + const std::add_pointer_t<${type}> constRawPtrInSwizzle${type_as_str}${size} = rawPtrInSwizzle${type_as_str}${size}; + std::add_pointer_t<${type}> rawPtrOutSwizzle${type_as_str}${size} = swizzleOutPtr${type_as_str}${size}.get_multi_ptr().get_raw(); + swizzledVec.load(0, constRawPtrInSwizzle${type_as_str}${size}); + swizzledVec.store(0, rawPtrOutSwizzle${type_as_str}${size}); + }); + }); + + } + check_array_equality<${type}, ${size}>(log, inputData${type_as_str}${size}, outputData${type_as_str}${size}); + check_array_equality<${type}, ${size}>(log, swizzleInputData${type_as_str}${size}, swizzleOutputData${type_as_str}${size}); + + testQueue.wait_and_throw(); + } + """) + +local_raw_ptr_load_store_test_template = Template( + """ + { + ${type} inputData${type_as_str}${size}[${size}] = {${in_order_vals}}; + ${type} outputData${type_as_str}${size}[${size}] = {${val}}; + ${type} swizzleInputData${type_as_str}${size}[${size}] = {${reverse_order_vals}}; + ${type} swizzleOutputData${type_as_str}${size}[${size}] = {${val}}; + { + sycl::buffer<${type}, 1> inBuffer${type_as_str}${size}(inputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> outBuffer${type_as_str}${size}(outputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleInBuffer${type_as_str}${size}(swizzleInputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleOutBuffer${type_as_str}${size}(swizzleOutputData${type_as_str}${size}, sycl::range<1>(${size})); + + testQueue.submit([&](sycl::handler &cgh) { + auto inPtr${type_as_str}${size} = inBuffer${type_as_str}${size}.get_access(cgh); + auto outPtr${type_as_str}${size} = outBuffer${type_as_str}${size}.get_access(cgh); + + auto swizzleInPtr${type_as_str}${size} = swizzleInBuffer${type_as_str}${size}.get_access(cgh); + auto swizzleOutPtr${type_as_str}${size} = swizzleOutBuffer${type_as_str}${size}.get_access(cgh); + + sycl::local_accessor<${type}> inLocalPtr${type_as_str}${size}(${size}, cgh); + sycl::local_accessor<${type}> outLocalPtr${type_as_str}${size}(${size}, cgh); + + sycl::local_accessor<${type}> swizzleInLocalPtr${type_as_str}${size}(${size}, cgh); + sycl::local_accessor<${type}> swizzleOutLocalPtr${type_as_str}${size}(${size}, cgh); + + cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(1), sycl::range<1>(1)), [=](sycl::nd_item<1>) { + for (unsigned i = 0; i < ${size}; ++i) { + inLocalPtr${type_as_str}${size}[i] = inPtr${type_as_str}${size}[i]; + swizzleInLocalPtr${type_as_str}${size}[i] = swizzleInPtr${type_as_str}${size}[i]; + } + + auto testVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + + std::add_pointer_t<${type}> rawPtrIn${type_as_str}${size} = inLocalPtr${type_as_str}${size}.get_multi_ptr().get_raw(); + const std::add_pointer_t<${type}> constRawPtrIn${type_as_str}${size} = rawPtrIn${type_as_str}${size}; + std::add_pointer_t<${type}> rawPtrOut${type_as_str}${size} = outLocalPtr${type_as_str}${size}.get_multi_ptr().get_raw(); + testVec${type_as_str}${size}.load(0, constRawPtrIn${type_as_str}${size}); + testVec${type_as_str}${size}.store(0, rawPtrOut${type_as_str}${size}); + + auto cleanVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + sycl::vec<${type}, ${size}> swizzledVec {cleanVec${type_as_str}${size}.template swizzle<${swizVals}>()}; + + std::add_pointer_t<${type}> rawPtrInSwizzle${type_as_str}${size} = swizzleInLocalPtr${type_as_str}${size}.get_multi_ptr().get_raw(); + const std::add_pointer_t<${type}> constRawPtrInSwizzle${type_as_str}${size} = rawPtrInSwizzle${type_as_str}${size}; + std::add_pointer_t<${type}> rawPtrOutSwizzle${type_as_str}${size} = swizzleOutLocalPtr${type_as_str}${size}.get_multi_ptr().get_raw(); + swizzledVec.load(0, constRawPtrInSwizzle${type_as_str}${size}); + swizzledVec.store(0, rawPtrOutSwizzle${type_as_str}${size}); + + for (unsigned i = 0; i < ${size}; ++i) { + outPtr${type_as_str}${size}[i] = outLocalPtr${type_as_str}${size}[i]; + swizzleOutPtr${type_as_str}${size}[i] = swizzleOutLocalPtr${type_as_str}${size}[i]; + } + }); + }); + + } + check_array_equality<${type}, ${size}>(log, inputData${type_as_str}${size}, outputData${type_as_str}${size}); + check_array_equality<${type}, ${size}>(log, swizzleInputData${type_as_str}${size}, swizzleOutputData${type_as_str}${size}); + + testQueue.wait_and_throw(); + } + """) + +private_raw_ptr_load_store_test_template = Template( + """ + { + ${type} inputData${type_as_str}${size}[${size}] = {${in_order_vals}}; + ${type} outputData${type_as_str}${size}[${size}] = {${val}}; + ${type} swizzleInputData${type_as_str}${size}[${size}] = {${reverse_order_vals}}; + ${type} swizzleOutputData${type_as_str}${size}[${size}] = {${val}}; + { + sycl::buffer<${type}, 1> inBuffer${type_as_str}${size}(inputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> outBuffer${type_as_str}${size}(outputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleInBuffer${type_as_str}${size}(swizzleInputData${type_as_str}${size}, sycl::range<1>(${size})); + sycl::buffer<${type}, 1> swizzleOutBuffer${type_as_str}${size}(swizzleOutputData${type_as_str}${size}, sycl::range<1>(${size})); + + testQueue.submit([&](sycl::handler &cgh) { + auto inPtr${type_as_str}${size} = inBuffer${type_as_str}${size}.get_access(cgh); + auto outPtr${type_as_str}${size} = outBuffer${type_as_str}${size}.get_access(cgh); + + auto swizzleInPtr${type_as_str}${size} = swizzleInBuffer${type_as_str}${size}.get_access(cgh); + auto swizzleOutPtr${type_as_str}${size} = swizzleOutBuffer${type_as_str}${size}.get_access(cgh); + + cgh.single_task([=]() { + ${type} inPrivatePtr${type_as_str}${size}[${size}]; + ${type} outPrivatePtr${type_as_str}${size}[${size}]; + + ${type} swizzleInPrivatePtr${type_as_str}${size}[${size}]; + ${type} swizzleOutPrivatePtr${type_as_str}${size}[${size}]; + + for (unsigned i = 0; i < ${size}; ++i) { + inPrivatePtr${type_as_str}${size}[i] = inPtr${type_as_str}${size}[i]; + swizzleInPrivatePtr${type_as_str}${size}[i] = swizzleInPtr${type_as_str}${size}[i]; + } + + auto testVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + + std::add_pointer_t<${type}> rawPtrIn${type_as_str}${size} = inPrivatePtr${type_as_str}${size}; + const std::add_pointer_t<${type}> constRawPtrIn${type_as_str}${size} = rawPtrIn${type_as_str}${size}; + std::add_pointer_t<${type}> rawPtrOut${type_as_str}${size} = outPrivatePtr${type_as_str}${size}; + testVec${type_as_str}${size}.load(0, constRawPtrIn${type_as_str}${size}); + testVec${type_as_str}${size}.store(0, rawPtrOut${type_as_str}${size}); + + auto cleanVec${type_as_str}${size} = sycl::vec<${type}, ${size}>(${val}); + sycl::vec<${type}, ${size}> swizzledVec {cleanVec${type_as_str}${size}.template swizzle<${swizVals}>()}; + + std::add_pointer_t<${type}> rawPtrInSwizzle${type_as_str}${size} = swizzleInPrivatePtr${type_as_str}${size}; + const std::add_pointer_t<${type}> constRawPtrInSwizzle${type_as_str}${size} = rawPtrInSwizzle${type_as_str}${size}; + std::add_pointer_t<${type}> rawPtrOutSwizzle${type_as_str}${size} = swizzleOutPrivatePtr${type_as_str}${size}; + swizzledVec.load(0, constRawPtrInSwizzle${type_as_str}${size}); + swizzledVec.store(0, rawPtrOutSwizzle${type_as_str}${size}); + + for (unsigned i = 0; i < ${size}; ++i) { + outPtr${type_as_str}${size}[i] = outPrivatePtr${type_as_str}${size}[i]; + swizzleOutPtr${type_as_str}${size}[i] = swizzleOutPrivatePtr${type_as_str}${size}[i]; + } + }); + }); + + } + check_array_equality<${type}, ${size}>(log, inputData${type_as_str}${size}, outputData${type_as_str}${size}); + check_array_equality<${type}, ${size}>(log, swizzleInputData${type_as_str}${size}, swizzleOutputData${type_as_str}${size}); + + testQueue.wait_and_throw(); + } + """) + def gen_kernel_name(type_str, size, decorated): return 'KERNEL_load_store_' + remove_namespaces_whitespaces(type_str) + str(size) + decorated @@ -93,7 +403,48 @@ def wrap_with_deprecated(test_string): def gen_load_store_test(type_str, size): no_whitespace_type_str = remove_namespaces_whitespaces(type_str) - test_string = load_store_test_template.substitute( + test_string = '' + for decoration in ['yes', 'no', 'legacy']: + multi_ptr_test_string = global_multi_ptr_load_store_test_template.substitute( + type=type_str, + type_as_str=no_whitespace_type_str, + size=size, + val=Data.value_default_dict[type_str], + in_order_vals=', '.join( + append_fp_postfix(type_str, Data.vals_list_dict[size])), + reverse_order_vals=', '.join( + append_fp_postfix(type_str, Data.vals_list_dict[size][::-1])), + kernelName=gen_kernel_name(type_str, size, decoration + '_global'), + swizVals=', '.join(Data.swizzle_elem_list_dict[size]), + decorated=('sycl::access::decorated::' + decoration)) + multi_ptr_test_string += local_multi_ptr_load_store_test_template.substitute( + type=type_str, + type_as_str=no_whitespace_type_str, + size=size, + val=Data.value_default_dict[type_str], + in_order_vals=', '.join( + append_fp_postfix(type_str, Data.vals_list_dict[size])), + reverse_order_vals=', '.join( + append_fp_postfix(type_str, Data.vals_list_dict[size][::-1])), + kernelName=gen_kernel_name(type_str, size, decoration + '_local'), + swizVals=', '.join(Data.swizzle_elem_list_dict[size]), + decorated=('sycl::access::decorated::' + decoration)) + multi_ptr_test_string += private_multi_ptr_load_store_test_template.substitute( + type=type_str, + type_as_str=no_whitespace_type_str, + size=size, + val=Data.value_default_dict[type_str], + in_order_vals=', '.join( + append_fp_postfix(type_str, Data.vals_list_dict[size])), + reverse_order_vals=', '.join( + append_fp_postfix(type_str, Data.vals_list_dict[size][::-1])), + kernelName=gen_kernel_name(type_str, size, decoration + '_private'), + swizVals=', '.join(Data.swizzle_elem_list_dict[size]), + decorated=('sycl::access::decorated::' + decoration)) + if decoration == 'legacy': + multi_ptr_test_string = wrap_with_deprecated(multi_ptr_test_string) + test_string += multi_ptr_test_string + test_string += global_raw_ptr_load_store_test_template.substitute( type=type_str, type_as_str=no_whitespace_type_str, size=size, @@ -102,10 +453,9 @@ def gen_load_store_test(type_str, size): append_fp_postfix(type_str, Data.vals_list_dict[size])), reverse_order_vals=', '.join( append_fp_postfix(type_str, Data.vals_list_dict[size][::-1])), - kernelName=gen_kernel_name(type_str, size, 'yes'), - swizVals=', '.join(Data.swizzle_elem_list_dict[size]), - decorated='sycl::access::decorated::yes') - test_string += load_store_test_template.substitute( + kernelName=gen_kernel_name(type_str, size, 'raw_global'), + swizVals=', '.join(Data.swizzle_elem_list_dict[size])) + test_string += local_raw_ptr_load_store_test_template.substitute( type=type_str, type_as_str=no_whitespace_type_str, size=size, @@ -114,10 +464,9 @@ def gen_load_store_test(type_str, size): append_fp_postfix(type_str, Data.vals_list_dict[size])), reverse_order_vals=', '.join( append_fp_postfix(type_str, Data.vals_list_dict[size][::-1])), - kernelName=gen_kernel_name(type_str, size, 'no'), - swizVals=', '.join(Data.swizzle_elem_list_dict[size]), - decorated='sycl::access::decorated::no') - test_string += wrap_with_deprecated(load_store_test_template.substitute( + kernelName=gen_kernel_name(type_str, size, 'raw_local'), + swizVals=', '.join(Data.swizzle_elem_list_dict[size])) + test_string += private_raw_ptr_load_store_test_template.substitute( type=type_str, type_as_str=no_whitespace_type_str, size=size, @@ -126,9 +475,8 @@ def gen_load_store_test(type_str, size): append_fp_postfix(type_str, Data.vals_list_dict[size])), reverse_order_vals=', '.join( append_fp_postfix(type_str, Data.vals_list_dict[size][::-1])), - kernelName=gen_kernel_name(type_str, size, 'legacy'), - swizVals=', '.join(Data.swizzle_elem_list_dict[size]), - decorated='sycl::access::decorated::legacy')) + kernelName=gen_kernel_name(type_str, size, 'raw_private'), + swizVals=', '.join(Data.swizzle_elem_list_dict[size])) return wrap_with_test_func(TEST_NAME, type_str, wrap_with_extension_checks( type_str, test_string), str(size))