From ad62d72eb7babeb158fb485b7a8149740a2bff1f Mon Sep 17 00:00:00 2001 From: Oskar Hubert Weber Date: Thu, 27 Nov 2025 16:25:21 +0100 Subject: [PATCH] test_mutable_cmdlist: add SLM scenario Signed-off-by: Oskar Hubert Weber --- .../kernels/test_mutable_cmdlist.cl | 13 ++- .../kernels/test_mutable_cmdlist.spv | Bin 11140 -> 12712 bytes .../src/test_mutable_cmdlist.cpp | 100 +++++++++++++++++- 3 files changed, 111 insertions(+), 2 deletions(-) diff --git a/conformance_tests/core/test_mutable_cmdlist/kernels/test_mutable_cmdlist.cl b/conformance_tests/core/test_mutable_cmdlist/kernels/test_mutable_cmdlist.cl index 97c79e998..989791685 100644 --- a/conformance_tests/core/test_mutable_cmdlist/kernels/test_mutable_cmdlist.cl +++ b/conformance_tests/core/test_mutable_cmdlist/kernels/test_mutable_cmdlist.cl @@ -1,6 +1,6 @@ /* * - * Copyright (C) 2024 Intel Corporation + * Copyright (C) 2024-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -85,4 +85,15 @@ kernel void testGlobalOffset(global int *globalOffsets) { globalOffsets[1] += get_global_offset(1); globalOffsets[2] += get_global_offset(2); } +} + +kernel void test_slm_mutation(global uint *out, local uint *slm_1, local uint *slm_2, uint value) { + uint gsize = get_global_size(0); + uint gid = get_global_id(0); + uint lid = get_local_id(0); + + slm_1[lid] = lid + value; + slm_2[lid] = -lid + value; + barrier(CLK_LOCAL_MEM_FENCE); + out[gid] = gsize + slm_1[lid] + slm_2[lid]; } \ No newline at end of file diff --git a/conformance_tests/core/test_mutable_cmdlist/kernels/test_mutable_cmdlist.spv b/conformance_tests/core/test_mutable_cmdlist/kernels/test_mutable_cmdlist.spv index 115ceb18d6459ff6ef0ab0f4ce65d0ec3e0f9b2d..0afd20f365b8e7793ed67d3e5a8dbc3fa5b50992 100644 GIT binary patch delta 2323 zcmZWrOK)3682#p6-?(Ywyo~MO#H48;;MfUq(=;)V(m+Xgl@Kb_4We>m7e)9%W4of+ z7$L+CBx`;_cU`clBDz6rSi`E_BEgDoSaj9$DCL~FchcKrr1AGWXC8OHFZt@ijjfpL zOhqN4WI%f4d^fh?q*eFhv@@Z)yXok&N~G4Vb~_1`?cN+ybKPI2)I#@OPA%(IrWx_f8>m+w@v&TC$8Sn*PR9RuCf1gQtBM+lpc4-)qBRb>88{Md>Q?< zJEGnX@{Z~s++Hs#!?MZz*(Y{6*WRwTX7UGr^Jd2}`G6dWLuZ@I=kt3%xp(_JQTYq0 zhZ5lP#Ff7x8|@bKi2MVc_fs&ji2NJaCj$E+}AS?CO&b7c6T=G__T7_x& z6zxu^N$fGfDeD^ukt~@-JsEnOw`}uie^}hl z+_-$XS!p$cEw8|DU*Boy;n(+SVB5!W2DW{D>k<=B-yBKmKitILjo~|PW?g#CL(hWx z&|}x_#O;9`?ktArAc%&PUQ6y~kHK{xt_UVWlV}7FAp0=~Fd5^F?fsb4Bj;F1986h4 z@i2u2M4f3YAgi-@sFRk3^FSP7!K|R)xR8%@cPur{l6eG+<4UkUmJR(j(^0yfN@iI+ z4nr6=l?6nk}=K{2evkUx&WqwASlV8qmlw}#_ zVM%hM^sh*#$82B(1_PGNo}9JHYi8UMvR=c4ELCode%>!{ED4?dgM2Juax!lC_zGFa z{6M@)z7*ClUkm*_DeaWf-=@>6X{a27nD$|@yzK*GY4ZVPAdA=w-{&*Xhaj0Zu7972 zuTDVY2(z#eSQhpGvK)uBEK|_OfMMRGVd0-tjxY;5hGk*5Aj|WRmgNQL)4(uq+OTAy zafDgeTPzFv3R&`ymSq+?8zu8{WKnba`b>P5eTO(+B6|*T1xSlK0euBA=FJ=2g2AyD zS==IY;!Z(Y+!FLpfn(kxafM}Q9N%MFBzqT;FF_XZ*?9LaL%#)CB&y1PE;zz%ft`&E ntRZ$btFx<7=W99xoy@C{MLA-aUxUszLp*WJXWm0mvW)pZb7DgM delta 822 zcmZ9K&rcIk5XWcUZo6o^rS?(T{y>l_#x%7h9+X56-c1^#i8qrX;9O+0vX z^Zo*F9=MPO{|dQr@M^+=Yfqp8`kiMU*|5pWcV@ox>D!s@*Ts)7Lr*mV@|g;hr8`O3 zRmYUPs;W~;wyFU$(duTLY?nEk?AAD!>^HcWjGA0FQ#_psC`T1D?`iX%Bd(eeM`o8b z*KF+vYIW3_JZ1Baiui`@=_$?Aw)P3N8S_Oo`6lv7bExV(YuIaZ1NnWvg}81$c$#lp zf5(e>4z<~!+wTvSHtr4k{!xD0rxPMtS?g|YKr`ggFX-D|jM%5MOWt|PQy4MLdB`0X zT=K?JUOw$T#Wf)c%klc_Ui^GzYiVFFzLsu6rGz|QUvw>~Wa;_^_kW_n@a{z5n-|e{ z)2X~09#lS;_}Kic>z0i80HcLa8SjJ?{e~{!SMtz5I7M5BZ$>rErrCM2Pb`F zmdlxE9o>?Z3Ri$9+Z3(tv|Q10vSr7~hK18m$4wx915Wx%!aXV@l4lRc ak@6jcoP;jm%g9C}#NVLeoncK6asB~=8;E-V diff --git a/conformance_tests/core/test_mutable_cmdlist/src/test_mutable_cmdlist.cpp b/conformance_tests/core/test_mutable_cmdlist/src/test_mutable_cmdlist.cpp index 6b226f975..f353cb016 100644 --- a/conformance_tests/core/test_mutable_cmdlist/src/test_mutable_cmdlist.cpp +++ b/conformance_tests/core/test_mutable_cmdlist/src/test_mutable_cmdlist.cpp @@ -1,6 +1,6 @@ /* * - * Copyright (C) 2024 Intel Corporation + * Copyright (C) 2024-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -199,6 +199,104 @@ LZT_TEST_F( lzt::destroy_function(addKernel); } +class zeMutableCommandListSLMTests + : public zeMutableCommandListTests, + public ::testing::WithParamInterface< + std::tuple> {}; + +LZT_TEST_P( + zeMutableCommandListSLMTests, + GivenMutationOfSLMKernelArgumentsWhenCommandListIsClosedThenArgumentsWereReplaced) { + if (!kernelArgumentsSupport || !groupSizeSupport) { + GTEST_SKIP() << "Not all required extensions are supported"; + } + uint32_t group_size_x = std::get<0>(GetParam()); + uint32_t mutated_group_size_x = std::get<1>(GetParam()); + uint32_t group_count_x = std::get<2>(GetParam()); + + uint32_t global_size = group_size_x * group_count_x; + uint32_t mutated_global_size = mutated_group_size_x * group_count_x; + + uint32_t verify_value = 21u; + uint32_t *output = reinterpret_cast( + lzt::allocate_host_memory(global_size * sizeof(uint32_t))); + std::memset(output, 0, global_size * sizeof(uint32_t)); + + ze_kernel_handle_t slm_kernel = + lzt::create_function(module, "test_slm_mutation"); + + lzt::set_group_size(slm_kernel, group_size_x, 1, 1); + lzt::set_argument_value(slm_kernel, 0, sizeof(void *), &output); + lzt::set_argument_value(slm_kernel, 1, group_size_x * sizeof(uint32_t), + nullptr); + lzt::set_argument_value(slm_kernel, 2, group_size_x * sizeof(uint32_t), + nullptr); + lzt::set_argument_value(slm_kernel, 3, sizeof(uint32_t), &verify_value); + + uint64_t command_id = 0; + commandIdDesc.flags = ZE_MUTABLE_COMMAND_EXP_FLAG_KERNEL_ARGUMENTS | + ZE_MUTABLE_COMMAND_EXP_FLAG_GROUP_SIZE; + EXPECT_ZE_RESULT_SUCCESS(zeCommandListGetNextCommandIdExp( + mutableCmdList, &commandIdDesc, &command_id)); + ze_group_count_t group_count{group_count_x, 1, 1}; + lzt::append_launch_function(mutableCmdList, slm_kernel, &group_count, nullptr, + 0, nullptr); + lzt::close_command_list(mutableCmdList); + lzt::execute_command_lists(queue, 1, &mutableCmdList, nullptr); + lzt::synchronize(queue, std::numeric_limits::max()); + + for (uint32_t i = 0; i < global_size; i++) { + EXPECT_EQ(output[i], global_size + verify_value * 2); + } + + lzt::free_memory(output); + output = reinterpret_cast( + lzt::allocate_host_memory(mutated_global_size * sizeof(uint32_t))); + std::memset(output, 0, mutated_global_size * sizeof(uint32_t)); + + ze_mutable_kernel_argument_exp_desc_t mutate_kernel_slm_arg_2 = { + ZE_STRUCTURE_TYPE_MUTABLE_KERNEL_ARGUMENT_EXP_DESC}; + mutate_kernel_slm_arg_2.commandId = command_id; + mutate_kernel_slm_arg_2.argIndex = 2; + mutate_kernel_slm_arg_2.argSize = mutated_group_size_x * sizeof(uint32_t); + mutate_kernel_slm_arg_2.pArgValue = nullptr; + ze_mutable_kernel_argument_exp_desc_t mutate_kernel_slm_arg_1 = { + ZE_STRUCTURE_TYPE_MUTABLE_KERNEL_ARGUMENT_EXP_DESC}; + mutate_kernel_slm_arg_1.commandId = command_id; + mutate_kernel_slm_arg_1.argIndex = 1; + mutate_kernel_slm_arg_1.argSize = mutated_group_size_x * sizeof(uint32_t); + mutate_kernel_slm_arg_1.pArgValue = nullptr; + mutate_kernel_slm_arg_1.pNext = &mutate_kernel_slm_arg_2; + ze_mutable_group_size_exp_desc_t mutate_group_size_desc = { + ZE_STRUCTURE_TYPE_MUTABLE_GROUP_SIZE_EXP_DESC}; + mutate_group_size_desc.commandId = command_id; + mutate_group_size_desc.groupSizeX = mutated_group_size_x; + mutate_group_size_desc.groupSizeY = 1; + mutate_group_size_desc.groupSizeZ = 1; + mutate_group_size_desc.pNext = &mutate_kernel_slm_arg_1; + mutableCmdDesc.pNext = &mutate_group_size_desc; + + EXPECT_ZE_RESULT_SUCCESS( + zeCommandListUpdateMutableCommandsExp(mutableCmdList, &mutableCmdDesc)); + + lzt::close_command_list(mutableCmdList); + lzt::execute_command_lists(queue, 1, &mutableCmdList, nullptr); + lzt::synchronize(queue, std::numeric_limits::max()); + + for (uint32_t i = 0; i < mutated_global_size; i++) { + EXPECT_EQ(output[i], mutated_global_size + verify_value * 2); + } + + lzt::free_memory(output); + lzt::destroy_function(slm_kernel); +} + +INSTANTIATE_TEST_SUITE_P(MutableCommandListSLMTests, + zeMutableCommandListSLMTests, + ::testing::Combine(::testing::Values(1, 16, 32, 64), + ::testing::Values(1, 16, 32, 64), + ::testing::Values(1, 2))); + LZT_TEST_F( zeMutableCommandListTests, GivenMutationOfGroupCountWhenCommandListIsClosedThenGlobalWorkSizeIsUpdated) {