From 472afd0d6002b85c6c2cc5cc196fef0bb8e5a592 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 4 Mar 2024 15:36:09 -0600 Subject: [PATCH 1/6] Switch DPCTLKernelArgTypes to C++11 types. --- dpctl/enum_types.py | 19 +++++++ .../include/dpctl_sycl_enum_types.h | 28 +++++----- .../source/dpctl_sycl_queue_interface.cpp | 51 +++++++------------ 3 files changed, 49 insertions(+), 49 deletions(-) diff --git a/dpctl/enum_types.py b/dpctl/enum_types.py index 102ae09015..bb8c54b7be 100644 --- a/dpctl/enum_types.py +++ b/dpctl/enum_types.py @@ -113,3 +113,22 @@ class global_mem_cache_type(Enum): none = auto() read_only = auto() read_write = auto() + + +class kernel_arg_type(Enum): + """ + An enumeration of supported kernel argument types in + :func:`dpctl.SyclQueue.submit` + """ + + dpctl_int8 = auto() + dpctl_uint8 = auto() + dpctl_int16 = auto() + dpctl_uint16 = auto() + dpctl_int32 = auto() + dpctl_uint32 = auto() + dpctl_int64 = auto() + dpctl_uint64 = auto() + dpctl_float32 = auto() + dpctl_float64 = auto() + dpctl_void_ptr = auto() diff --git a/libsyclinterface/include/dpctl_sycl_enum_types.h b/libsyclinterface/include/dpctl_sycl_enum_types.h index 2a1da04ee1..4edcce45df 100644 --- a/libsyclinterface/include/dpctl_sycl_enum_types.h +++ b/libsyclinterface/include/dpctl_sycl_enum_types.h @@ -87,22 +87,18 @@ typedef enum */ typedef enum { - DPCTL_CHAR, - DPCTL_SIGNED_CHAR, - DPCTL_UNSIGNED_CHAR, - DPCTL_SHORT, - DPCTL_INT, - DPCTL_UNSIGNED_INT, - DPCTL_UNSIGNED_INT8, - DPCTL_LONG, - DPCTL_UNSIGNED_LONG, - DPCTL_LONG_LONG, - DPCTL_UNSIGNED_LONG_LONG, - DPCTL_SIZE_T, - DPCTL_FLOAT, - DPCTL_DOUBLE, - DPCTL_LONG_DOUBLE, - DPCTL_VOID_PTR + DPCTL_INT8_T, + DPCTL_UINT8_T, + DPCTL_INT16_T, + DPCTL_UINT16_T, + DPCTL_INT32_T, + DPCTL_UINT32_T, + DPCTL_INT64_T, + DPCTL_UINT64_T, + DPCTL_FLOAT32_T, + DPCTL_FLOAT64_T, + DPCTL_VOID_PTR, + DPCTL_UNSUPPORTED_KERNEL_ARG } DPCTLKernelArgType; /*! diff --git a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp index 63c4720ff3..8ffd579b31 100644 --- a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp @@ -65,51 +65,36 @@ bool set_kernel_arg(handler &cgh, bool arg_set = true; switch (ArgTy) { - case DPCTL_CHAR: - cgh.set_arg(idx, *(char *)Arg); + case DPCTL_INT8_T: + cgh.set_arg(idx, *(int8_t *)Arg); break; - case DPCTL_SIGNED_CHAR: - cgh.set_arg(idx, *(signed char *)Arg); - break; - case DPCTL_UNSIGNED_CHAR: - cgh.set_arg(idx, *(unsigned char *)Arg); - break; - case DPCTL_SHORT: - cgh.set_arg(idx, *(short *)Arg); - break; - case DPCTL_INT: - cgh.set_arg(idx, *(int *)Arg); - break; - case DPCTL_UNSIGNED_INT: - cgh.set_arg(idx, *(unsigned int *)Arg); - break; - case DPCTL_UNSIGNED_INT8: + case DPCTL_UINT8_T: cgh.set_arg(idx, *(uint8_t *)Arg); break; - case DPCTL_LONG: - cgh.set_arg(idx, *(long *)Arg); + case DPCTL_INT16_T: + cgh.set_arg(idx, *(int16_t *)Arg); + break; + case DPCTL_UINT16_T: + cgh.set_arg(idx, *(uint16_t *)Arg); break; - case DPCTL_UNSIGNED_LONG: - cgh.set_arg(idx, *(unsigned long *)Arg); + case DPCTL_INT32_T: + cgh.set_arg(idx, *(int32_t *)Arg); break; - case DPCTL_LONG_LONG: - cgh.set_arg(idx, *(long long *)Arg); + case DPCTL_UINT32_T: + cgh.set_arg(idx, *(uint32_t *)Arg); break; - case DPCTL_UNSIGNED_LONG_LONG: - cgh.set_arg(idx, *(unsigned long long *)Arg); + case DPCTL_INT64_T: + cgh.set_arg(idx, *(int64_t *)Arg); break; - case DPCTL_SIZE_T: - cgh.set_arg(idx, *(size_t *)Arg); + case DPCTL_UINT64_T: + cgh.set_arg(idx, *(uint64_t *)Arg); break; - case DPCTL_FLOAT: + case DPCTL_FLOAT32_T: cgh.set_arg(idx, *(float *)Arg); break; - case DPCTL_DOUBLE: + case DPCTL_FLOAT64_T: cgh.set_arg(idx, *(double *)Arg); break; - case DPCTL_LONG_DOUBLE: - cgh.set_arg(idx, *(long double *)Arg); - break; case DPCTL_VOID_PTR: cgh.set_arg(idx, Arg); break; From e5c055760c694cc59dbaf276ac9e7021e2fa73f3 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 4 Mar 2024 15:40:35 -0600 Subject: [PATCH 2/6] Refactor SubmitRange/SubmitNdRange to fix warning. --- .../source/dpctl_sycl_queue_interface.cpp | 159 ++++++++++++------ 1 file changed, 104 insertions(+), 55 deletions(-) diff --git a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp index 8ffd579b31..dc512fd126 100644 --- a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp @@ -51,6 +51,17 @@ typedef struct complex uint64_t imag; } complexNumber; +void set_dependent_events(handler &cgh, + __dpctl_keep const DPCTLSyclEventRef *DepEvents, + size_t NDepEvents) +{ + for (auto i = 0ul; i < NDepEvents; ++i) { + auto ei = unwrap(DepEvents[i]); + if (ei) + cgh.depends_on(*ei); + } +} + /*! * @brief Set the kernel arg object * @@ -107,6 +118,21 @@ bool set_kernel_arg(handler &cgh, return arg_set; } +void set_kernel_args(handler &cgh, + __dpctl_keep void **Args, + __dpctl_keep const DPCTLKernelArgType *ArgTypes, + size_t NArgs) +{ + for (auto i = 0ul; i < NArgs; ++i) { + if (!set_kernel_arg(cgh, i, Args[i], ArgTypes[i])) { + error_handler("Kernel argument could not be created.", __FILE__, + __func__, __LINE__); + throw std::invalid_argument( + "Kernel argument could not be created."); + } + } +} + std::unique_ptr create_property_list(int properties) { std::unique_ptr propList; @@ -341,39 +367,52 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, event e; try { - e = Queue->submit([&](handler &cgh) { - // Depend on any event that was specified by the caller. - if (NDepEvents) - for (auto i = 0ul; i < NDepEvents; ++i) - cgh.depends_on(*unwrap(DepEvents[i])); - - for (auto i = 0ul; i < NArgs; ++i) { - // \todo add support for Sycl buffers - if (!set_kernel_arg(cgh, i, Args[i], ArgTypes[i])) - exit(1); - } - switch (NDims) { - case 1: + switch (NDims) { + case 1: + { + e = Queue->submit([&](handler &cgh) { + // Depend on any event that was specified by the caller. + set_dependent_events(cgh, DepEvents, NDepEvents); + set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(range<1>{Range[0]}, *Kernel); - break; - case 2: + }); + return wrap(new event(std::move(e))); + } + case 2: + { + e = Queue->submit([&](handler &cgh) { + // Depend on any event that was specified by the caller. + set_dependent_events(cgh, DepEvents, NDepEvents); + set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(range<2>{Range[0], Range[1]}, *Kernel); - break; - case 3: + }); + return wrap(new event(std::move(e))); + } + case 3: + { + e = Queue->submit([&](handler &cgh) { + // Depend on any event that was specified by the caller. + set_dependent_events(cgh, DepEvents, NDepEvents); + set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(range<3>{Range[0], Range[1], Range[2]}, *Kernel); - break; - default: - throw std::runtime_error("Range cannot be greater than three " - "dimensions."); - } - }); + }); + return wrap(new event(std::move(e))); + } + default: + error_handler("Range cannot be greater than three " + "dimensions.", + __FILE__, __func__, __LINE__, error_level::error); + return nullptr; + } } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__); + error_handler(e, __FILE__, __func__, __LINE__, error_level::error); + return nullptr; + } catch (...) { + error_handler("Unknown exception encountered", __FILE__, __func__, + __LINE__, error_level::error); return nullptr; } - - return wrap(new event(std::move(e))); } __dpctl_give DPCTLSyclEventRef @@ -393,46 +432,56 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, event e; try { - e = Queue->submit([&](handler &cgh) { - // Depend on any event that was specified by the caller. - if (DepEvents) - for (auto i = 0ul; i < NDepEvents; ++i) { - auto ei = unwrap(DepEvents[i]); - if (ei) - cgh.depends_on(*ei); - } - - for (auto i = 0ul; i < NArgs; ++i) { - // \todo add support for Sycl buffers - if (!set_kernel_arg(cgh, i, Args[i], ArgTypes[i])) - exit(1); - } - switch (NDims) { - case 1: + switch (NDims) { + case 1: + { + e = Queue->submit([&](handler &cgh) { + // Depend on any event that was specified by the caller. + set_dependent_events(cgh, DepEvents, NDepEvents); + set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(nd_range<1>{{gRange[0]}, {lRange[0]}}, *Kernel); - break; - case 2: + }); + return wrap(new event(std::move(e))); + } + case 2: + { + e = Queue->submit([&](handler &cgh) { + // Depend on any event that was specified by the caller. + set_dependent_events(cgh, DepEvents, NDepEvents); + set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for( nd_range<2>{{gRange[0], gRange[1]}, {lRange[0], lRange[1]}}, *Kernel); - break; - case 3: + }); + return wrap(new event(std::move(e))); + } + case 3: + { + e = Queue->submit([&](handler &cgh) { + // Depend on any event that was specified by the caller. + set_dependent_events(cgh, DepEvents, NDepEvents); + set_kernel_args(cgh, Args, ArgTypes, NArgs); cgh.parallel_for(nd_range<3>{{gRange[0], gRange[1], gRange[2]}, {lRange[0], lRange[1], lRange[2]}}, *Kernel); - break; - default: - throw std::runtime_error("Range cannot be greater than three " - "dimensions."); - } - }); + }); + return wrap(new event(std::move(e))); + } + default: + error_handler("Range cannot be greater than three " + "dimensions.", + __FILE__, __func__, __LINE__, error_level::error); + return nullptr; + } } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__); + error_handler(e, __FILE__, __func__, __LINE__, error_level::error); + return nullptr; + } catch (...) { + error_handler("Unknown exception encountered", __FILE__, __func__, + __LINE__, error_level::error); return nullptr; } - - return wrap(new event(std::move(e))); } void DPCTLQueue_Wait(__dpctl_keep DPCTLSyclQueueRef QRef) From fe43eb8d812df17b4968a213a0b900863d39cbb2 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 4 Mar 2024 23:58:49 -0600 Subject: [PATCH 3/6] Unit tests for one-D kernel submission --- libsyclinterface/tests/CMakeLists.txt | 123 +++--- .../tests/oneD_range_kernel_fp64.spv | Bin 0 -> 5180 bytes .../tests/oneD_range_kernel_inttys_fp32.spv | Bin 0 -> 17808 bytes .../tests/test_sycl_queue_submit.cpp | 402 ++++++++++++------ 4 files changed, 346 insertions(+), 179 deletions(-) create mode 100644 libsyclinterface/tests/oneD_range_kernel_fp64.spv create mode 100644 libsyclinterface/tests/oneD_range_kernel_inttys_fp32.spv diff --git a/libsyclinterface/tests/CMakeLists.txt b/libsyclinterface/tests/CMakeLists.txt index 13f36c39d7..a0be739b2e 100644 --- a/libsyclinterface/tests/CMakeLists.txt +++ b/libsyclinterface/tests/CMakeLists.txt @@ -1,4 +1,5 @@ find_package(GTest REQUIRED) + # We need thread support for gtest find_package(Threads REQUIRED) @@ -16,7 +17,12 @@ include_directories( link_directories(${GTEST_LIB_DIR}) # Copy the spir-v input files to test build directory -set(spirv-test-files multi_kernel.spv) +set(spirv-test-files + multi_kernel.spv + oneD_range_kernel_inttys_fp32.spv + oneD_range_kernel_fp64.spv +) + foreach(tf ${spirv-test-files}) file(COPY ${tf} DESTINATION ${CMAKE_CURRENT_BINARY_DIR}) endforeach() @@ -32,38 +38,39 @@ add_executable(dpctl_c_api_tests ${sources} ) add_sycl_to_target( - TARGET dpctl_c_api_tests - SOURCES - ${CMAKE_CURRENT_SOURCE_DIR}/test_helper.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_context_interface.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_invalid_filters.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_subdevices.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_interface.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_manager.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_selector_interface.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_aspects.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_event_interface.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_platform_interface.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_kernel_interface.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_kernel_bundle_interface.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_platform_invalid_filters.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_manager.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_submit.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_interface.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_usm_interface.cpp -) -if (_dpctl_sycl_targets) -# make fat binary -target_compile_options( - dpctl_c_api_tests - PRIVATE - -fsycl-targets=${_dpctl_sycl_targets} -) -target_link_options( - dpctl_c_api_tests - PRIVATE - -fsycl-targets=${_dpctl_sycl_targets} + TARGET dpctl_c_api_tests + SOURCES + ${CMAKE_CURRENT_SOURCE_DIR}/test_helper.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_context_interface.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_invalid_filters.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_subdevices.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_interface.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_manager.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_selector_interface.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_aspects.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_event_interface.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_platform_interface.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_kernel_interface.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_kernel_bundle_interface.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_platform_invalid_filters.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_manager.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_submit.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_queue_interface.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_usm_interface.cpp ) + +if(_dpctl_sycl_targets) + # make fat binary + target_compile_options( + dpctl_c_api_tests + PRIVATE + -fsycl-targets=${_dpctl_sycl_targets} + ) + target_link_options( + dpctl_c_api_tests + PRIVATE + -fsycl-targets=${_dpctl_sycl_targets} + ) endif() if(DPCTL_GENERATE_COVERAGE) @@ -82,21 +89,21 @@ if(DPCTL_GENERATE_COVERAGE) add_custom_target(llvm-cov COMMAND ${CMAKE_COMMAND} -E env DPCTL_VERBOSITY=warning ${CMAKE_CURRENT_BINARY_DIR}/dpctl_c_api_tests COMMAND ${LLVMProfdata_EXE} - merge - -sparse default.profraw - -o - dpctl.profdata + merge + -sparse default.profraw + -o + dpctl.profdata COMMAND ${LLVMCov_EXE} - export - -format=lcov - -ignore-filename-regex=/tmp/icpx* - -instr-profile=dpctl.profdata - "${object_arg}$,;${object_arg}>" - > dpctl.lcov + export + -format=lcov + -ignore-filename-regex=/tmp/icpx* + -instr-profile=dpctl.profdata + "${object_arg}$,;${object_arg}>" + > dpctl.lcov COMMAND ${LLVMCov_EXE} - report - -instr-profile=dpctl.profdata - "${object_arg}$,;${object_arg}>" + report + -instr-profile=dpctl.profdata + "${object_arg}$,;${object_arg}>" WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} COMMAND_EXPAND_LISTS DEPENDS dpctl_c_api_tests @@ -105,21 +112,21 @@ if(DPCTL_GENERATE_COVERAGE) add_custom_target(lcov-genhtml COMMAND ${CMAKE_COMMAND} -E env DPCTL_VERBOSITY=warning ${CMAKE_CURRENT_BINARY_DIR}/dpctl_c_api_tests COMMAND ${LLVMProfdata_EXE} - merge - -sparse default.profraw - -o - dpctl.profdata + merge + -sparse default.profraw + -o + dpctl.profdata COMMAND ${LLVMCov_EXE} - export - -format=lcov - -instr-profile=dpctl.profdata - "${object_arg}$,;${object_arg}>" - > dpctl.lcov + export + -format=lcov + -instr-profile=dpctl.profdata + "${object_arg}$,;${object_arg}>" + > dpctl.lcov COMMAND ${GENHTML_EXE} - ${CMAKE_CURRENT_BINARY_DIR}/dpctl.lcov - --no-source - --output-directory - ${COVERAGE_OUTPUT_DIR}/dpctl-c-api-coverage + ${CMAKE_CURRENT_BINARY_DIR}/dpctl.lcov + --no-source + --output-directory + ${COVERAGE_OUTPUT_DIR}/dpctl-c-api-coverage WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} COMMAND_EXPAND_LISTS DEPENDS dpctl_c_api_tests diff --git a/libsyclinterface/tests/oneD_range_kernel_fp64.spv b/libsyclinterface/tests/oneD_range_kernel_fp64.spv new file mode 100644 index 0000000000000000000000000000000000000000..32f5438ff18618badfaa170ab3d80f4dbc8d514e GIT binary patch literal 5180 zcmai#d3Tdl5XNt5*h&Fe3WA6!D4-}(K$apOfdDL@IpWmB1%}b2>=FFLyXO=s6?%bQ&F$>4LSdE+F#<`zd>6_-pxR^xv zSkK3KUg!BF&lh>V*z*SsrMO>bJ6ETiJ-hY=9Xs~4b_Lm7F_XPL_6T&z*WRIPe~LN1jb3N{a>(#4KUdpg^jNOxq0vVDnS zDx2xhR^(U8UE#Ywn?HIupB>EMG*%s&@NfCLQki5T?{n2D+dtMF&DE1SMN5@$f+}>b zFReB&qhHOcd&8l6$mf zn2T(K}$x$ao#w@5bbG5sR0ymrj#xo}ZCdf8qd_PJEsf6tS`!n0a0 zJ~*u`4|#c01V=96F}FqX5)t!FeOn*W>s_1ls_R$GsocW&?c(M7W3ID8?|MkDb=}gf zUSz9xZv;oSdJjmB)_br*=dFpZuWKI8o4WT`yH1L`RCPa*`y`ieFPFrlxh_j*21>c= z9e1v3u8Mmnw=myz@o2sq75Xi&$Grc&3j6(t-k`H$rdmXN=#ApcG#~nD@i>aUrMjv3 zPZqHUKEvpz#TiDO%mkld%OyV=B4un%Ngsik-$u!-#b;PTa&3q>x5C%T+3+p(`anq? zjY)YOW-*)r-)zN@k9FX2>GX~F$n^1E$7hB-`468buPNfzi}~`dmkx%^T}CJ0RB>{E z&5})5$+d{EO&4b^u=#qe=oOd|IDRw4$$`C55i`Z9gWT8}mG?1^;Xhw8{?v!Q!0U}( ze^Q*YviwhZJ2?I;ypDaQvaA(vRK{)MtZVhb@u45R;;fetk)uzX8o<$yinAs@joRxO zadLr!F=y=Uig*+!_L?wq;fsFL>!sO@x!cmoZT7of$G%?vcBb^&^l_3+CpYn?<3G{s ztA$cegNTpbuTj70J?}6wXTY2Ln|8u%uhWXeJT96iIxU?uc|t_Y2JN3c$3sBiyiOcvwAf=7*QOeY;~gP(iO6D~@wmkm|wb-EyI~h%A@&XtOZZG5nm?rM4|1Y%QV-9@{FJo?^3gS|xKvwhlUemhX2- z)(^PNX}jb$(F0ebKG!I|H1o`Ejfl5}9=;*Y{pl0$7jbWXmM>>>Sd|AZ$J`qL@tT9<+ z=lZE+;%(mW1FMYxOqj(-=k0Uh#N*HU_<>ox&D$5kqVu*+yWz~f7F`vwUwk+(>k;Qg zPw6MT^1M2P@of|FKR8`{dqnPx$X#BZCjPRBoIAwr4caN0_s06YOY&F|`-qFIU%Mq^ z7kii{ z{S^`G7;f+R0dFJ5-t&W!!Rd|tr@bnfHP(puzeBegUXx5utp;>**jm=t*M+r-D$kC4 zI!04>Mx3)dDn8<6a-8!r_xv3Z=k<*^_wkzebrJUwop<|2M87Hdme;vYwG;?_ew2(4HDV)wv_|Z#W0=+WH<1M$%K!iX literal 0 HcmV?d00001 diff --git a/libsyclinterface/tests/oneD_range_kernel_inttys_fp32.spv b/libsyclinterface/tests/oneD_range_kernel_inttys_fp32.spv new file mode 100644 index 0000000000000000000000000000000000000000..08cef17e0492086c1430e2c2743d2a34266d1fb4 GIT binary patch literal 17808 zcmcJWcbr{S)rL=)2_*y~T@>Pk7$72n2@puAf*}M#OAr)+aY!aGFqsLHNkWhku!3C? zyQ0`D_FhoH-h1!8cj@2r+`HeIGso-aKXQJv*lWFey=$+%*E;*&BtzF0v$v|cr&QZl zTUJ+ZQMH~Osx7K+G{>oB+_H?@mT}iI?p?-x%6QK%ZQd7ltEy?lTyff&eM?U{WAU=S zk+oyPBf|suO{=zQ;U}H2yl>#rwF4{1`uf+cTffG+ZCP;&s|LpU2ZwqOTRyUWc-6qF<^97K44gDD zIy^A6^pq6`_8mMpHn3*tvccZP%g*mzym+x%Z*SeJAC$1wjB1mR#F;yy$+&V_hSUk9W*t-K0FeaX;iS*2v?Ij(Lnt%427m=zrZG zgN;0nENe{ekHJZKoPIy#aZw|Wv5t9MG%1g#-4A&THS&0S$2^86<#Fr%kjI)v9>3_A z$C^ob+T-O8hjJIy=x#wSnV#`D(0w{P9r!O;zU$F3h78e2NN zWN764{-LGA8%9?4j}49tFXbx4uYET)e9sygz4(IBk@ahJn#v8k^?!=DY;bs>f3(D! zMznuhcRbdL!OI4iCf?R$SXC>BxTfcM-!59jg}Z3cs8?)zz3H9i+c)x}Mg60r{hQi; z4;c6B>$_lJcwn@DY+zMi?G|pWH?5k<*w;5WHr6+?diBsq|Ej)?gMI79`bWq5HjehM zT{|$^j^|{DZ=S$ar}H= zLWjz?_IuRqhokqj>}^}0lO5uZ_mGX2xPERYuDWp_)QZ#HiZc;+daHjs#BJwY>*2(< z-xFKi*n5@O+QxQVQcZ95e23WMIj8RvdmDI(t#0hyTd}%Zu_p5Fs_gG+W$uBDem|mA zt+Ujg$fLUzqdsrUiTUZRcc7i;|NkC9YwgQ-p`(A}eQ5i4S4&Y`uf5Oi>cob9Qo~+` zZNBk8x7gXA_C)^O)p?D1?fU#x#ZJ7<4g2bbeGRtj``D+x(>d!Gqu0+;yS~;-8uQxq z`A-x(d%3k?-`22i$F85nH;WyC2UDf7s zy}NokdObhwiS@gyXEf%uC+54V*J5Yi>dEyDHC#RTn``)6#`W&%9q1EryQ_CL>~~GF zyQ&Y2bK832f2fA52mfdd|Jb_h(tYUly=m9i_HwV!YfsE~RS&|>+Ui+j z&l;|tz0E@3m-ooe5#Ob~8umWe#!t`g+py6PDcI>u^z&tdsiy!DNE z8yfb;V(09xXv|+(vv=WyF{@(?&ZEuM*ZK|-`*|qV$;bKaIf$N?RaN`(%}M@ti09)I zY$x=sTF3j-IISPlaR+OrmU-Gk$+CJ_4j~ln! zoa2kFPdN6A&u+vpo~C0Dwt4#9HMZxitFt%ba*4LcL8|I%8eWhn*aTq^Ipo z_vUAeXLdM&d_Rt1^!uIf#T&`TT09cl&t^Z5J#XzJks}bfj_KdU=<=&6U6r~$e?0Z? zVO&z?{2W&{5zF`V1a#-pCriJEq?T(B<22>h*qG zUFMqob^*9^IjalNJN4TjxN|<%TaK7ez27bZFbCJE_uIwT^4+(5{lxg0#2*3(zus?a zz~$>NUq3PJdcO?=)cb8AyRl~L!2R=+`_;#K<&3PCXX<%6HtTf)e4izVpL07hF0JV& z*7TE#?)|z9G3LpPxr0tY_rA$_pNih;bKo>^=X`R1pN=k``*JzDSnhB0T7j6Gex8%I zer&6C_Tip^EjRbK+(#kqC0uR&!Zm02btYom;E8!wnUj0`8?D1 zZa0Us5zjC=Xd5GIF+t{8xYUc_3UqL{mkFHCB3*2U2ghY?xl$58m_i};pR*> zff+Y=VqOO3oSgLcrKzXnN9J`3G7 z(y!XNf1Zu*872p9V`QzI<8#0|)#^pyJ}*S9&kV*F7ySi{FF~xyOBoaE73kI^_s@;! za@?=$z7n|+QLo=WufjH-yjL^6rr7dsLVs<+!pr^hI&isjk#Jv+E*|b1(8ZeEH-gFC z9|`wO=;GnN8C|T&eG8b}UPSI&8Q+G8%Y7c>%>@fr-23S5$W%nU2T7m41Kl~F)c2j} z;;FB8=H6B2JX7x@?Ux}hN34sw{cbMrMvP;BX}`B%?_u;TJ>&Nwp0j@TyS9Gjsf7Ch zbmP4r3HO8Oo^!a``h}Y_`4E_KgD2*P!JLz0znkMn5OL2j>wFa5nq(bq{SwdnDEoIz zUw#a|Q;&Zf+~*d=9>1RP6E*$Tntof+&Fyx?GyWuF*7y|qOeA&wGjf?k+pJ;UkB?{t8ak&)bFEjVp})wqi-SBI_&9c)1sTRBU;^D?diW!gG(av44Wdor{F~(}IWl zGina(@LTw-=G?ee`QYT&_L)n}UTa-UaqsWGW)wgXBK? z9lCQqsqgR6#ZzDH%>ALvd8Xb++TTIGi&z)+^!p#tjg$WR6S`RX{U-E3Bc8K<_Pe%z z=Bb4H7j(Jlce#H>Jm+w=^$Ry=@;5N!22aetgE=QB{r(ShanCXL(Ld3xN!HQUFY&yO zvVX_)<-gE7_4pm&KDQvbkM6ALch&T}i*9Z2K|JGo8MDTH==OfD3ZftYY&-!S${J`q_tCcKo?&v(Hb&OUIZg-bR4acE)yIEl$oHdNitfKR z_`4U@WH-jddJwub@g419d@!~g_v^ZQpo^*3@1s4zj3>{3v*;8u{eY6+2 z+}uZUXCva_KD1y>?i_HrxsT-TjfjW4Pr;hpeZl4CKAMaEFhpFge;2i1!NL{yK6*HK zr#{^u+&Q1rcOJHQ>Z_f(1InCd>V2gBKxB8sx~SXl=5ion9Q#Xqe!=`VP|wmcJ_zxg z^|Rl#^)vr&B-}&5jOV|XhPwdWa}HNqzi@LVhk_Y5cw+kdInK$k-_3C$;u$)Zbrzvp zldPkyU*dTmW&e)p%frz-_4twCK1VRNdYtj7ntpUmKc?vB_6Wo?K9(_S9EV=Nj~i_yc`QK}OMU#exf2j`)6e>7>zBQUyA(`r>Ld3=#Jz;8tzWpQ&q-j$ z4W5|Gz?_ql`kag|?pdTKPeFH$^sBbDGlx?V&oDV?8zXDw98c4SvuXa>`gCxg`hBz< z%({6Wtw5~F8H|Z_7JB_YdNj5i_v^ZkK^If6-$!SI887$IV~Z`%cjX*JEIju(n;7SU z$*tc<=V6P7+lMaJF zUI6Z#PwIOiws`8Row>m>=b3sRX+H`%6R|Gp>GzA!jg$Vm7+oy=?imjup0j@TyS9Gj zU%!vmfXPk2%N<5M=Ww<43pZyn0%qLc$!9H?b8^z}m!Nxw&gDKDMYkqdM_a$d^FGS{ z9n+WV&^z_`7`RXUK3Wf!yJZ(__Em!a41qsL=+s`KUG&iSN1S73|h zJg!6+OMPabUxk>Pe%41@zwABS&0umAp{lZOst^qS{@Wgxqm~(PcpC_V= zdlu=*C!xDW`c*sk(UZ|V!{nfCjI5P&dI#fW&gFDY1) z`%-YZxsT+&3=t3aV2gBEacgUby2sUT>rJ`#@P!4vb%V9v?0zny;zx@YEG)_E(sHOM;J`X!$CPxkAWK71Q`r=EU0 zxX;au_ViZh@2KhTtm*G6x@)}~-7|d;W7c>tdj0-+A9km@z8~B&d3+Gt zT*R%1wtm@bxF0G$sfXMTgPD7{+WLi?dVB=TxWNF7DpbbGM+oM*30P zTA9Pg(LJ-|plyt-m2>+9Sf|?D3hqSkt=~6a!xj(s z>*!)l?l-{X*6*8dVvA?5-$EA)SKRyN+hCpg@;l(p`J}$z#THL}wKMmkIT+cW(E;<<|3-`e^m-*A6eeA3@?e}s6h;cDv_ZqDGxV8#uem_Gq? zPEPv!r|6!UbGdJRhHeeAj<$Y@=Y5m?I;IbQj^3%Ke*xY)U&h=wzpUxMs_DNjx@-Lg z-820yW7hZ`di}onJ$9$M{sG)MpX~pS*y5>&@9Uqi%|+aLXzQ1~hWqE@lX}Sg3z)fw ztF2$SsmEWzj2k>L{|4rqoYdp*=;H1@J@*fE*GNBV=f3$Tx@VRgw2hIqa&G?u>r|V2 z!F}#w%=er7ihj3q1hNKG7!#`-y?)!(Ziy|%&!q1+TL~cEH`_2yE4I9?(YGy_ zu-}Vv-%JOPoBKxYc8GYm+ZT*s5_blG+}t;EAApF5yF zd*-j>S{Or%EAD->D?q2d+zs40pVaq(*y5?LcII|3bDpX9jrJBu7h+x1(@*~Y8^*D} z*Wzn_55^YvEIs2r5YJgZ`(0Zfeqc<*+7mz@`#rH{Vv8r%L(o0<#M0J>pH{4#Q4fIY z1izYf&21Kdb8?5>=y?!6&K?ySK#B)`y=-^*RK=b%Li}3jmyxlX@MBF7DZ+XAeVn zjr6m2a#)D&nI;EqWAM|;A?LaXpi}MqJwI-v*6&?gJ3kj63#M+2bBew___^rTYmCv~ z=eq%HJ-T9_c+3sJzi3DIivi?eO=C&eO-ZG-`ADJ&OWZf)+hVejINJ4I%j-- N#`AXGb%G_w{{uM`Vj2Jd literal 0 HcmV?d00001 diff --git a/libsyclinterface/tests/test_sycl_queue_submit.cpp b/libsyclinterface/tests/test_sycl_queue_submit.cpp index 34f6f71099..1e8f2ae4b3 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit.cpp @@ -35,7 +35,9 @@ #include #include #include +#include #include +#include namespace { @@ -44,154 +46,312 @@ static_assert(SIZE % 8 == 0); using namespace dpctl::syclinterface; +template +void submit_kernel(DPCTLSyclQueueRef QRef, + DPCTLSyclKernelBundleRef KBRef, + std::vector spirvBuffer, + size_t spirvFileSize, + DPCTLKernelArgType kernelArgTy, + std::string kernelName) +{ + T scalarVal = 3; + constexpr size_t NARGS = 4; + constexpr size_t RANGE_NDIMS = 1; + + ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, kernelName.c_str())); + auto kernel = DPCTLKernelBundle_GetKernel(KBRef, kernelName.c_str()); + + // Create the input args + auto a = DPCTLmalloc_shared(SIZE * sizeof(T), QRef); + ASSERT_TRUE(a != nullptr); + auto b = DPCTLmalloc_shared(SIZE * sizeof(T), QRef); + ASSERT_TRUE(b != nullptr); + auto c = DPCTLmalloc_shared(SIZE * sizeof(T), QRef); + ASSERT_TRUE(c != nullptr); + + // Create kernel args for vector_add + size_t Range[] = {SIZE}; + void *args[NARGS] = {unwrap(a), unwrap(b), unwrap(c), + (void *)&scalarVal}; + DPCTLKernelArgType addKernelArgTypes[] = {DPCTL_VOID_PTR, DPCTL_VOID_PTR, + DPCTL_VOID_PTR, kernelArgTy}; + auto ERef = DPCTLQueue_SubmitRange(kernel, QRef, args, addKernelArgTypes, + NARGS, Range, RANGE_NDIMS, nullptr, 0); + ASSERT_TRUE(ERef != nullptr); + DPCTLQueue_Wait(QRef); + + // clean ups + DPCTLEvent_Delete(ERef); + DPCTLKernel_Delete(kernel); + DPCTLfree_with_queue((DPCTLSyclUSMRef)a, QRef); + DPCTLfree_with_queue((DPCTLSyclUSMRef)b, QRef); + DPCTLfree_with_queue((DPCTLSyclUSMRef)c, QRef); +} + } /* end of anonymous namespace */ +/* +// The oneD_range_kernel spv files were generated from the below SYCL program. +// To compile: +// icpx -fsycl oneD_range_kernel.cpp +// IGC_ShaderDumpEnable=1 IGC_DumpToCustomDir=dump ./a.out + +// The generated spv files should be inspected using spirv-dis to identify +// kernel names. When these files were generated using dpcpp 2024.1, a single +// spv file was generated for all integer types and float32_t and a separate +// file generated for float64_t. + +#include +#include +#include + +template +class Range1DKernel +{ +private: + T *a_ = nullptr; + T *b_ = nullptr; + T *c_ = nullptr; + T scalarVal_; + +public: + RangeKernel(T *a, T *b, T *c, T scalarVal) + : a_(a), b_(b), c_(c), scalarVal_(scalarVal) + { + } + + void operator()(sycl::item<1> it) const + { + auto i = it.get_id(); + a_[i] = i + 1; + b_[i] = i + 2; + c_[i] = scalarVal_ * (a_[i] + b_[i]); + } +}; + +template +void submit_kernel( + sycl::queue q, + const unsigned long N, + T *a, + T *b, + T *c, + T scalarVal) +{ + // clang-format off + q.submit([&](auto &h) { + h.parallel_for(sycl::range(N), RangeKernel(a, b, c, scalarVal)); + }); + // clang-format on +} + +template +void driver(size_t N) +{ + sycl::queue q; + auto *a = sycl::malloc_shared(N, q); + auto *b = sycl::malloc_shared(N, q); + auto *c = sycl::malloc_shared(N, q); + T scalarVal = 3; + + submit_kernel(q, N, a, b, c, scalarVal); + q.wait(); + + std::cout << "C[0] : " << (size_t)c[0] << " " << std::endl; + sycl::free(a, q); +} + +int main(int argc, const char **argv) +{ + size_t N = 0; + std::cout << "Enter problem size in N:\n"; + std::cin >> N; + std::cout << "Executing with N = " << N << std::endl; + + driver(N); + driver(N); + driver(N); + driver(N); + driver(N); + driver(N); + driver(N); + driver(N); + driver(N); + driver(N); + + return 0; +} +*/ + struct TestQueueSubmit : public ::testing::Test { std::ifstream spirvFile; - size_t spirvFileSize; - std::vector spirvBuffer; + size_t spirvFileSize_; + std::vector spirvBuffer_; + DPCTLSyclQueueRef QRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; TestQueueSubmit() { - spirvFile.open("./multi_kernel.spv", std::ios::binary | std::ios::ate); - spirvFileSize = std::filesystem::file_size("./multi_kernel.spv"); - spirvBuffer.reserve(spirvFileSize); + DPCTLSyclDeviceSelectorRef DSRef = nullptr; + DPCTLSyclDeviceRef DRef = nullptr; + + spirvFile.open("./oneD_range_kernel_inttys_fp32.spv", + std::ios::binary | std::ios::ate); + spirvFileSize_ = + std::filesystem::file_size("./oneD_range_kernel_inttys_fp32.spv"); + spirvBuffer_.reserve(spirvFileSize_); spirvFile.seekg(0, std::ios::beg); - spirvFile.read(spirvBuffer.data(), spirvFileSize); + spirvFile.read(spirvBuffer_.data(), spirvFileSize_); + + DSRef = DPCTLDefaultSelector_Create(); + DRef = DPCTLDevice_CreateFromSelector(DSRef); + QRef = + DPCTLQueue_CreateForDevice(DRef, nullptr, DPCTL_DEFAULT_PROPERTY); + auto CRef = DPCTLQueue_GetContext(QRef); + + KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + DPCTLDevice_Delete(DRef); + DPCTLDeviceSelector_Delete(DSRef); } ~TestQueueSubmit() { spirvFile.close(); + DPCTLQueue_Delete(QRef); + DPCTLKernelBundle_Delete(KBRef); } }; -TEST_F(TestQueueSubmit, CheckSubmitRange_saxpy) -{ - DPCTLSyclDeviceSelectorRef DSRef = nullptr; - DPCTLSyclDeviceRef DRef = nullptr; - - EXPECT_NO_FATAL_FAILURE(DSRef = DPCTLDefaultSelector_Create()); - EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); - DPCTLDeviceMgr_PrintDeviceInfo(DRef); - ASSERT_TRUE(DRef); - auto QRef = - DPCTLQueue_CreateForDevice(DRef, nullptr, DPCTL_DEFAULT_PROPERTY); - ASSERT_TRUE(QRef); - auto CRef = DPCTLQueue_GetContext(QRef); - ASSERT_TRUE(CRef); - auto KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr); - ASSERT_TRUE(KBRef != nullptr); - ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); - auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); +struct TestQueueSubmitFP64 : public ::testing::Test +{ + std::ifstream spirvFile; + size_t spirvFileSize_; + std::vector spirvBuffer_; + DPCTLSyclQueueRef QRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; - // Create the input args - auto a = DPCTLmalloc_shared(SIZE * sizeof(float), QRef); - ASSERT_TRUE(a != nullptr); - auto b = DPCTLmalloc_shared(SIZE * sizeof(float), QRef); - ASSERT_TRUE(b != nullptr); - auto c = DPCTLmalloc_shared(SIZE * sizeof(float), QRef); - ASSERT_TRUE(c != nullptr); + TestQueueSubmitFP64() + { + DPCTLSyclDeviceSelectorRef DSRef = nullptr; + DPCTLSyclDeviceRef DRef = nullptr; + + spirvFile.open("./oneD_range_kernel_fp64.spv", + std::ios::binary | std::ios::ate); + spirvFileSize_ = + std::filesystem::file_size("./oneD_range_kernel_fp64.spv"); + spirvBuffer_.reserve(spirvFileSize_); + spirvFile.seekg(0, std::ios::beg); + spirvFile.read(spirvBuffer_.data(), spirvFileSize_); + DSRef = DPCTLDefaultSelector_Create(); + DRef = DPCTLDevice_CreateFromSelector(DSRef); + QRef = + DPCTLQueue_CreateForDevice(DRef, nullptr, DPCTL_DEFAULT_PROPERTY); + auto CRef = DPCTLQueue_GetContext(QRef); - auto a_ptr = reinterpret_cast(unwrap(a)); - auto b_ptr = reinterpret_cast(unwrap(b)); - // Initialize a,b - for (auto i = 0ul; i < SIZE; ++i) { - a_ptr[i] = i + 1.0; - b_ptr[i] = i + 2.0; + KBRef = DPCTLKernelBundle_CreateFromSpirv( + CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr); + DPCTLDevice_Delete(DRef); + DPCTLDeviceSelector_Delete(DSRef); } - // Create kernel args for axpy - float d = 10.0; - size_t Range[] = {SIZE}; - void *args2[4] = {unwrap(a), unwrap(b), unwrap(c), - (void *)&d}; - DPCTLKernelArgType addKernelArgTypes[] = {DPCTL_VOID_PTR, DPCTL_VOID_PTR, - DPCTL_VOID_PTR, DPCTL_FLOAT}; - auto ERef = DPCTLQueue_SubmitRange( - AxpyKernel, QRef, args2, addKernelArgTypes, 4, Range, 1, nullptr, 0); - ASSERT_TRUE(ERef != nullptr); - DPCTLQueue_Wait(QRef); + ~TestQueueSubmitFP64() + { + spirvFile.close(); + DPCTLQueue_Delete(QRef); + DPCTLKernelBundle_Delete(KBRef); + } +}; - // clean ups - DPCTLEvent_Delete(ERef); - DPCTLKernel_Delete(AxpyKernel); - DPCTLfree_with_queue((DPCTLSyclUSMRef)a, QRef); - DPCTLfree_with_queue((DPCTLSyclUSMRef)b, QRef); - DPCTLfree_with_queue((DPCTLSyclUSMRef)c, QRef); - DPCTLQueue_Delete(QRef); - DPCTLContext_Delete(CRef); - DPCTLKernelBundle_Delete(KBRef); - DPCTLDevice_Delete(DRef); - DPCTLDeviceSelector_Delete(DSRef); -} - -TEST_F(TestQueueSubmit, CheckSubmitNDRange_saxpy) -{ - DPCTLSyclDeviceSelectorRef DSRef = nullptr; - DPCTLSyclDeviceRef DRef = nullptr; - - EXPECT_NO_FATAL_FAILURE(DSRef = DPCTLDefaultSelector_Create()); - EXPECT_NO_FATAL_FAILURE(DRef = DPCTLDevice_CreateFromSelector(DSRef)); - DPCTLDeviceMgr_PrintDeviceInfo(DRef); - ASSERT_TRUE(DRef); - auto QRef = - DPCTLQueue_CreateForDevice(DRef, nullptr, DPCTL_DEFAULT_PROPERTY); - ASSERT_TRUE(QRef); - auto CRef = DPCTLQueue_GetContext(QRef); - ASSERT_TRUE(CRef); - auto KBRef = DPCTLKernelBundle_CreateFromSpirv( - CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr); - ASSERT_TRUE(KBRef != nullptr); - ASSERT_TRUE(DPCTLKernelBundle_HasKernel(KBRef, "axpy")); - auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); +TEST_F(TestQueueSubmit, CheckForInt8) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_INT8_T, + "_ZTS11RangeKernelIaE"); +} - // Create the input args - auto a = DPCTLmalloc_shared(SIZE * sizeof(float), QRef); - ASSERT_TRUE(a != nullptr); - auto b = DPCTLmalloc_shared(SIZE * sizeof(float), QRef); - ASSERT_TRUE(b != nullptr); - auto c = DPCTLmalloc_shared(SIZE * sizeof(float), QRef); - ASSERT_TRUE(c != nullptr); +TEST_F(TestQueueSubmit, CheckForUInt8) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_UINT8_T, + "_ZTS11RangeKernelIhE"); +} - auto a_ptr = reinterpret_cast(unwrap(a)); - auto b_ptr = reinterpret_cast(unwrap(b)); - // Initialize a,b - for (auto i = 0ul; i < SIZE; ++i) { - a_ptr[i] = i + 1.0; - b_ptr[i] = i + 2.0; - } +TEST_F(TestQueueSubmit, CheckForInt16) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_INT16_T, + "_ZTS11RangeKernelIsE"); +} - // Create kernel args for axpy - float d = 10.0; - size_t gRange[] = {1, 1, SIZE}; - size_t lRange[] = {1, 1, 8}; - void *args2[4] = {unwrap(a), unwrap(b), unwrap(c), - (void *)&d}; - DPCTLKernelArgType addKernelArgTypes[] = {DPCTL_VOID_PTR, DPCTL_VOID_PTR, - DPCTL_VOID_PTR, DPCTL_FLOAT}; - DPCTLSyclEventRef events[1]; - events[0] = DPCTLEvent_Create(); +TEST_F(TestQueueSubmit, CheckForUInt16) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_UINT16_T, + "_ZTS11RangeKernelItE"); +} - auto ERef = - DPCTLQueue_SubmitNDRange(AxpyKernel, QRef, args2, addKernelArgTypes, 4, - gRange, lRange, 3, events, 1); - ASSERT_TRUE(ERef != nullptr); - DPCTLQueue_Wait(QRef); +TEST_F(TestQueueSubmit, CheckForInt32) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_INT32_T, + "_ZTS11RangeKernelIiE"); +} - // clean ups - DPCTLEvent_Delete(ERef); - DPCTLKernel_Delete(AxpyKernel); - DPCTLfree_with_queue((DPCTLSyclUSMRef)a, QRef); - DPCTLfree_with_queue((DPCTLSyclUSMRef)b, QRef); - DPCTLfree_with_queue((DPCTLSyclUSMRef)c, QRef); - DPCTLQueue_Delete(QRef); - DPCTLContext_Delete(CRef); - DPCTLKernelBundle_Delete(KBRef); - DPCTLDevice_Delete(DRef); - DPCTLDeviceSelector_Delete(DSRef); +TEST_F(TestQueueSubmit, CheckForUInt32) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_UINT32_T, + "_ZTS11RangeKernelIjE"); +} + +TEST_F(TestQueueSubmit, CheckForInt64) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_INT64_T, + "_ZTS11RangeKernelIlE"); +} + +TEST_F(TestQueueSubmit, CheckForUInt64) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_UINT64_T, + "_ZTS11RangeKernelImE"); +} + +TEST_F(TestQueueSubmit, CheckForFloat) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_FLOAT32_T, + "_ZTS11RangeKernelIfE"); +} + +TEST_F(TestQueueSubmitFP64, CheckForDouble) +{ + submit_kernel(QRef, KBRef, spirvBuffer_, spirvFileSize_, + DPCTLKernelArgType::DPCTL_FLOAT64_T, + "_ZTS11RangeKernelIdE"); +} + +TEST_F(TestQueueSubmit, CheckForUnsupportedArgTy) +{ + + int scalarVal = 3; + size_t Range[] = {SIZE}; + size_t RANGE_NDIMS = 1; + constexpr size_t NARGS = 4; + + auto kernel = DPCTLKernelBundle_GetKernel(KBRef, "_ZTS11RangeKernelIdE"); + void *args[NARGS] = {unwrap(nullptr), unwrap(nullptr), + unwrap(nullptr), (void *)&scalarVal}; + DPCTLKernelArgType addKernelArgTypes[] = {DPCTL_VOID_PTR, DPCTL_VOID_PTR, + DPCTL_VOID_PTR, + DPCTL_UNSUPPORTED_KERNEL_ARG}; + auto ERef = DPCTLQueue_SubmitRange(kernel, QRef, args, addKernelArgTypes, + NARGS, Range, RANGE_NDIMS, nullptr, 0); + + ASSERT_TRUE(ERef == nullptr); } struct TestQueueSubmitBarrier : public ::testing::Test From 733be44060afbff6f60106fb911a7e39d9b5ad6a Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Tue, 5 Mar 2024 00:12:28 -0600 Subject: [PATCH 4/6] Update the dbg_build.sh --- libsyclinterface/dbg_build.sh | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/libsyclinterface/dbg_build.sh b/libsyclinterface/dbg_build.sh index eaa1a02785..f036796a2e 100755 --- a/libsyclinterface/dbg_build.sh +++ b/libsyclinterface/dbg_build.sh @@ -2,30 +2,31 @@ set +xe rm -rf build mkdir build -pushd build +pushd build || exit 1 -INSTALL_PREFIX=`pwd`/../install +INSTALL_PREFIX=$(pwd)/../install rm -rf ${INSTALL_PREFIX} cmake \ -DCMAKE_BUILD_TYPE=Debug \ -DCMAKE_C_COMPILER=icx \ - -DCMAKE_CXX_COMPILER=dpcpp \ + -DCMAKE_CXX_COMPILER=icpx \ + -DCMAKE_CXX_FLAGS=-fsycl \ -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} \ -DDPCTL_ENABLE_L0_PROGRAM_CREATION=ON \ -DDPCTL_BUILD_CAPI_TESTS=ON \ - -DDPCTL_GENERATE_COVERAGE=ON \ .. make V=1 -n -j 4 && make check && make install -# Turn on to generate coverage report html files -make lcov-genhtml +# Turn on to generate coverage report html files reconfigure with +# -DDPCTL_GENERATE_COVERAGE=ON and then +# make lcov-genhtml # For more verbose tests use: # cd tests # ctest -V --progress --output-on-failure -j 4 # cd .. -popd +popd || exit 1 From e36f5b4a160dd9dc3bc5aa778f85b9ca32f65670 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Tue, 5 Mar 2024 10:56:49 -0600 Subject: [PATCH 5/6] Update dpctl Cython API for new DPCTLKernelArgType --- dpctl/_backend.pxd | 25 ++++++++++--------------- dpctl/_sycl_queue.pyx | 32 +++++++++++++------------------- 2 files changed, 23 insertions(+), 34 deletions(-) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 2c0605fd33..e8627e8241 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -57,21 +57,16 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h": _UNKNOWN_DEVICE 'DPCTL_UNKNOWN_DEVICE' ctypedef enum _arg_data_type 'DPCTLKernelArgType': - _CHAR 'DPCTL_CHAR', - _SIGNED_CHAR 'DPCTL_SIGNED_CHAR', - _UNSIGNED_CHAR 'DPCTL_UNSIGNED_CHAR', - _SHORT 'DPCTL_SHORT', - _INT 'DPCTL_INT', - _UNSIGNED_INT 'DPCTL_UNSIGNED_INT', - _UNSIGNED_INT8 'DPCTL_UNSIGNED_INT8', - _LONG 'DPCTL_LONG', - _UNSIGNED_LONG 'DPCTL_UNSIGNED_LONG', - _LONG_LONG 'DPCTL_LONG_LONG', - _UNSIGNED_LONG_LONG 'DPCTL_UNSIGNED_LONG_LONG', - _SIZE_T 'DPCTL_SIZE_T', - _FLOAT 'DPCTL_FLOAT', - _DOUBLE 'DPCTL_DOUBLE', - _LONG_DOUBLE 'DPCTL_DOUBLE', + _INT8_T 'DPCTL_INT8_T', + _UINT8_T 'DPCTL_UINT8_T', + _INT16_T 'DPCTL_INT16_T', + _UINT16_T 'DPCTL_UINT16_T', + _INT32_T 'DPCTL_INT32_T', + _UINT32_T 'DPCTL_UINT32_T', + _INT64_T 'DPCTL_INT64_T', + _UINT64_T 'DPCTL_UINT64_T', + _FLOAT 'DPCTL_FLOAT32_T', + _DOUBLE 'DPCTL_FLOAT64_T', _VOID_PTR 'DPCTL_VOID_PTR' ctypedef enum _queue_property_type 'DPCTLQueuePropertyType': diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 5add749403..542b7b5a47 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -631,34 +631,28 @@ cdef class SyclQueue(_SyclQueue): for idx, arg in enumerate(args): if isinstance(arg, ctypes.c_char): kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._CHAR - elif isinstance(arg, ctypes.c_int): + kargty[idx] = _arg_data_type._INT8_T + elif isinstance(arg, ctypes.c_uint8): kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._INT - elif isinstance(arg, ctypes.c_uint): + kargty[idx] = _arg_data_type._UINT8_T + elif isinstance(arg, ctypes.c_short): kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._UNSIGNED_INT - elif isinstance(arg, ctypes.c_uint8): + kargty[idx] = _arg_data_type._INT16_T + elif isinstance(arg, ctypes.c_ushort): kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._UNSIGNED_INT8 - elif isinstance(arg, ctypes.c_long): + kargty[idx] = _arg_data_type._UINT16_T + elif isinstance(arg, ctypes.c_int): kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._LONG - elif isinstance(arg, ctypes.c_ulong): + kargty[idx] = _arg_data_type._INT32_T + elif isinstance(arg, ctypes.c_uint): kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._UNSIGNED_LONG + kargty[idx] = _arg_data_type._UINT32_T elif isinstance(arg, ctypes.c_longlong): kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._LONG_LONG + kargty[idx] = _arg_data_type._INT64_T elif isinstance(arg, ctypes.c_ulonglong): kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._UNSIGNED_LONG_LONG - elif isinstance(arg, ctypes.c_short): - kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._SHORT - elif isinstance(arg, ctypes.c_size_t): - kargs[idx] = (ctypes.addressof(arg)) - kargty[idx] = _arg_data_type._SIZE_T + kargty[idx] = _arg_data_type._UINT64_T elif isinstance(arg, ctypes.c_float): kargs[idx] = (ctypes.addressof(arg)) kargty[idx] = _arg_data_type._FLOAT From ac95daa1781dd88bc1fd06ddbe967feacc5d7a38 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Wed, 6 Mar 2024 13:07:54 -0600 Subject: [PATCH 6/6] Update comment to clarify how the SPV files were generated. --- .../tests/test_sycl_queue_submit.cpp | 23 +++++++++++-------- 1 file changed, 14 insertions(+), 9 deletions(-) diff --git a/libsyclinterface/tests/test_sycl_queue_submit.cpp b/libsyclinterface/tests/test_sycl_queue_submit.cpp index 1e8f2ae4b3..cc9bc836ce 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit.cpp @@ -91,15 +91,20 @@ void submit_kernel(DPCTLSyclQueueRef QRef, } /* end of anonymous namespace */ /* -// The oneD_range_kernel spv files were generated from the below SYCL program. -// To compile: -// icpx -fsycl oneD_range_kernel.cpp -// IGC_ShaderDumpEnable=1 IGC_DumpToCustomDir=dump ./a.out - -// The generated spv files should be inspected using spirv-dis to identify -// kernel names. When these files were generated using dpcpp 2024.1, a single -// spv file was generated for all integer types and float32_t and a separate -// file generated for float64_t. +// The oneD_range_kernel spv files were generated from the SYCL program included +// in this comment. The program can be compiled using +// `icpx -fsycl oneD_range_kernel.cpp`. After that if the generated executable +// is run with the environment variable `SYCL_DUMP_IMAGES=1`, icpx runtime +// will dump all offload sections of fat binary to the current working +// directory. When tested with DPC++ 2024.0 the kernels are split across two +// separate SPV files. One contains all kernels for integers and FP32 +// data type, and another contains the kernel for FP64. +// +// Note that, `SYCL_DUMP_IMAGES=1` will also generate extra SPV files that +// contain the code for built in functions such as indexing and barriers. To +// figure which SPV file contains the kernels, use `spirv-dis` from the +// spirv-tools package to translate the SPV binary format to a human-readable +// textual format. #include #include