diff --git a/sycl/include/sycl/detail/pi.def b/sycl/include/sycl/detail/pi.def index 195beb0a24861..ba8a245cc6dfd 100644 --- a/sycl/include/sycl/detail/pi.def +++ b/sycl/include/sycl/detail/pi.def @@ -105,6 +105,8 @@ _PI_API(piSamplerCreate) _PI_API(piSamplerGetInfo) _PI_API(piSamplerRetain) _PI_API(piSamplerRelease) +_PI_API(piextSamplerGetNativeHandle) +_PI_API(piextSamplerCreateWithNativeHandle) // Queue commands _PI_API(piEnqueueKernelLaunch) _PI_API(piextEnqueueCooperativeKernelLaunch) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index f3e99c32eb3c9..05ffd1f735790 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -1768,6 +1768,13 @@ __SYCL_EXPORT pi_result piSamplerRetain(pi_sampler sampler); __SYCL_EXPORT pi_result piSamplerRelease(pi_sampler sampler); +__SYCL_EXPORT pi_result piextSamplerCreateWithNativeHandle( + pi_native_handle nativeHandle, pi_context context, + const bool ownNativeHandle, pi_sampler *sampler); + +__SYCL_EXPORT pi_result +piextSamplerGetNativeHandle(pi_sampler sampler, pi_native_handle *nativeHandle); + // // Queue Commands // diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 8bf4eea26620c..a0393313ea5b1 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -654,6 +654,19 @@ pi_result piSamplerRelease(pi_sampler Sampler) { return pi2ur::piSamplerRelease(Sampler); } +pi_result piextSamplerGetNativeHandle(pi_sampler sampler, + pi_native_handle *nativeHandle) { + return pi2ur::piextSamplerGetNativeHandle(sampler, nativeHandle); +} + +pi_result piextSamplerCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, + const bool ownNativeHandle, + pi_sampler *sampler) { + return pi2ur::piextSamplerCreateWithNativeHandle(nativeHandle, context, + ownNativeHandle, sampler); +} + pi_result piEnqueueEventsWait(pi_queue Queue, pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList, pi_event *OutEvent) { diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index 609750a4892b7..ff81ec15d5483 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -657,6 +657,19 @@ pi_result piSamplerRelease(pi_sampler Sampler) { return pi2ur::piSamplerRelease(Sampler); } +pi_result piextSamplerGetNativeHandle(pi_sampler sampler, + pi_native_handle *nativeHandle) { + return pi2ur::piextSamplerGetNativeHandle(sampler, nativeHandle); +} + +pi_result piextSamplerCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, + const bool ownNativeHandle, + pi_sampler *sampler) { + return pi2ur::piextSamplerCreateWithNativeHandle(nativeHandle, context, + ownNativeHandle, sampler); +} + pi_result piEnqueueEventsWait(pi_queue Queue, pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList, pi_event *OutEvent) { diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 8e6224ba5794a..5f6ea4ddc068c 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -676,6 +676,19 @@ pi_result piSamplerRelease(pi_sampler Sampler) { return pi2ur::piSamplerRelease(Sampler); } +pi_result piextSamplerGetNativeHandle(pi_sampler sampler, + pi_native_handle *nativeHandle) { + return pi2ur::piextSamplerGetNativeHandle(sampler, nativeHandle); +} + +pi_result piextSamplerCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, + const bool ownNativeHandle, + pi_sampler *sampler) { + return pi2ur::piextSamplerCreateWithNativeHandle(nativeHandle, context, + ownNativeHandle, sampler); +} + // // Queue Commands // diff --git a/sycl/plugins/native_cpu/pi_native_cpu.cpp b/sycl/plugins/native_cpu/pi_native_cpu.cpp index 35f17a5316bac..97a9460071ac5 100644 --- a/sycl/plugins/native_cpu/pi_native_cpu.cpp +++ b/sycl/plugins/native_cpu/pi_native_cpu.cpp @@ -654,6 +654,19 @@ pi_result piSamplerRelease(pi_sampler Sampler) { return pi2ur::piSamplerRelease(Sampler); } +pi_result piextSamplerGetNativeHandle(pi_sampler sampler, + pi_native_handle *nativeHandle) { + return pi2ur::piextSamplerGetNativeHandle(sampler, nativeHandle); +} + +pi_result piextSamplerCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, + const bool ownNativeHandle, + pi_sampler *sampler) { + return pi2ur::piextSamplerCreateWithNativeHandle(nativeHandle, context, + ownNativeHandle, sampler); +} + // // Queue Commands // diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 9441e29804021..ce3d091162cf0 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -633,6 +633,19 @@ pi_result piSamplerRelease(pi_sampler Sampler) { return pi2ur::piSamplerRelease(Sampler); } +pi_result piextSamplerGetNativeHandle(pi_sampler sampler, + pi_native_handle *nativeHandle) { + return pi2ur::piextSamplerGetNativeHandle(sampler, nativeHandle); +} + +pi_result piextSamplerCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, + const bool ownNativeHandle, + pi_sampler *sampler) { + return pi2ur::piextSamplerCreateWithNativeHandle(nativeHandle, context, + ownNativeHandle, sampler); +} + pi_result piEnqueueEventsWait(pi_queue Queue, pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList, pi_event *OutEvent) { diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index 74a4d461127b6..04406ffea6ecd 100644 --- a/sycl/plugins/unified_runtime/CMakeLists.txt +++ b/sycl/plugins/unified_runtime/CMakeLists.txt @@ -81,14 +81,14 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") - # commit 1f4b703d8a136d49e4e98ca5f017727f3aaa41dc - # Merge: 6883118 a187fbb - # Author: Piotr Balcer - # Date: Mon Mar 25 11:22:45 2024 +0100 - # Merge pull request #1471 from pbalcer/remove-coverage - # [CI] remove coverage workflow - set(UNIFIED_RUNTIME_TAG 1f4b703d8a136d49e4e98ca5f017727f3aaa41dc) + set(UNIFIED_RUNTIME_REPO "https://github.com/omarahmed1111/unified-runtime.git") + # commit 6513abc404979fa109d64500bf899e632d511291 + # Merge: 09be0881 6d586094 + # Author: Kenneth Benzie (Benie) + # Date: Thu Mar 14 22:38:53 2024 +0000 + # Merge pull request #1410 from kbenzie/benie/cmake-external-adapter-source-dirs + # [CMake] Support external adapter source dirs + set(UNIFIED_RUNTIME_TAG 820779f6ede3291227581c2b3f8acaf71a91ee0c) if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO) set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}") diff --git a/sycl/plugins/unified_runtime/pi2ur.hpp b/sycl/plugins/unified_runtime/pi2ur.hpp index f396441f0de5d..94463111b85f7 100644 --- a/sycl/plugins/unified_runtime/pi2ur.hpp +++ b/sycl/plugins/unified_runtime/pi2ur.hpp @@ -4471,6 +4471,38 @@ inline pi_result piSamplerRelease(pi_sampler Sampler) { return PI_SUCCESS; } +pi_result piextSamplerGetNativeHandle(pi_sampler sampler, + pi_native_handle *nativeHandle) { + auto UrSampler = reinterpret_cast(sampler); + + ur_native_handle_t UrNativeHandle{}; + HANDLE_ERRORS(urSamplerGetNativeHandle(UrSampler, &UrNativeHandle)); + + *nativeHandle = reinterpret_cast(UrNativeHandle); + + return PI_SUCCESS; +} + +pi_result piextSamplerCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, + bool ownNativeHandle, + pi_sampler *sampler) { + ur_native_handle_t UrNativeSampler = + reinterpret_cast(nativeHandle); + + ur_context_handle_t UrContext = + reinterpret_cast(context); + + ur_sampler_handle_t *URSampler = + reinterpret_cast(sampler); + ur_sampler_native_properties_t Properties{}; + Properties.isNativeHandleOwned = ownNativeHandle; + HANDLE_ERRORS(urSamplerCreateWithNativeHandle(UrNativeSampler, UrContext, + &Properties, URSampler)); + + return PI_SUCCESS; +} + // Sampler /////////////////////////////////////////////////////////////////////////////// diff --git a/sycl/plugins/unified_runtime/pi_unified_runtime.cpp b/sycl/plugins/unified_runtime/pi_unified_runtime.cpp index 0f42f21d39093..68d10c89f9600 100644 --- a/sycl/plugins/unified_runtime/pi_unified_runtime.cpp +++ b/sycl/plugins/unified_runtime/pi_unified_runtime.cpp @@ -919,6 +919,18 @@ __SYCL_EXPORT pi_result piSamplerRetain(pi_sampler Sampler) { __SYCL_EXPORT pi_result piSamplerRelease(pi_sampler Sampler) { return pi2ur::piSamplerRelease(Sampler); } +/* clang-format off */ +__SYCL_EXPORT pi_result piextSamplerGetNativeHandle( + pi_sampler sampler, pi_native_handle *nativeHandle) { + return pi2ur::piextSamplerGetNativeHandle(sampler, nativeHandle); +} +/* clang-format on */ +__SYCL_EXPORT pi_result piextSamplerCreateWithNativeHandle( + pi_native_handle nativeHandle, pi_context context, + const bool ownNativeHandle, pi_sampler *sampler) { + return pi2ur::piextSamplerCreateWithNativeHandle(nativeHandle, context, + ownNativeHandle, sampler); +} __SYCL_EXPORT pi_result piMemImageGetInfo(pi_mem Image, pi_image_info ParamName, size_t ParamValueSize, @@ -1517,6 +1529,8 @@ __SYCL_EXPORT pi_result piPluginInit(pi_plugin *PluginInit) { _PI_API(piSamplerGetInfo) _PI_API(piSamplerRetain) _PI_API(piSamplerRelease) + _PI_API(piextSamplerCreateWithNativeHandle) + _PI_API(piextSamplerGetNativeHandle) // Peer to Peer _PI_API(piextEnablePeerAccess) diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index ead8f2c83ab71..55327a3bd83b4 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -139,6 +139,9 @@ function(add_sycl_rt_library LIB_NAME LIB_OBJ_NAME) target_link_libraries(${LIB_OBJ_NAME} PRIVATE OpenCL-Headers ) + target_link_libraries(${LIB_NAME} + PRIVATE OpenCL::OpenCL + ) if(SYCL_ENABLE_KERNEL_FUSION) target_link_libraries(${LIB_NAME} PRIVATE sycl-fusion) diff --git a/sycl/source/backend.cpp b/sycl/source/backend.cpp index ed0539f266ee2..591c9931c096d 100644 --- a/sycl/source/backend.cpp +++ b/sycl/source/backend.cpp @@ -94,6 +94,11 @@ __SYCL_EXPORT device make_device(pi_native_handle NativeHandle, __SYCL_EXPORT context make_context(pi_native_handle NativeHandle, const async_handler &Handler, backend Backend) { + if (Backend == backend::opencl) { + sycl::context Context(detail::pi::cast(NativeHandle), Handler); + return Context; + } + const auto &Plugin = getPlugin(Backend); pi::PiContext PiContext = nullptr; @@ -127,7 +132,11 @@ __SYCL_EXPORT queue make_queue(pi_native_handle NativeHandle, make_error_code(errc::invalid), "Queue create using make_queue cannot have compute_index property."); } - + if (Backend == backend::opencl) { + sycl::queue Queue(detail::pi::cast(NativeHandle), Context, + Handler); + return Queue; + } // Create PI queue first. pi::PiQueue PiQueue = nullptr; Plugin->call( @@ -276,6 +285,19 @@ kernel make_kernel(const context &TargetContext, *KernelBundle.begin(); const auto &DeviceImageImpl = getSyclObjImpl(DeviceImage); PiProgram = DeviceImageImpl->get_program_ref(); + } else if (Backend == backend::opencl) { + cl_kernel CLKernel = detail::pi::cast(NativeHandle); + cl_program CLProgram; + size_t Ret = clGetKernelInfo(CLKernel, CL_KERNEL_PROGRAM, sizeof(CLProgram), + &CLProgram, nullptr); + if (Ret != CL_SUCCESS) { + throw runtime_error( + "Failed to retrieve program associated with the kernel", + PI_ERROR_INVALID_KERNEL); + } + Plugin->call( + detail::pi::cast(CLProgram), + ContextImpl->getHandleRef(), false, &PiProgram); } // Create PI kernel first. diff --git a/sycl/source/context.cpp b/sycl/source/context.cpp index 5ae0bba6dadcc..79633afdac9ca 100644 --- a/sycl/source/context.cpp +++ b/sycl/source/context.cpp @@ -82,11 +82,35 @@ context::context(const std::vector &DeviceList, PropList); } } + context::context(cl_context ClContext, async_handler AsyncHandler) { const auto &Plugin = sycl::detail::pi::getPlugin(); - impl = std::make_shared( - detail::pi::cast(ClContext), AsyncHandler, - Plugin); + uint32_t DeviceCount = 0; + size_t Ret = clGetContextInfo(ClContext, CL_CONTEXT_NUM_DEVICES, + sizeof(DeviceCount), &DeviceCount, nullptr); + if (Ret != CL_SUCCESS) { + throw runtime_error( + "Failed to retrieve device count associated with the context", + PI_ERROR_INVALID_CONTEXT); + } + std::vector CLDevices(DeviceCount); + Ret = clGetContextInfo(ClContext, CL_CONTEXT_DEVICES, sizeof(CLDevices), + CLDevices.data(), nullptr); + if (Ret != CL_SUCCESS) { + throw runtime_error( + "Failed to retrieve devices associated with the context", + PI_ERROR_INVALID_CONTEXT); + } + std::vector Devices(DeviceCount); + for (uint32_t i = 0; i < DeviceCount; i++) { + Plugin->call( + detail::pi::cast(CLDevices[i]), nullptr, &Devices[i]); + } + sycl::detail::pi::PiContext Context; + Plugin->call( + detail::pi::cast(ClContext), DeviceCount, + Devices.data(), false, &Context); + impl = std::make_shared(Context, AsyncHandler, Plugin); } template diff --git a/sycl/source/detail/buffer_impl.cpp b/sycl/source/detail/buffer_impl.cpp index 835c732a40bf9..9c89dbfd1a801 100644 --- a/sycl/source/detail/buffer_impl.cpp +++ b/sycl/source/detail/buffer_impl.cpp @@ -54,7 +54,10 @@ void buffer_impl::addInteropObject( const PluginPtr &Plugin = getPlugin(); Plugin->call( pi::cast(MInteropMemObject)); - Handles.push_back(pi::cast(MInteropMemObject)); + pi_native_handle Handle; + Plugin->call(MInteropMemObject, + nullptr, &Handle); + Handles.push_back(Handle); } } } diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp index 9c58335a069c2..9e6bd72de7c83 100644 --- a/sycl/source/detail/context_impl.cpp +++ b/sycl/source/detail/context_impl.cpp @@ -137,9 +137,12 @@ cl_context context_impl::get() const { "This instance of context doesn't support OpenCL interoperability.", PI_ERROR_INVALID_CONTEXT); } - // TODO catch an exception and put it to list of asynchronous exceptions getPlugin()->call(MContext); - return pi::cast(MContext); + // TODO catch an exception and put it to list of asynchronous exceptions + pi_native_handle NativeContext; + getPlugin()->call(MContext, + &NativeContext); + return pi::cast(NativeContext); } bool context_impl::is_host() const { return MHostContext; } diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp index 89b240b816ff9..7de4f290d8ce8 100644 --- a/sycl/source/detail/kernel_impl.cpp +++ b/sycl/source/detail/kernel_impl.cpp @@ -45,16 +45,6 @@ kernel_impl::kernel_impl(sycl::detail::pi::PiKernel Kernel, MCreatedFromSource(IsCreatedFromSource), MKernelBundleImpl(std::move(KernelBundleImpl)), MKernelArgMaskPtr{ArgMask} { - - sycl::detail::pi::PiContext Context = nullptr; - // Using the plugin from the passed ContextImpl - getPlugin()->call( - MKernel, PI_KERNEL_INFO_CONTEXT, sizeof(Context), &Context, nullptr); - if (ContextImpl->getHandleRef() != Context) - throw sycl::invalid_parameter_error( - "Input context must be the same as the context of cl_kernel", - PI_ERROR_INVALID_CONTEXT); - MIsInterop = MProgramImpl->isInterop(); } diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp index e5952fd4d22c7..2a93d0a60d737 100644 --- a/sycl/source/detail/kernel_impl.hpp +++ b/sycl/source/detail/kernel_impl.hpp @@ -107,7 +107,11 @@ class kernel_impl { PI_ERROR_INVALID_KERNEL); } getPlugin()->call(MKernel); - return pi::cast(MKernel); + // TODO catch an exception and put it to list of asynchronous exceptions + pi_native_handle NativeKernel; + getPlugin()->call(MKernel, + &NativeKernel); + return pi::cast(NativeKernel); } /// Check if the associated SYCL context is a SYCL host context. diff --git a/sycl/source/detail/platform_impl.hpp b/sycl/source/detail/platform_impl.hpp index 2948907aa61f4..2f22d1a419e3c 100644 --- a/sycl/source/detail/platform_impl.hpp +++ b/sycl/source/detail/platform_impl.hpp @@ -106,7 +106,7 @@ class platform_impl { "This instance of platform doesn't support OpenCL interoperability.", PI_ERROR_INVALID_PLATFORM); } - return pi::cast(MPlatform); + return pi::cast(getNative()); } /// Returns raw underlying plug-in platform handle. diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp index d65f3163b961f..2873175eadf67 100644 --- a/sycl/source/detail/program_impl.cpp +++ b/sycl/source/detail/program_impl.cpp @@ -222,7 +222,11 @@ cl_program program_impl::get() const { PI_ERROR_INVALID_PROGRAM); } getPlugin()->call(MProgram); - return pi::cast(MProgram); + // TODO catch an exception and put it to list of asynchronous exceptions + pi_native_handle NativeProgram; + getPlugin()->call(MProgram, + &NativeProgram); + return pi::cast(NativeProgram); } void program_impl::compile_with_kernel_name(std::string KernelName, @@ -389,24 +393,23 @@ std::pair program_impl::get_pi_kernel_arg_mask_pair(const std::string &KernelName) const { std::pair Result; - const PluginPtr &Plugin = getPlugin(); - sycl::detail::pi::PiResult Err = - Plugin->call_nocheck( - MProgram, KernelName.c_str(), &Result.first); - if (Err == PI_ERROR_INVALID_KERNEL_NAME) { - throw invalid_object_error( - "This instance of program does not contain the kernel requested", - Err); - } - Plugin->checkPiResult(Err); + const PluginPtr &Plugin = getPlugin(); + sycl::detail::pi::PiResult Err = + Plugin->call_nocheck( + MProgram, KernelName.c_str(), &Result.first); + if (Err == PI_ERROR_INVALID_KERNEL_NAME) { + throw invalid_object_error( + "This instance of program does not contain the kernel requested", Err); + } + Plugin->checkPiResult(Err); - // Some PI Plugins (like OpenCL) require this call to enable USM - // For others, PI will turn this into a NOP. - if (getContextImplPtr()->getPlatformImpl()->supports_usm()) - Plugin->call( - Result.first, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE); + // Some PI Plugins (like OpenCL) require this call to enable USM + // For others, PI will turn this into a NOP. + if (getContextImplPtr()->getPlatformImpl()->supports_usm()) + Plugin->call( + Result.first, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE); - return Result; + return Result; } std::vector diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index 890891644bbac..9e3046ba02244 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -359,7 +359,11 @@ class queue_impl { PI_ERROR_INVALID_QUEUE); } getPlugin()->call(MQueues[0]); - return pi::cast(MQueues[0]); + // TODO catch an exception and put it to list of asynchronous exceptions + pi_native_handle NativeQueue; + getPlugin()->call( + MQueues[0], &NativeQueue, nullptr); + return pi::cast(NativeQueue); } /// \return an associated SYCL context. diff --git a/sycl/source/detail/sampler_impl.cpp b/sycl/source/detail/sampler_impl.cpp index c2af7884a164c..4c59de3b4a639 100644 --- a/sycl/source/detail/sampler_impl.cpp +++ b/sycl/source/detail/sampler_impl.cpp @@ -23,11 +23,14 @@ sampler_impl::sampler_impl(coordinate_normalization_mode normalizationMode, sampler_impl::sampler_impl(cl_sampler clSampler, const context &syclContext) { - sycl::detail::pi::PiSampler Sampler = - pi::cast(clSampler); + sycl::detail::pi::PiSampler Sampler; + + auto Plugin = sycl::detail::pi::getPlugin(); + Plugin->call( + pi::cast(clSampler), + getSyclObjImpl(syclContext)->getHandleRef(), false, &Sampler); + MContextToSampler[syclContext] = Sampler; - const PluginPtr &Plugin = getSyclObjImpl(syclContext)->getPlugin(); - Plugin->call(Sampler); Plugin->call( Sampler, PI_SAMPLER_INFO_NORMALIZED_COORDS, sizeof(pi_bool), &MCoordNormMode, nullptr); diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 7781fb7e1cd1e..2d50226a757ed 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -2299,23 +2299,11 @@ void SetArgBasedOnType( getMemAllocationFunc ? (sycl::detail::pi::PiMem)getMemAllocationFunc(Req) : nullptr; - if (Context.get_backend() == backend::opencl) { - // clSetKernelArg (corresponding to piKernelSetArg) returns an error - // when MemArg is null, which is the case when zero-sized buffers are - // handled. Below assignment provides later call to clSetKernelArg with - // acceptable arguments. - if (!MemArg) - MemArg = sycl::detail::pi::PiMem(); - - Plugin->call( - Kernel, NextTrueIndex, sizeof(sycl::detail::pi::PiMem), &MemArg); - } else { - pi_mem_obj_property MemObjData{}; - MemObjData.mem_access = AccessModeToPi(Req->MAccessMode); - MemObjData.type = PI_KERNEL_ARG_MEM_OBJ_ACCESS; - Plugin->call(Kernel, NextTrueIndex, - &MemObjData, &MemArg); - } + pi_mem_obj_property MemObjData{}; + MemObjData.mem_access = AccessModeToPi(Req->MAccessMode); + MemObjData.type = PI_KERNEL_ARG_MEM_OBJ_ACCESS; + Plugin->call(Kernel, NextTrueIndex, + &MemObjData, &MemArg); break; } case kernel_param_kind_t::kind_std_layout: { diff --git a/sycl/source/event.cpp b/sycl/source/event.cpp index 897c48ca7e890..b4941ff258e4e 100644 --- a/sycl/source/event.cpp +++ b/sycl/source/event.cpp @@ -24,13 +24,14 @@ inline namespace _V1 { event::event() : impl(std::make_shared(std::nullopt)) {} -event::event(cl_event ClEvent, const context &SyclContext) - : impl(std::make_shared( - detail::pi::cast(ClEvent), SyclContext)) { - // This is a special interop constructor for OpenCL, so the event must be - // retained. - impl->getPlugin()->call( - detail::pi::cast(ClEvent)); +event::event(cl_event ClEvent, const context &SyclContext) { + sycl::detail::pi::PiEvent Event; + auto Plugin = sycl::detail::pi::getPlugin(); + Plugin->call( + detail::pi::cast(ClEvent), + detail::getSyclObjImpl(SyclContext)->getHandleRef(), false, &Event); + + impl = std::make_shared(Event, SyclContext); } bool event::operator==(const event &rhs) const { return rhs.impl == impl; } diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp index d981833f94958..bab6a43537701 100644 --- a/sycl/source/kernel.cpp +++ b/sycl/source/kernel.cpp @@ -16,15 +16,33 @@ namespace sycl { inline namespace _V1 { -kernel::kernel(cl_kernel ClKernel, const context &SyclContext) - : impl(std::make_shared( - detail::pi::cast(ClKernel), - detail::getSyclObjImpl(SyclContext), nullptr, nullptr)) { - // This is a special interop constructor for OpenCL, so the kernel must be - // retained. - if (get_backend() == backend::opencl) { - impl->getPlugin()->call( - detail::pi::cast(ClKernel)); +kernel::kernel(cl_kernel ClKernel, const context &SyclContext) { + try { + sycl::detail::pi::PiKernel Kernel; + auto Context = detail::getSyclObjImpl(SyclContext); + auto Plugin = sycl::detail::pi::getPlugin(); + cl_program CLProgram; + size_t Ret = clGetKernelInfo(ClKernel, CL_KERNEL_PROGRAM, sizeof(CLProgram), + &CLProgram, nullptr); + if (Ret != CL_SUCCESS) { + throw runtime_error( + "Failed to retrieve program associated with the kernel", + PI_ERROR_INVALID_KERNEL); + } + sycl::detail::pi::PiProgram Program; + Plugin->call( + detail::pi::cast(CLProgram), Context->getHandleRef(), + false, &Program); + + Plugin->call( + detail::pi::cast(ClKernel), Context->getHandleRef(), + Program, false, &Kernel); + impl = std::make_shared(Kernel, Context, nullptr, + nullptr); + } catch (sycl::runtime_error &) { + throw sycl::invalid_parameter_error( + "Input context must be the same as the context of cl_kernel", + PI_ERROR_INVALID_CONTEXT); } } diff --git a/sycl/source/platform.cpp b/sycl/source/platform.cpp index c991f27eda51b..d4f5622f47a08 100644 --- a/sycl/source/platform.cpp +++ b/sycl/source/platform.cpp @@ -21,9 +21,12 @@ inline namespace _V1 { platform::platform() : platform(default_selector_v) {} platform::platform(cl_platform_id PlatformId) { - impl = detail::platform_impl::getOrMakePlatformImpl( - detail::pi::cast(PlatformId), - sycl::detail::pi::getPlugin()); + sycl::detail::pi::PiPlatform Platform; + auto Plugin = sycl::detail::pi::getPlugin(); + Plugin->call( + detail::pi::cast(PlatformId), &Platform); + + impl = detail::platform_impl::getOrMakePlatformImpl(Platform, Plugin); } // protected constructor for internal use diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index acaecf2696629..6f8cd773224c9 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -64,9 +64,29 @@ queue::queue(const context &SyclContext, const device &SyclDevice, queue::queue(cl_command_queue clQueue, const context &SyclContext, const async_handler &AsyncHandler) { const property_list PropList{}; - impl = std::make_shared( - reinterpret_cast(clQueue), - detail::getSyclObjImpl(SyclContext), AsyncHandler, PropList); + sycl::detail::pi::PiQueue Queue; + auto Context = detail::getSyclObjImpl(SyclContext); + auto Plugin = sycl::detail::pi::getPlugin(); + + cl_device_id CLDevice; + size_t Ret = clGetCommandQueueInfo(clQueue, CL_QUEUE_DEVICE, sizeof(CLDevice), + &CLDevice, nullptr); + if (Ret) { + throw runtime_error("Failed to retrieve device associated with the queue", + PI_ERROR_INVALID_QUEUE); + } + sycl::detail::pi::PiDevice Device; + Plugin->call( + detail::pi::cast(CLDevice), nullptr, &Device); + + sycl::detail::pi::PiQueueProperties Properties[] = {PI_QUEUE_FLAGS, 0, 0, 0, + 0}; + Plugin->call( + detail::pi::cast(clQueue), 0, Context->getHandleRef(), + Device, false, Properties, &Queue); + + impl = std::make_shared(Queue, Context, AsyncHandler, + PropList); } cl_command_queue queue::get() const { return impl->get(); } diff --git a/sycl/test-e2e/DeprecatedFeatures/set_arg_interop.cpp b/sycl/test-e2e/DeprecatedFeatures/set_arg_interop.cpp index 986f9314823ab..70e938b1a2815 100644 --- a/sycl/test-e2e/DeprecatedFeatures/set_arg_interop.cpp +++ b/sycl/test-e2e/DeprecatedFeatures/set_arg_interop.cpp @@ -6,6 +6,7 @@ #include #include +#include using namespace sycl; @@ -15,7 +16,7 @@ int main() { cl_context ClContext = Context.get(); - const size_t CountSources = 3; + const size_t CountSources = 4; const char *Sources[CountSources] = { "kernel void foo1(global float* Array, global int* Value) { *Array = " "42; *Value = 1; }\n", @@ -23,6 +24,7 @@ int main() { "Array[id] = id; }\n", "kernel void foo3(global float* Array, local float* LocalArray) { " "(void)LocalArray; (void)Array; }\n", + "kernel void foo4(global int* Value) {}\n", }; cl_int Err; @@ -42,12 +44,16 @@ int main() { cl_kernel ThirdCLKernel = clCreateKernel(ClProgram, "foo3", &Err); assert(Err == CL_SUCCESS); + cl_kernel FourthCLKernel = clCreateKernel(ClProgram, "foo4", &Err); + assert(Err == CL_SUCCESS); + const size_t Count = 100; float Array[Count]; kernel FirstKernel(FirstCLKernel, Context); kernel SecondKernel(SecondCLKernel, Context); kernel ThirdKernel(ThirdCLKernel, Context); + kernel FourthKernel(FourthCLKernel, Context); int Value; { buffer FirstBuffer(Array, range<1>(1)); @@ -110,10 +116,37 @@ int main() { } Queue.wait_and_throw(); + // Enqueuing an interop kernel while avoid calls to piKernelSetArg from + // different threads on the same kernel. + { + constexpr std::size_t NArgs = 16; + constexpr std::size_t ThreadCount = 4; + constexpr std::size_t LaunchCount = 8; + auto TestLambda = [&](int ThreadId) { + Queue + .submit([&](sycl::handler &CGH) { + for (std::size_t I = 0; I < NArgs; ++I) + CGH.set_arg(I, &ThreadId); + }) + .wait(); + }; + + std::vector threadPool; + threadPool.reserve(ThreadCount); + for (size_t tid = 0; tid < ThreadCount; ++tid) { + threadPool.push_back(std::thread(TestLambda, tid)); + } + + for (auto ¤tThread : threadPool) { + currentThread.join(); + } + } + clReleaseContext(ClContext); clReleaseKernel(FirstCLKernel); clReleaseKernel(SecondCLKernel); clReleaseKernel(ThirdCLKernel); + clReleaseKernel(FourthCLKernel); clReleaseProgram(ClProgram); return 0; } diff --git a/sycl/test/abi/pi_cuda_symbol_check.dump b/sycl/test/abi/pi_cuda_symbol_check.dump index 0a7db9e19498c..061f4b0ea8cfb 100644 --- a/sycl/test/abi/pi_cuda_symbol_check.dump +++ b/sycl/test/abi/pi_cuda_symbol_check.dump @@ -81,6 +81,8 @@ piSamplerCreate piSamplerGetInfo piSamplerRelease piSamplerRetain +piextSamplerGetNativeHandle +piextSamplerCreateWithNativeHandle piTearDown piextBindlessImageSamplerCreate piextCommandBufferAdviseUSM diff --git a/sycl/test/abi/pi_hip_symbol_check.dump b/sycl/test/abi/pi_hip_symbol_check.dump index ab85eb32b8ce5..fb78b7842994e 100644 --- a/sycl/test/abi/pi_hip_symbol_check.dump +++ b/sycl/test/abi/pi_hip_symbol_check.dump @@ -81,6 +81,8 @@ piSamplerCreate piSamplerGetInfo piSamplerRelease piSamplerRetain +piextSamplerGetNativeHandle +piextSamplerCreateWithNativeHandle piTearDown piextBindlessImageSamplerCreate piextCommandBufferAdviseUSM diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index 2ebc6b56078a4..cf764ada4876b 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -80,6 +80,8 @@ piSamplerCreate piSamplerGetInfo piSamplerRelease piSamplerRetain +piextSamplerGetNativeHandle +piextSamplerCreateWithNativeHandle piTearDown piextBindlessImageSamplerCreate piextCommandBufferAdviseUSM diff --git a/sycl/test/abi/pi_nativecpu_symbol_check.dump b/sycl/test/abi/pi_nativecpu_symbol_check.dump index 6198c8aeb5832..9134ceb8eda1c 100644 --- a/sycl/test/abi/pi_nativecpu_symbol_check.dump +++ b/sycl/test/abi/pi_nativecpu_symbol_check.dump @@ -81,6 +81,8 @@ piSamplerCreate piSamplerGetInfo piSamplerRelease piSamplerRetain +piextSamplerGetNativeHandle +piextSamplerCreateWithNativeHandle piTearDown piextBindlessImageSamplerCreate piextCommandBufferAdviseUSM diff --git a/sycl/test/abi/pi_opencl_symbol_check.dump b/sycl/test/abi/pi_opencl_symbol_check.dump index 86860b50e57b6..20551fa837da5 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -80,6 +80,8 @@ piSamplerCreate piSamplerGetInfo piSamplerRelease piSamplerRetain +piextSamplerGetNativeHandle +piextSamplerCreateWithNativeHandle piTearDown piextBindlessImageSamplerCreate piextCommandBufferAdviseUSM diff --git a/sycl/unittests/assert/assert.cpp b/sycl/unittests/assert/assert.cpp index b94d43925d374..63d0a1c8f7525 100644 --- a/sycl/unittests/assert/assert.cpp +++ b/sycl/unittests/assert/assert.cpp @@ -217,19 +217,6 @@ static pi_result redefinedEventsWaitPositive(pi_uint32 num_events, return PI_SUCCESS; } -static pi_result redefinedEventsWaitNegative(pi_uint32 num_events, - const pi_event *event_list) { - // For negative tests we do not expect the copier kernel to be used, so - // instead we accept whatever amount we get. - // This output here is to reduce amount of time requried to debug/reproduce - // a failing test upon feature break - printf("Waiting for %i events ", num_events); - for (size_t I = 0; I < num_events; ++I) - printf("%i, ", reinterpret_cast(event_list[I])[0]); - printf("\n"); - return PI_SUCCESS; -} - static pi_result redefinedEnqueueMemBufferMapAfter( pi_queue command_queue, pi_mem buffer, pi_bool blocking_map, pi_map_flags map_flags, size_t offset, size_t size, @@ -256,153 +243,6 @@ static void setupMock(sycl::unittest::PiMock &Mock) { Mock.redefineBefore(redefinedEventsWaitPositive); } -namespace TestInteropKernel { -const sycl::context *Context = nullptr; -const sycl::device *Device = nullptr; -int KernelLaunchCounter = ::KernelLaunchCounterBase; - -static pi_result redefinedKernelGetInfo(pi_kernel Kernel, - pi_kernel_info ParamName, - size_t ParamValueSize, void *ParamValue, - size_t *ParamValueSizeRet) { - if (PI_KERNEL_INFO_CONTEXT == ParamName) { - pi_context PiContext = - sycl::detail::getSyclObjImpl(*Context)->getHandleRef(); - - if (ParamValue) - memcpy(ParamValue, &PiContext, sizeof(PiContext)); - if (ParamValueSizeRet) - *ParamValueSizeRet = sizeof(PiContext); - - return PI_SUCCESS; - } - - if (PI_KERNEL_INFO_PROGRAM == ParamName) { - pi_program PIProgram = nullptr; - pi_result Res = mock_piProgramCreate(/*pi_context=*/0x0, /**il*/ nullptr, - /*length=*/0, &PIProgram); - EXPECT_TRUE(PI_SUCCESS == Res); - - if (ParamValue) - memcpy(ParamValue, &PIProgram, sizeof(PIProgram)); - if (ParamValueSizeRet) - *ParamValueSizeRet = sizeof(PIProgram); - - return PI_SUCCESS; - } - - if (PI_KERNEL_INFO_FUNCTION_NAME == ParamName) { - static const char FName[] = "TestFnName"; - if (ParamValue) { - size_t L = strlen(FName) + 1; - if (L < ParamValueSize) - L = ParamValueSize; - - memcpy(ParamValue, FName, L); - } - if (ParamValueSizeRet) - *ParamValueSizeRet = strlen(FName) + 1; - - return PI_SUCCESS; - } - - return PI_ERROR_UNKNOWN; -} - -static pi_result redefinedEnqueueKernelLaunch(pi_queue, pi_kernel, pi_uint32, - const size_t *, const size_t *, - const size_t *LocalSize, - pi_uint32 N, const pi_event *Deps, - pi_event *RetEvent) { - int Val = KernelLaunchCounter++; - // This output here is to reduce amount of time requried to debug/reproduce a - // failing test upon feature break - printf("Enqueued %i\n", Val); - - return PI_SUCCESS; -} - -static pi_result redefinedProgramGetInfo(pi_program P, - pi_program_info ParamName, - size_t ParamValueSize, - void *ParamValue, - size_t *ParamValueSizeRet) { - if (PI_PROGRAM_INFO_NUM_DEVICES == ParamName) { - static const int V = 1; - - if (ParamValue) - memcpy(ParamValue, &V, sizeof(V)); - if (ParamValueSizeRet) - *ParamValueSizeRet = sizeof(V); - - return PI_SUCCESS; - } - - if (PI_PROGRAM_INFO_DEVICES == ParamName) { - EXPECT_EQ(ParamValueSize, 1 * sizeof(pi_device)); - - pi_device Dev = sycl::detail::getSyclObjImpl(*Device)->getHandleRef(); - - if (ParamValue) - memcpy(ParamValue, &Dev, sizeof(Dev)); - if (ParamValueSizeRet) - *ParamValueSizeRet = sizeof(Dev); - - return PI_SUCCESS; - } - - return PI_ERROR_UNKNOWN; -} - -static pi_result redefinedProgramGetBuildInfo(pi_program P, pi_device D, - pi_program_build_info ParamName, - size_t ParamValueSize, - void *ParamValue, - size_t *ParamValueSizeRet) { - if (PI_PROGRAM_BUILD_INFO_BINARY_TYPE == ParamName) { - static const pi_program_binary_type T = PI_PROGRAM_BINARY_TYPE_EXECUTABLE; - if (ParamValue) - memcpy(ParamValue, &T, sizeof(T)); - if (ParamValueSizeRet) - *ParamValueSizeRet = sizeof(T); - return PI_SUCCESS; - } - - if (PI_PROGRAM_BUILD_INFO_OPTIONS == ParamName) { - if (ParamValueSizeRet) - *ParamValueSizeRet = 0; - return PI_SUCCESS; - } - - return PI_ERROR_UNKNOWN; -} - -} // namespace TestInteropKernel - -static void setupMockForInterop(sycl::unittest::PiMock &Mock, - const sycl::context &Ctx, - const sycl::device &Dev) { - using namespace sycl::detail; - - TestInteropKernel::KernelLaunchCounter = ::KernelLaunchCounterBase; - TestInteropKernel::Device = &Dev; - TestInteropKernel::Context = &Ctx; - - Mock.redefineAfter( - redefinedKernelGetGroupInfoAfter); - Mock.redefineBefore( - TestInteropKernel::redefinedEnqueueKernelLaunch); - Mock.redefineAfter( - redefinedEnqueueMemBufferMapAfter); - Mock.redefineBefore(redefinedEventsWaitNegative); - Mock.redefineBefore( - TestInteropKernel::redefinedKernelGetInfo); - Mock.redefineBefore( - TestInteropKernel::redefinedProgramGetInfo); - Mock.redefineBefore( - TestInteropKernel::redefinedProgramGetBuildInfo); -} - #ifndef _WIN32 void ChildProcess(int StdErrFD) { static constexpr int StandardStdErrFD = 2; @@ -528,54 +368,3 @@ TEST(Assert, TestAssertServiceKernelHidden) { EXPECT_TRUE(NoFoundServiceKernelID); } - -TEST(Assert, TestInteropKernelNegative) { - sycl::unittest::PiMock Mock; - sycl::platform Plt = Mock.getPlatform(); - - const sycl::device Dev = Plt.get_devices()[0]; - sycl::context Ctx{Dev}; - - setupMockForInterop(Mock, Ctx, Dev); - - sycl::queue Queue{Ctx, Dev}; - - pi_kernel PIKernel = nullptr; - - pi_result Res = mock_piKernelCreate( - /*pi_program=*/0x0, /*kernel_name=*/"dummy_kernel", &PIKernel); - EXPECT_TRUE(PI_SUCCESS == Res); - - // TODO use make_kernel. This requires a fix in backend.cpp to get plugin - // from context instead of free getPlugin to alllow for mocking of its methods - sycl::kernel KInterop((cl_kernel)PIKernel, Ctx); - - Queue.submit([&](sycl::handler &H) { H.single_task(KInterop); }); - - EXPECT_EQ(TestInteropKernel::KernelLaunchCounter, - KernelLaunchCounterBase + 1); -} - -TEST(Assert, TestInteropKernelFromProgramNegative) { - sycl::unittest::PiMock Mock; - sycl::platform Plt = Mock.getPlatform(); - - const sycl::device Dev = Plt.get_devices()[0]; - sycl::context Ctx{Dev}; - - setupMockForInterop(Mock, Ctx, Dev); - - sycl::queue Queue{Ctx, Dev}; - - sycl::kernel_bundle Bundle = - sycl::get_kernel_bundle(Ctx); - sycl::kernel KOrig = Bundle.get_kernel(sycl::get_kernel_id()); - - cl_kernel CLKernel = sycl::get_native(KOrig); - sycl::kernel KInterop{CLKernel, Ctx}; - - Queue.submit([&](sycl::handler &H) { H.single_task(KInterop); }); - - EXPECT_EQ(TestInteropKernel::KernelLaunchCounter, - KernelLaunchCounterBase + 1); -} diff --git a/sycl/unittests/handler/CMakeLists.txt b/sycl/unittests/handler/CMakeLists.txt index eb7fc559ab73c..3f90404ab35b9 100644 --- a/sycl/unittests/handler/CMakeLists.txt +++ b/sycl/unittests/handler/CMakeLists.txt @@ -1,4 +1,3 @@ add_sycl_unittest(HandlerTests OBJECT - SetArgForLocalAccessor.cpp require.cpp ) diff --git a/sycl/unittests/handler/SetArgForLocalAccessor.cpp b/sycl/unittests/handler/SetArgForLocalAccessor.cpp deleted file mode 100644 index ab7d8f387b761..0000000000000 --- a/sycl/unittests/handler/SetArgForLocalAccessor.cpp +++ /dev/null @@ -1,53 +0,0 @@ -//==------- SetArgForLocalAccessor.cpp --- Handler unit tests --------------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include -#include -#include - -#include - -// This test checks that we pass the correct buffer size value when setting -// local_accessor as an argument through handler::set_arg to a kernel created -// using OpenCL interoperability methods. - -namespace { - -size_t LocalBufferArgSize = 0; - -pi_result redefined_piKernelSetArg(pi_kernel kernel, pi_uint32 arg_index, - size_t arg_size, const void *arg_value) { - LocalBufferArgSize = arg_size; - - return PI_SUCCESS; -} - -TEST(HandlerSetArg, LocalAccessor) { - sycl::unittest::PiMock Mock; - redefineMockForKernelInterop(Mock); - Mock.redefine( - redefined_piKernelSetArg); - - constexpr size_t Size = 128; - sycl::queue Q; - - DummyHandleT handle; - auto KernelCL = reinterpret_cast::template input_type>(&handle); - auto Kernel = - sycl::make_kernel(KernelCL, Q.get_context()); - - Q.submit([&](sycl::handler &CGH) { - sycl::local_accessor Acc(Size, CGH); - CGH.set_arg(0, Acc); - CGH.single_task(Kernel); - }).wait(); - - ASSERT_EQ(LocalBufferArgSize, Size * sizeof(float)); -} -} // namespace diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index a0f267bd97d50..108262608fb13 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -974,6 +974,21 @@ inline pi_result mock_piSamplerRelease(pi_sampler sampler) { return PI_SUCCESS; } +inline pi_result +mock_piextSamplerGetNativeHandle(pi_sampler sampler, + pi_native_handle *nativeHandle) { + *nativeHandle = reinterpret_cast(sampler); + return PI_SUCCESS; +} + +inline pi_result mock_piextSamplerCreateWithNativeHandle( + pi_native_handle nativeHandle, pi_context context, + const bool ownNativeHandle, pi_sampler *sampler) { + *sampler = reinterpret_cast(nativeHandle); + retainDummyHandle(*sampler); + return PI_SUCCESS; +} + // // Queue Commands // diff --git a/sycl/unittests/thread_safety/CMakeLists.txt b/sycl/unittests/thread_safety/CMakeLists.txt index 8b725af8b4dd4..78dc6f2190178 100644 --- a/sycl/unittests/thread_safety/CMakeLists.txt +++ b/sycl/unittests/thread_safety/CMakeLists.txt @@ -1,4 +1,3 @@ add_sycl_unittest(ThreadSafetyTests OBJECT HostAccessorDeadLock.cpp - InteropKernelEnqueue.cpp ) diff --git a/sycl/unittests/thread_safety/InteropKernelEnqueue.cpp b/sycl/unittests/thread_safety/InteropKernelEnqueue.cpp deleted file mode 100644 index 79f19504abea1..0000000000000 --- a/sycl/unittests/thread_safety/InteropKernelEnqueue.cpp +++ /dev/null @@ -1,65 +0,0 @@ -//==-------- InteropKernelEnqueue.cpp --- Thread safety unit tests ---------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include -#include -#include -#include -#include - -#include "ThreadUtils.h" - -namespace { -using namespace sycl; - -constexpr std::size_t NArgs = 16; -constexpr std::size_t ThreadCount = 4; -constexpr std::size_t LaunchCount = 8; - -pi_uint32 LastArgSet = -1; -std::size_t LastThread = -1; -pi_result redefined_piKernelSetArg(pi_kernel kernel, pi_uint32 arg_index, - size_t arg_size, const void *arg_value) { - EXPECT_EQ((LastArgSet + 1) % NArgs, arg_index); - LastArgSet = arg_index; - std::size_t ArgValue = *static_cast(arg_value); - if (arg_index == 0) - LastThread = ArgValue; - else - EXPECT_EQ(LastThread, ArgValue); - return PI_SUCCESS; -} - -TEST(KernelEnqueue, InteropKernel) { - unittest::PiMock Mock; - redefineMockForKernelInterop(Mock); - Mock.redefine( - redefined_piKernelSetArg); - - platform Plt = Mock.getPlatform(); - queue Q; - - DummyHandleT Handle; - auto KernelCL = reinterpret_cast::template input_type>(&Handle); - auto Kernel = - sycl::make_kernel(KernelCL, Q.get_context()); - - auto TestLambda = [&](std::size_t ThreadId) { - Q.submit([&](sycl::handler &CGH) { - for (std::size_t I = 0; I < NArgs; ++I) - CGH.set_arg(I, ThreadId); - CGH.single_task(Kernel); - }).wait(); - }; - - for (std::size_t I = 0; I < LaunchCount; ++I) { - ThreadPool Pool(ThreadCount, TestLambda); - } -} -} // namespace