Skip to content

Commit

Permalink
[SYCL][UR] Bump UR version and add missing values to pi2ur (#10049)
Browse files Browse the repository at this point in the history 
Includes a large number of missing `pi_result` mappings, so previously
reported errors should no longer map to just `PI_ERROR_UNKNOWN`. NFCI
for the adapters.
  • Loading branch information
@callumfare
callumfare authored Jun 26, 2023
1 parent 7e98cd4 commit 9fcb2ac
Show file tree
Hide file tree
Showing 9 changed files with 187 additions and 84 deletions.
2 changes: 1 addition & 1 deletion sycl/plugins/unified_runtime/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@ if (NOT DEFINED UNIFIED_RUNTIME_LIBRARY OR NOT DEFINED UNIFIED_RUNTIME_INCLUDE_D
include(FetchContent)

set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
set(UNIFIED_RUNTIME_TAG 4136fbb19c37a8aa9d368559a738e2e7cc35033e)
set(UNIFIED_RUNTIME_TAG 7e16bb37cbb12450637e595749c3617151cbe851)

message(STATUS "Will fetch Unified Runtime from ${UNIFIED_RUNTIME_REPO}")
FetchContent_Declare(unified-runtime
Expand Down
161 changes: 130 additions & 31 deletions sycl/plugins/unified_runtime/pi2ur.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,7 @@

#include "ur_api.h"
#include <cstdarg>
#include <sycl/detail/cuda_definitions.hpp>
#include <sycl/detail/pi.h>
#include <ur/ur.hpp>

Expand All @@ -18,38 +19,128 @@ static pi_result ur2piResult(ur_result_t urResult) {
return PI_SUCCESS;

switch (urResult) {
case UR_RESULT_ERROR_UNKNOWN:
return PI_ERROR_UNKNOWN;
case UR_RESULT_ERROR_DEVICE_LOST:
return PI_ERROR_DEVICE_NOT_FOUND;
case UR_RESULT_ERROR_INVALID_OPERATION:
return PI_ERROR_INVALID_OPERATION;
case UR_RESULT_ERROR_INVALID_PLATFORM:
return PI_ERROR_INVALID_PLATFORM;
case UR_RESULT_ERROR_INVALID_ARGUMENT:
return PI_ERROR_INVALID_ARG_VALUE;
case UR_RESULT_ERROR_INVALID_QUEUE_PROPERTIES:
return PI_ERROR_INVALID_QUEUE_PROPERTIES;
case UR_RESULT_ERROR_INVALID_QUEUE:
return PI_ERROR_INVALID_QUEUE;
case UR_RESULT_ERROR_INVALID_VALUE:
return PI_ERROR_INVALID_VALUE;
case UR_RESULT_ERROR_INVALID_EVENT:
return PI_ERROR_INVALID_EVENT;
case UR_RESULT_ERROR_INVALID_CONTEXT:
return PI_ERROR_INVALID_CONTEXT;
case UR_RESULT_ERROR_INVALID_PLATFORM:
return PI_ERROR_INVALID_PLATFORM;
case UR_RESULT_ERROR_INVALID_BINARY:
return PI_ERROR_INVALID_BINARY;
case UR_RESULT_ERROR_INVALID_KERNEL_NAME:
return PI_ERROR_INVALID_KERNEL_NAME;
case UR_RESULT_ERROR_INVALID_FUNCTION_NAME:
return PI_ERROR_BUILD_PROGRAM_FAILURE;
case UR_RESULT_ERROR_INVALID_PROGRAM:
return PI_ERROR_INVALID_PROGRAM;
case UR_RESULT_ERROR_INVALID_SAMPLER:
return PI_ERROR_INVALID_SAMPLER;
case UR_RESULT_ERROR_INVALID_MEM_OBJECT:
return PI_ERROR_INVALID_MEM_OBJECT;
case UR_RESULT_ERROR_INVALID_EVENT:
return PI_ERROR_INVALID_EVENT;
case UR_RESULT_ERROR_INVALID_EVENT_WAIT_LIST:
return PI_ERROR_INVALID_EVENT_WAIT_LIST;
case UR_RESULT_ERROR_MISALIGNED_SUB_BUFFER_OFFSET:
return PI_ERROR_MISALIGNED_SUB_BUFFER_OFFSET;
case UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE:
return PI_ERROR_INVALID_WORK_GROUP_SIZE;
case UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY:
return PI_ERROR_OUT_OF_RESOURCES;
case UR_RESULT_ERROR_COMPILER_NOT_AVAILABLE:
return PI_ERROR_COMPILER_NOT_AVAILABLE;
case UR_RESULT_ERROR_PROFILING_INFO_NOT_AVAILABLE:
return PI_ERROR_PROFILING_INFO_NOT_AVAILABLE;
case UR_RESULT_ERROR_DEVICE_NOT_FOUND:
return PI_ERROR_DEVICE_NOT_FOUND;
case UR_RESULT_ERROR_INVALID_DEVICE:
return PI_ERROR_INVALID_DEVICE;
case UR_RESULT_ERROR_DEVICE_REQUIRES_RESET:
case UR_RESULT_ERROR_DEVICE_LOST:
return PI_ERROR_DEVICE_NOT_AVAILABLE;
case UR_RESULT_ERROR_DEVICE_PARTITION_FAILED:
return PI_ERROR_DEVICE_PARTITION_FAILED;
case UR_RESULT_ERROR_INVALID_DEVICE_PARTITION_COUNT:
return PI_ERROR_INVALID_DEVICE_PARTITION_COUNT;
case UR_RESULT_ERROR_INVALID_WORK_ITEM_SIZE:
return PI_ERROR_INVALID_WORK_ITEM_SIZE;
case UR_RESULT_ERROR_INVALID_WORK_DIMENSION:
return PI_ERROR_INVALID_WORK_DIMENSION;
case UR_RESULT_ERROR_INVALID_KERNEL_ARGS:
return PI_ERROR_INVALID_KERNEL_ARGS;
case UR_RESULT_ERROR_INVALID_KERNEL:
return PI_ERROR_INVALID_KERNEL;
case UR_RESULT_ERROR_INVALID_KERNEL_NAME:
return PI_ERROR_INVALID_KERNEL_NAME;
case UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX:
return PI_ERROR_INVALID_ARG_INDEX;
case UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE:
return PI_ERROR_INVALID_ARG_SIZE;
case UR_RESULT_ERROR_INVALID_KERNEL_ATTRIBUTE_VALUE:
return PI_ERROR_INVALID_VALUE;
case UR_RESULT_ERROR_INVALID_IMAGE_SIZE:
return PI_ERROR_INVALID_IMAGE_SIZE;
case UR_RESULT_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR:
return PI_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR;
case UR_RESULT_ERROR_IMAGE_FORMAT_NOT_SUPPORTED:
return PI_ERROR_IMAGE_FORMAT_NOT_SUPPORTED;
case UR_RESULT_ERROR_MEM_OBJECT_ALLOCATION_FAILURE:
return PI_ERROR_MEM_OBJECT_ALLOCATION_FAILURE;
case UR_RESULT_ERROR_INVALID_PROGRAM_EXECUTABLE:
return PI_ERROR_INVALID_PROGRAM_EXECUTABLE;
case UR_RESULT_ERROR_UNINITIALIZED:
return PI_ERROR_UNINITIALIZED;
case UR_RESULT_ERROR_OUT_OF_HOST_MEMORY:
return PI_ERROR_OUT_OF_HOST_MEMORY;
case UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY:
case UR_RESULT_ERROR_OUT_OF_RESOURCES:
return PI_ERROR_OUT_OF_RESOURCES;
case UR_RESULT_ERROR_PROGRAM_BUILD_FAILURE:
return PI_ERROR_BUILD_PROGRAM_FAILURE;
case UR_RESULT_ERROR_UNINITIALIZED:
return PI_ERROR_UNINITIALIZED;
case UR_RESULT_ERROR_PROGRAM_LINK_FAILURE:
return PI_ERROR_LINK_PROGRAM_FAILURE;
case UR_RESULT_ERROR_UNSUPPORTED_VERSION:
case UR_RESULT_ERROR_UNSUPPORTED_FEATURE:
case UR_RESULT_ERROR_INVALID_ARGUMENT:
case UR_RESULT_ERROR_INVALID_NULL_HANDLE:
case UR_RESULT_ERROR_HANDLE_OBJECT_IN_USE:
case UR_RESULT_ERROR_INVALID_NULL_POINTER:
return PI_ERROR_INVALID_VALUE;
case UR_RESULT_ERROR_INVALID_SIZE:
case UR_RESULT_ERROR_UNSUPPORTED_SIZE:
return PI_ERROR_INVALID_BUFFER_SIZE;
case UR_RESULT_ERROR_UNSUPPORTED_ALIGNMENT:
return PI_ERROR_INVALID_VALUE;
case UR_RESULT_ERROR_INVALID_SYNCHRONIZATION_OBJECT:
case UR_RESULT_ERROR_INVALID_ENUMERATION:
case UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION:
return PI_ERROR_INVALID_VALUE;
case UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT:
return PI_ERROR_IMAGE_FORMAT_NOT_SUPPORTED;
case UR_RESULT_ERROR_INVALID_NATIVE_BINARY:
return PI_ERROR_INVALID_BINARY;
case UR_RESULT_ERROR_INVALID_GLOBAL_NAME:
return PI_ERROR_INVALID_VALUE;
case UR_RESULT_ERROR_INVALID_FUNCTION_NAME:
return PI_ERROR_FUNCTION_ADDRESS_IS_NOT_AVAILABLE;
case UR_RESULT_ERROR_INVALID_GROUP_SIZE_DIMENSION:
return PI_ERROR_INVALID_WORK_DIMENSION;
case UR_RESULT_ERROR_INVALID_GLOBAL_WIDTH_DIMENSION:
return PI_ERROR_INVALID_VALUE;

case UR_RESULT_ERROR_PROGRAM_UNLINKED:
return PI_ERROR_INVALID_PROGRAM_EXECUTABLE;
case UR_RESULT_ERROR_OVERLAPPING_REGIONS:
return PI_ERROR_MEM_COPY_OVERLAP;
case UR_RESULT_ERROR_INVALID_HOST_PTR:
return PI_ERROR_INVALID_HOST_PTR;
case UR_RESULT_ERROR_INVALID_USM_SIZE:
return PI_ERROR_INVALID_BUFFER_SIZE;
case UR_RESULT_ERROR_OBJECT_ALLOCATION_FAILURE:
return PI_ERROR_OUT_OF_RESOURCES;
case UR_RESULT_ERROR_ADAPTER_SPECIFIC:
return PI_ERROR_PLUGIN_SPECIFIC_ERROR;
case UR_RESULT_ERROR_UNKNOWN:
default:
return PI_ERROR_UNKNOWN;
};
Expand Down Expand Up @@ -253,6 +344,10 @@ inline pi_result ur2piDeviceInfoValue(ur_device_info_t ParamName,
return PI_QUEUE_FLAG_ON_DEVICE;
case UR_QUEUE_FLAG_ON_DEVICE_DEFAULT:
return PI_QUEUE_FLAG_ON_DEVICE_DEFAULT;
case UR_QUEUE_FLAG_SYNC_WITH_DEFAULT_STREAM:
return static_cast<uint64_t>(__SYCL_PI_CUDA_SYNC_WITH_DEFAULT);
case UR_QUEUE_FLAG_USE_DEFAULT_STREAM:
return static_cast<uint64_t>(__SYCL_PI_CUDA_USE_DEFAULT_STREAM);
default:
die("UR_DEVICE_INFO_QUEUE_PROPERTIES: unhandled value");
}
Expand Down Expand Up @@ -1014,7 +1109,7 @@ inline pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
break;
}
case PI_EXT_CODEPLAY_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP: {
InfoType = UR_EXT_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP;
InfoType = UR_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP;
break;
}
default:
Expand Down Expand Up @@ -1380,6 +1475,10 @@ inline pi_result piextQueueCreate(pi_context Context, pi_device Device,
UrProperties.flags |= UR_QUEUE_FLAG_PRIORITY_LOW;
if (Properties[1] & PI_EXT_ONEAPI_QUEUE_FLAG_PRIORITY_HIGH)
UrProperties.flags |= UR_QUEUE_FLAG_PRIORITY_HIGH;
if (Properties[1] & __SYCL_PI_CUDA_SYNC_WITH_DEFAULT)
UrProperties.flags |= UR_QUEUE_FLAG_SYNC_WITH_DEFAULT_STREAM;
if (Properties[1] & __SYCL_PI_CUDA_USE_DEFAULT_STREAM)
UrProperties.flags |= UR_QUEUE_FLAG_USE_DEFAULT_STREAM;

ur_queue_index_properties_t IndexProperties{};
IndexProperties.stype = UR_STRUCTURE_TYPE_QUEUE_INDEX_PROPERTIES;
Expand Down Expand Up @@ -1943,15 +2042,16 @@ inline pi_result piKernelSetArg(pi_kernel Kernel, pi_uint32 ArgIndex,

ur_kernel_handle_t UrKernel = reinterpret_cast<ur_kernel_handle_t>(Kernel);

HANDLE_ERRORS(urKernelSetArgValue(UrKernel, ArgIndex, ArgSize, ArgValue));
HANDLE_ERRORS(
urKernelSetArgValue(UrKernel, ArgIndex, ArgSize, nullptr, ArgValue));
return PI_SUCCESS;
}

inline pi_result piKernelSetArgPointer(pi_kernel Kernel, pi_uint32 ArgIndex,
size_t ArgSize, const void *ArgValue) {
std::ignore = ArgSize;
ur_kernel_handle_t UrKernel = reinterpret_cast<ur_kernel_handle_t>(Kernel);
HANDLE_ERRORS(urKernelSetArgPointer(UrKernel, ArgIndex, ArgValue));
HANDLE_ERRORS(urKernelSetArgPointer(UrKernel, ArgIndex, nullptr, ArgValue));

return PI_SUCCESS;
}
Expand Down Expand Up @@ -2013,17 +2113,15 @@ inline pi_result piKernelSetExecInfo(pi_kernel Kernel,
break;
}
case PI_EXT_KERNEL_EXEC_INFO_CACHE_CONFIG: {
PropName = UR_EXT_KERNEL_EXEC_INFO_CACHE_CONFIG;
PropName = UR_KERNEL_EXEC_INFO_CACHE_CONFIG;
auto Param = (*(static_cast<const pi_kernel_cache_config *>(ParamValue)));
if (Param == PI_EXT_KERNEL_EXEC_INFO_CACHE_LARGE_SLM) {
PropValue =
static_cast<uint64_t>(UR_EXT_KERNEL_EXEC_INFO_CACHE_LARGE_SLM);
PropValue = static_cast<uint64_t>(UR_KERNEL_CACHE_CONFIG_LARGE_SLM);
} else if (Param == PI_EXT_KERNEL_EXEC_INFO_CACHE_LARGE_DATA) {
PropValue =
static_cast<uint64_t>(UR_EXT_KERNEL_EXEC_INFO_CACHE_LARGE_DATA);
PropValue = static_cast<uint64_t>(UR_KERNEL_CACHE_CONFIG_LARGE_DATA);
break;
} else if (Param == PI_EXT_KERNEL_EXEC_INFO_CACHE_DEFAULT) {
PropValue = static_cast<uint64_t>(UR_EXT_KERNEL_EXEC_INFO_CACHE_DEFAULT);
PropValue = static_cast<uint64_t>(UR_KERNEL_CACHE_CONFIG_DEFAULT);
} else {
die("piKernelSetExecInfo: unsupported ParamValue\n");
}
Expand All @@ -2032,8 +2130,8 @@ inline pi_result piKernelSetExecInfo(pi_kernel Kernel,
default:
die("piKernelSetExecInfo: unsupported ParamName\n");
}
HANDLE_ERRORS(
urKernelSetExecInfo(UrKernel, PropName, ParamValueSize, &PropValue));
HANDLE_ERRORS(urKernelSetExecInfo(UrKernel, PropName, ParamValueSize, nullptr,
&PropValue));

return PI_SUCCESS;
}
Expand Down Expand Up @@ -2211,7 +2309,8 @@ inline pi_result piextKernelSetArgPointer(pi_kernel Kernel, pi_uint32 ArgIndex,
const void *ArgValue) {
ur_kernel_handle_t UrKernel = reinterpret_cast<ur_kernel_handle_t>(Kernel);

HANDLE_ERRORS(urKernelSetArgValue(UrKernel, ArgIndex, ArgSize, ArgValue));
HANDLE_ERRORS(
urKernelSetArgValue(UrKernel, ArgIndex, ArgSize, nullptr, ArgValue));

return PI_SUCCESS;
}
Expand Down Expand Up @@ -3919,7 +4018,7 @@ inline pi_result piextKernelSetArgSampler(pi_kernel Kernel, pi_uint32 ArgIndex,
ur_sampler_handle_t UrSampler =
reinterpret_cast<ur_sampler_handle_t>(*ArgValue);

HANDLE_ERRORS(urKernelSetArgSampler(UrKernel, ArgIndex, UrSampler));
HANDLE_ERRORS(urKernelSetArgSampler(UrKernel, ArgIndex, nullptr, UrSampler));

return PI_SUCCESS;
}
Expand Down
2 changes: 1 addition & 1 deletion sycl/plugins/unified_runtime/ur/adapters/cuda/device.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -968,7 +968,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,

return ReturnValue(ILVersion.data(), ILVersion.size());
}
case UR_EXT_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP: {
case UR_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP: {
// Maximum number of 32-bit registers available to a thread block.
// Note: This number is shared by all thread blocks simultaneously resident
// on a multiprocessor.
Expand Down
25 changes: 17 additions & 8 deletions sycl/plugins/unified_runtime/ur/adapters/cuda/kernel.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -186,9 +186,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelGetNativeHandle(
return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
}

UR_APIEXPORT ur_result_t UR_APICALL
urKernelSetArgValue(ur_kernel_handle_t hKernel, uint32_t argIndex,
size_t argSize, const void *pArgValue) {
UR_APIEXPORT ur_result_t UR_APICALL urKernelSetArgValue(
ur_kernel_handle_t hKernel, uint32_t argIndex, size_t argSize,
const ur_kernel_arg_value_properties_t *pProperties,
const void *pArgValue) {
std::ignore = pProperties;
UR_ASSERT(hKernel, UR_RESULT_ERROR_INVALID_NULL_HANDLE);
UR_ASSERT(argSize, UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE);

Expand Down Expand Up @@ -289,8 +291,11 @@ urKernelGetSubGroupInfo(ur_kernel_handle_t hKernel, ur_device_handle_t hDevice,
return UR_RESULT_ERROR_INVALID_ENUMERATION;
}

UR_APIEXPORT ur_result_t UR_APICALL urKernelSetArgPointer(
ur_kernel_handle_t hKernel, uint32_t argIndex, const void *pArgValue) {
UR_APIEXPORT ur_result_t UR_APICALL
urKernelSetArgPointer(ur_kernel_handle_t hKernel, uint32_t argIndex,
const ur_kernel_arg_pointer_properties_t *pProperties,
const void *pArgValue) {
std::ignore = pProperties;
hKernel->setKernelArg(argIndex, sizeof(pArgValue), pArgValue);
return UR_RESULT_SUCCESS;
}
Expand Down Expand Up @@ -339,10 +344,12 @@ urKernelSetArgMemObj(ur_kernel_handle_t hKernel, uint32_t argIndex,
}

// A NOP for the CUDA backend
UR_APIEXPORT ur_result_t UR_APICALL
urKernelSetExecInfo(ur_kernel_handle_t hKernel, ur_kernel_exec_info_t propName,
size_t propSize, const void *pPropValue) {
UR_APIEXPORT ur_result_t UR_APICALL urKernelSetExecInfo(
ur_kernel_handle_t hKernel, ur_kernel_exec_info_t propName, size_t propSize,
const ur_kernel_exec_info_properties_t *pProperties,
const void *pPropValue) {
std::ignore = propSize;
std::ignore = pProperties;
UR_ASSERT(hKernel, UR_RESULT_ERROR_INVALID_NULL_HANDLE);
UR_ASSERT(pPropValue, UR_RESULT_ERROR_INVALID_NULL_POINTER);
switch (propName) {
Expand Down Expand Up @@ -370,8 +377,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelCreateWithNativeHandle(

UR_APIEXPORT ur_result_t UR_APICALL
urKernelSetArgSampler(ur_kernel_handle_t hKernel, uint32_t argIndex,
const ur_kernel_arg_sampler_properties_t *pProperties,
ur_sampler_handle_t hArgValue) {
UR_ASSERT(hKernel, UR_RESULT_ERROR_INVALID_NULL_HANDLE);
std::ignore = pProperties;

ur_result_t Result = UR_RESULT_SUCCESS;
try {
Expand Down
8 changes: 4 additions & 4 deletions sycl/plugins/unified_runtime/ur/adapters/cuda/queue.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -134,9 +134,9 @@ urQueueCreate(ur_context_handle_t hContext, ur_device_handle_t hDevice,
bool IsOutOfOrder = false;
if (pProps && pProps->stype == UR_STRUCTURE_TYPE_QUEUE_PROPERTIES) {
URFlags = pProps->flags;
if (URFlags == __SYCL_UR_CUDA_USE_DEFAULT_STREAM) {
if (URFlags == UR_QUEUE_FLAG_USE_DEFAULT_STREAM) {
Flags = CU_STREAM_DEFAULT;
} else if (URFlags == __SYCL_UR_CUDA_SYNC_WITH_DEFAULT) {
} else if (URFlags == UR_QUEUE_FLAG_SYNC_WITH_DEFAULT_STREAM) {
Flags = 0;
}

Expand Down Expand Up @@ -261,9 +261,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueCreateWithNativeHandle(

ur_queue_flags_t Flags = 0;
if (CuFlags == CU_STREAM_DEFAULT)
Flags = __SYCL_UR_CUDA_USE_DEFAULT_STREAM;
Flags = UR_QUEUE_FLAG_USE_DEFAULT_STREAM;
else if (CuFlags == CU_STREAM_NON_BLOCKING)
Flags = __SYCL_UR_CUDA_SYNC_WITH_DEFAULT;
Flags = UR_QUEUE_FLAG_SYNC_WITH_DEFAULT_STREAM;
else
sycl::detail::ur::die("Unknown cuda stream");

Expand Down
12 changes: 6 additions & 6 deletions sycl/plugins/unified_runtime/ur/adapters/level_zero/ur_level_zero_event.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -70,7 +70,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWait(
ur_event_handle_t InternalEvent;
bool IsInternal = OutEvent == nullptr;
ur_event_handle_t *Event = OutEvent ? OutEvent : &InternalEvent;
UR_CALL(createEventAndAssociateQueue(Queue, Event, UR_EXT_COMMAND_TYPE_USER,
UR_CALL(createEventAndAssociateQueue(Queue, Event, UR_COMMAND_EVENTS_WAIT,
CommandList, IsInternal));

ZeEvent = (*Event)->ZeEvent;
Expand Down Expand Up @@ -98,10 +98,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWait(
std::scoped_lock<ur_shared_mutex> lock(Queue->Mutex);

if (OutEvent) {
UR_CALL(createEventAndAssociateQueue(Queue, OutEvent,
UR_EXT_COMMAND_TYPE_USER,
Queue->CommandListMap.end(),
/* IsInternal */ false));
UR_CALL(createEventAndAssociateQueue(
Queue, OutEvent, UR_COMMAND_EVENTS_WAIT, Queue->CommandListMap.end(),
/* IsInternal */ false));
}

Queue->synchronize();
Expand Down Expand Up @@ -161,7 +160,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWaitWithBarrier(
}

UR_CALL(createEventAndAssociateQueue(
Queue, &Event, UR_EXT_COMMAND_TYPE_USER, CmdList, IsInternal));
Queue, &Event, UR_COMMAND_EVENTS_WAIT_WITH_BARRIER, CmdList,
IsInternal));

Event->WaitList = EventWaitList;

Expand Down
Loading

0 comments on commit 9fcb2ac

Please sign in to comment.