From 38b85189e081b26044decd4a3c34056527659963 Mon Sep 17 00:00:00 2001 From: "Zhao, Maosu" Date: Tue, 11 Nov 2025 07:47:51 +0100 Subject: [PATCH 1/3] [DevASAN] Move memory alloc info into DeviceInfo Since CPU/GPU device both support USM indirect access, we need to poison shadow of whole allocated memory in the device instead of only one context. --- .../loader/layers/sanitizer/asan/asan_ddi.cpp | 26 ++++++++++++++++- .../sanitizer/asan/asan_interceptor.cpp | 28 +++++++++---------- .../sanitizer/asan/asan_interceptor.hpp | 22 +++++++-------- .../layers/sanitizer/asan/asan_validator.cpp | 6 ++-- .../layers/sanitizer/asan/asan_validator.hpp | 3 +- .../loader/layers/validation/ur_valddi.cpp | 3 -- 6 files changed, 55 insertions(+), 33 deletions(-) diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_ddi.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_ddi.cpp index 3df7897f406c4..3caa93871aa20 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_ddi.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_ddi.cpp @@ -56,7 +56,6 @@ ur_result_t setupContext(ur_context_handle_t Context, uint32_t numDevices, (void *)DI->Handle, (void *)Context); DI->Shadow = ShadowMemory; CI->DeviceList.emplace_back(hDevice); - CI->AllocInfosMap[hDevice]; } } return UR_RESULT_SUCCESS; @@ -1623,6 +1622,30 @@ __urdlllocal ur_result_t UR_APICALL urKernelSetArgPointer( return result; } +__urdlllocal ur_result_t UR_APICALL urKernelSetExecInfo( + /// [in] handle of the kernel object + ur_kernel_handle_t hKernel, + /// [in] name of the execution attribute + ur_kernel_exec_info_t propName, + /// [in] size in byte the attribute value + size_t propSize, + /// [in][optional] pointer to execution info properties. + const ur_kernel_exec_info_properties_t *pProperties, + /// [in][typename(propName, propSize)] pointer to memory location holding + /// the property value. + const void *pPropValue) { + UR_LOG_L(getContext()->logger, DEBUG, "==== urKernelSetExecInfo"); + + UR_CALL(getContext()->urDdiTable.Kernel.pfnSetExecInfo( + hKernel, propName, propSize, pProperties, pPropValue)); + auto &KI = getAsanInterceptor()->getOrCreateKernelInfo(hKernel); + if (propName == UR_KERNEL_EXEC_INFO_USM_INDIRECT_ACCESS) { + KI.IsIndirectAccess = *ur_cast(pPropValue); + } + + return UR_RESULT_SUCCESS; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urDeviceGetInfo __urdlllocal ur_result_t UR_APICALL urDeviceGetInfo( @@ -1928,6 +1951,7 @@ __urdlllocal ur_result_t UR_APICALL urGetKernelProcAddrTable( pDdiTable->pfnSetArgMemObj = ur_sanitizer_layer::asan::urKernelSetArgMemObj; pDdiTable->pfnSetArgLocal = ur_sanitizer_layer::asan::urKernelSetArgLocal; pDdiTable->pfnSetArgPointer = ur_sanitizer_layer::asan::urKernelSetArgPointer; + pDdiTable->pfnSetExecInfo = ur_sanitizer_layer::asan::urKernelSetExecInfo; return result; } diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp index 45eed900f9150..9e52425869906 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp @@ -143,10 +143,11 @@ ur_result_t AsanInterceptor::allocateMemory(ur_context_handle_t Context, AI->print(); // For updating shadow memory - if (Device) { // Device/Shared USM - ContextInfo->insertAllocInfo({Device}, AI); + if (DeviceInfo) { // Device/Shared USM + DeviceInfo->insertAllocInfo(AI); } else { // Host USM - ContextInfo->insertAllocInfo(ContextInfo->DeviceList, AI); + for (const auto &Device : ContextInfo->DeviceList) + getDeviceInfo(Device)->insertAllocInfo(AI); } // For memory release @@ -212,9 +213,10 @@ ur_result_t AsanInterceptor::releaseMemory(ur_context_handle_t Context, AllocInfo->ReleaseStack = GetCurrentBacktrace(); if (AllocInfo->Type == AllocType::HOST_USM) { - ContextInfo->insertAllocInfo(ContextInfo->DeviceList, AllocInfo); + for (const auto &Device : ContextInfo->DeviceList) + getDeviceInfo(Device)->insertAllocInfo(AllocInfo); } else { - ContextInfo->insertAllocInfo({AllocInfo->Device}, AllocInfo); + getDeviceInfo(AllocInfo->Device)->insertAllocInfo(AllocInfo); } // If quarantine is disabled, USM is freed immediately @@ -279,7 +281,7 @@ ur_result_t AsanInterceptor::preLaunchKernel(ur_kernel_handle_t Kernel, (void)ArgPointer; } } - UR_CALL(updateShadowMemory(ContextInfo, DeviceInfo, InternalQueue)); + UR_CALL(updateShadowMemory(DeviceInfo, InternalQueue)); UR_CALL(prepareLaunch(ContextInfo, DeviceInfo, InternalQueue, Kernel, LaunchInfo)); @@ -423,16 +425,14 @@ AsanInterceptor::enqueueAllocInfo(std::shared_ptr &DeviceInfo, } ur_result_t -AsanInterceptor::updateShadowMemory(std::shared_ptr &ContextInfo, - std::shared_ptr &DeviceInfo, +AsanInterceptor::updateShadowMemory(std::shared_ptr &DeviceInfo, ur_queue_handle_t Queue) { - auto &AllocInfos = ContextInfo->AllocInfosMap[DeviceInfo->Handle]; - std::scoped_lock Guard(AllocInfos.Mutex); + std::scoped_lock Guard(DeviceInfo->AllocInfos.Mutex); - for (auto &AI : AllocInfos.List) { + for (auto &AI : DeviceInfo->AllocInfos.List) { UR_CALL(enqueueAllocInfo(DeviceInfo, Queue, AI)); } - AllocInfos.List.clear(); + DeviceInfo->AllocInfos.List.clear(); return UR_RESULT_SUCCESS; } @@ -585,7 +585,7 @@ AsanInterceptor::registerDeviceGlobals(ur_program_handle_t Program) { GetCurrentBacktrace(), {}}); - ContextInfo->insertAllocInfo({Device}, AI); + getDeviceInfo(Device)->insertAllocInfo(AI); ProgramInfo->AllocInfoForGlobals.emplace(AI); std::scoped_lock Guard(m_AllocationMapMutex); @@ -754,7 +754,7 @@ ur_result_t AsanInterceptor::prepareLaunch( continue; } if (auto ValidateResult = ValidateUSMPointer( - ContextInfo->Handle, DeviceInfo->Handle, (uptr)Ptr)) { + Kernel, ContextInfo->Handle, DeviceInfo->Handle, (uptr)Ptr)) { ReportInvalidKernelArgument(Kernel, ArgIndex, (uptr)Ptr, ValidateResult, PtrPair.second); if (ValidateResult.Type != ValidateUSMResult::MAYBE_HOST_POINTER) { diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp index 7e4bc8edee7e3..ba7b8d301e442 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp @@ -54,9 +54,16 @@ struct DeviceInfo { std::queue> Quarantine; size_t QuarantineSize = 0; + AllocInfoList AllocInfos; + // Device handles are special and alive in the whole process lifetime, // so we needn't retain&release here. explicit DeviceInfo(ur_device_handle_t Device) : Handle(Device) {} + + void insertAllocInfo(std::shared_ptr &AI) { + std::scoped_lock Guard(AllocInfos.Mutex); + AllocInfos.List.emplace_back(AI); + } }; struct QueueInfo { @@ -88,6 +95,8 @@ struct KernelInfo { bool IsInstrumented = false; // check shadow bounds bool IsCheckShadowBounds = false; + // might have indirect access + bool IsIndirectAccess = false; // lock this mutex if following fields are accessed ur_shared_mutex Mutex; @@ -147,7 +156,6 @@ struct ContextInfo { std::atomic RefCount = 1; std::vector DeviceList; - std::unordered_map AllocInfosMap; ur_shared_mutex InternalQueueMapMutex; std::unordered_map> @@ -169,15 +177,6 @@ struct ContextInfo { ~ContextInfo(); - void insertAllocInfo(const std::vector &Devices, - std::shared_ptr &AI) { - for (auto Device : Devices) { - auto &AllocInfos = AllocInfosMap[Device]; - std::scoped_lock Guard(AllocInfos.Mutex); - AllocInfos.List.emplace_back(AI); - } - } - ur_usm_pool_handle_t getUSMPool(); ur_queue_handle_t getInternalQueue(ur_device_handle_t); @@ -375,8 +374,7 @@ class AsanInterceptor { ur_shared_mutex KernelLaunchMutex; private: - ur_result_t updateShadowMemory(std::shared_ptr &ContextInfo, - std::shared_ptr &DeviceInfo, + ur_result_t updateShadowMemory(std::shared_ptr &DeviceInfo, ur_queue_handle_t Queue); ur_result_t enqueueAllocInfo(std::shared_ptr &DeviceInfo, diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_validator.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_validator.cpp index 2ff16ae0bfd8e..32b61d58d1c8b 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_validator.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_validator.cpp @@ -36,7 +36,8 @@ bool IsSameDevice(ur_device_handle_t Device1, ur_device_handle_t Device2) { } // namespace -ValidateUSMResult ValidateUSMPointer(ur_context_handle_t Context, +ValidateUSMResult ValidateUSMPointer(ur_kernel_handle_t Kernel, + ur_context_handle_t Context, ur_device_handle_t Device, uptr Ptr) { assert(Ptr != 0 && "Don't validate nullptr here"); @@ -53,7 +54,8 @@ ValidateUSMResult ValidateUSMPointer(ur_context_handle_t Context, auto AllocInfo = AllocInfoItOp.value()->second; - if (AllocInfo->Context != Context) { + auto &KI = getAsanInterceptor()->getOrCreateKernelInfo(Kernel); + if (!KI.IsIndirectAccess && AllocInfo->Context != Context) { return ValidateUSMResult::fail(ValidateUSMResult::BAD_CONTEXT, AllocInfo); } diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_validator.hpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_validator.hpp index 0979a4b601f88..5e26bd3f64ee7 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_validator.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_validator.hpp @@ -46,7 +46,8 @@ struct ValidateUSMResult { } }; -ValidateUSMResult ValidateUSMPointer(ur_context_handle_t Context, +ValidateUSMResult ValidateUSMPointer(ur_kernel_handle_t Kernel, + ur_context_handle_t Context, ur_device_handle_t Device, uptr Ptr); } // namespace asan diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index 464acb714b2a5..244a9f28179d9 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -3864,9 +3864,6 @@ __urdlllocal ur_result_t UR_APICALL urKernelGetSuggestedLocalWorkSize( } if (getContext()->enableParameterValidation) { - if (NULL == pGlobalWorkOffset) - return UR_RESULT_ERROR_INVALID_NULL_POINTER; - if (NULL == pGlobalWorkSize) return UR_RESULT_ERROR_INVALID_NULL_POINTER; From 74716bd3cb5629a594b6b627d6f1265b87e87be0 Mon Sep 17 00:00:00 2001 From: "Zhao, Maosu" Date: Wed, 12 Nov 2025 04:13:51 +0100 Subject: [PATCH 2/3] reverted validation layer changes --- .../layers/sanitizer/asan/asan_interceptor.cpp | 2 +- .../layers/sanitizer/asan/asan_interceptor.hpp | 13 ++++++++++--- .../layers/sanitizer/msan/msan_interceptor.cpp | 2 +- .../layers/sanitizer/msan/msan_interceptor.hpp | 13 ++++++++++--- .../layers/sanitizer/tsan/tsan_interceptor.cpp | 2 +- .../layers/sanitizer/tsan/tsan_interceptor.hpp | 13 ++++++++++--- .../source/loader/layers/validation/ur_valddi.cpp | 3 +++ 7 files changed, 36 insertions(+), 12 deletions(-) diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp index 9e52425869906..75de16d79a8a3 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.cpp @@ -801,7 +801,7 @@ ur_result_t AsanInterceptor::prepareLaunch( if (LaunchInfo.LocalWorkSize.empty()) { LaunchInfo.LocalWorkSize.resize(LaunchInfo.WorkDim); auto URes = getContext()->urDdiTable.Kernel.pfnGetSuggestedLocalWorkSize( - Kernel, Queue, LaunchInfo.WorkDim, LaunchInfo.GlobalWorkOffset, + Kernel, Queue, LaunchInfo.WorkDim, LaunchInfo.GlobalWorkOffset.data(), LaunchInfo.GlobalWorkSize, LaunchInfo.LocalWorkSize.data()); if (URes != UR_RESULT_SUCCESS) { if (URes != UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { diff --git a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp index ba7b8d301e442..bee76e839e1cc 100644 --- a/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/asan/asan_interceptor.hpp @@ -248,7 +248,7 @@ struct LaunchInfo { ur_context_handle_t Context = nullptr; ur_device_handle_t Device = nullptr; const size_t *GlobalWorkSize = nullptr; - const size_t *GlobalWorkOffset = nullptr; + std::vector GlobalWorkOffset; std::vector LocalWorkSize; uint32_t WorkDim = 0; @@ -258,12 +258,19 @@ struct LaunchInfo { const size_t *GlobalWorkSize, const size_t *LocalWorkSize, const size_t *GlobalWorkOffset, uint32_t WorkDim) : Context(Context), Device(Device), GlobalWorkSize(GlobalWorkSize), - GlobalWorkOffset(GlobalWorkOffset), WorkDim(WorkDim), - Data(Context, Device) { + WorkDim(WorkDim), Data(Context, Device) { if (LocalWorkSize) { this->LocalWorkSize = std::vector(LocalWorkSize, LocalWorkSize + WorkDim); } + // UR doesn't allow GlobalWorkOffset is null, we need to construct a zero + // value array if user doesn't specify its value. + if (GlobalWorkOffset) { + this->GlobalWorkOffset = + std::vector(GlobalWorkOffset, GlobalWorkOffset + WorkDim); + } else { + this->GlobalWorkOffset = std::vector(WorkDim, 0); + } [[maybe_unused]] auto Result = getContext()->urDdiTable.Context.pfnRetain(Context); assert(Result == UR_RESULT_SUCCESS); diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp index 347dce609085b..806fd9f9638b5 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp @@ -509,7 +509,7 @@ ur_result_t MsanInterceptor::prepareLaunch( if (LaunchInfo.LocalWorkSize.empty()) { LaunchInfo.LocalWorkSize.resize(LaunchInfo.WorkDim); auto URes = getContext()->urDdiTable.Kernel.pfnGetSuggestedLocalWorkSize( - Kernel, Queue, LaunchInfo.WorkDim, LaunchInfo.GlobalWorkOffset, + Kernel, Queue, LaunchInfo.WorkDim, LaunchInfo.GlobalWorkOffset.data(), LaunchInfo.GlobalWorkSize, LaunchInfo.LocalWorkSize.data()); if (URes != UR_RESULT_SUCCESS) { if (URes != UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { diff --git a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.hpp b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.hpp index 6c020974a0004..37708e5ca80b7 100644 --- a/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.hpp @@ -225,7 +225,7 @@ struct USMLaunchInfo { ur_context_handle_t Context = nullptr; ur_device_handle_t Device = nullptr; const size_t *GlobalWorkSize = nullptr; - const size_t *GlobalWorkOffset = nullptr; + std::vector GlobalWorkOffset; std::vector LocalWorkSize; uint32_t WorkDim = 0; @@ -233,12 +233,19 @@ struct USMLaunchInfo { const size_t *GlobalWorkSize, const size_t *LocalWorkSize, const size_t *GlobalWorkOffset, uint32_t WorkDim) : Data(Context, Device), Context(Context), Device(Device), - GlobalWorkSize(GlobalWorkSize), GlobalWorkOffset(GlobalWorkOffset), - WorkDim(WorkDim) { + GlobalWorkSize(GlobalWorkSize), WorkDim(WorkDim) { if (LocalWorkSize) { this->LocalWorkSize = std::vector(LocalWorkSize, LocalWorkSize + WorkDim); } + // UR doesn't allow GlobalWorkOffset is null, we need to construct a zero + // value array if user doesn't specify its value. + if (GlobalWorkOffset) { + this->GlobalWorkOffset = + std::vector(GlobalWorkOffset, GlobalWorkOffset + WorkDim); + } else { + this->GlobalWorkOffset = std::vector(WorkDim, 0); + } } ~USMLaunchInfo(); diff --git a/unified-runtime/source/loader/layers/sanitizer/tsan/tsan_interceptor.cpp b/unified-runtime/source/loader/layers/sanitizer/tsan/tsan_interceptor.cpp index 3f9248489f52b..d872025b7d9a9 100644 --- a/unified-runtime/source/loader/layers/sanitizer/tsan/tsan_interceptor.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/tsan/tsan_interceptor.cpp @@ -408,7 +408,7 @@ ur_result_t TsanInterceptor::prepareLaunch(std::shared_ptr &, if (LaunchInfo.LocalWorkSize.empty()) { LaunchInfo.LocalWorkSize.resize(LaunchInfo.WorkDim); auto URes = getContext()->urDdiTable.Kernel.pfnGetSuggestedLocalWorkSize( - Kernel, Queue, LaunchInfo.WorkDim, LaunchInfo.GlobalWorkOffset, + Kernel, Queue, LaunchInfo.WorkDim, LaunchInfo.GlobalWorkOffset.data(), LaunchInfo.GlobalWorkSize, LaunchInfo.LocalWorkSize.data()); if (URes != UR_RESULT_SUCCESS) { if (URes != UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { diff --git a/unified-runtime/source/loader/layers/sanitizer/tsan/tsan_interceptor.hpp b/unified-runtime/source/loader/layers/sanitizer/tsan/tsan_interceptor.hpp index eefcba4036c08..b7500da233296 100644 --- a/unified-runtime/source/loader/layers/sanitizer/tsan/tsan_interceptor.hpp +++ b/unified-runtime/source/loader/layers/sanitizer/tsan/tsan_interceptor.hpp @@ -181,7 +181,7 @@ struct LaunchInfo { ur_context_handle_t Context = nullptr; ur_device_handle_t Device = nullptr; const size_t *GlobalWorkSize = nullptr; - const size_t *GlobalWorkOffset = nullptr; + std::vector GlobalWorkOffset; std::vector LocalWorkSize; uint32_t WorkDim = 0; TsanRuntimeDataWrapper Data; @@ -190,8 +190,7 @@ struct LaunchInfo { const size_t *GlobalWorkSize, const size_t *LocalWorkSize, const size_t *GlobalWorkOffset, uint32_t WorkDim) : Context(Context), Device(Device), GlobalWorkSize(GlobalWorkSize), - GlobalWorkOffset(GlobalWorkOffset), WorkDim(WorkDim), - Data(Context, Device) { + WorkDim(WorkDim), Data(Context, Device) { [[maybe_unused]] auto Result = getContext()->urDdiTable.Context.pfnRetain(Context); assert(Result == UR_RESULT_SUCCESS); @@ -201,6 +200,14 @@ struct LaunchInfo { this->LocalWorkSize = std::vector(LocalWorkSize, LocalWorkSize + WorkDim); } + // UR doesn't allow GlobalWorkOffset is null, we need to construct a zero + // value array if user doesn't specify its value. + if (GlobalWorkOffset) { + this->GlobalWorkOffset = + std::vector(GlobalWorkOffset, GlobalWorkOffset + WorkDim); + } else { + this->GlobalWorkOffset = std::vector(WorkDim, 0); + } } ~LaunchInfo() { diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index 244a9f28179d9..464acb714b2a5 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -3864,6 +3864,9 @@ __urdlllocal ur_result_t UR_APICALL urKernelGetSuggestedLocalWorkSize( } if (getContext()->enableParameterValidation) { + if (NULL == pGlobalWorkOffset) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + if (NULL == pGlobalWorkSize) return UR_RESULT_ERROR_INVALID_NULL_POINTER; From 02c491888afed90b253a5a0e3578dc92efe82b2b Mon Sep 17 00:00:00 2001 From: "Zhao, Maosu" Date: Wed, 12 Nov 2025 06:32:34 +0100 Subject: [PATCH 3/3] update test --- .../invalid-argument/{bad-context.cpp => bad-device.cpp} | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) rename sycl/test-e2e/AddressSanitizer/invalid-argument/{bad-context.cpp => bad-device.cpp} (98%) diff --git a/sycl/test-e2e/AddressSanitizer/invalid-argument/bad-context.cpp b/sycl/test-e2e/AddressSanitizer/invalid-argument/bad-device.cpp similarity index 98% rename from sycl/test-e2e/AddressSanitizer/invalid-argument/bad-context.cpp rename to sycl/test-e2e/AddressSanitizer/invalid-argument/bad-device.cpp index 7b255417f09e4..36743610d7187 100644 --- a/sycl/test-e2e/AddressSanitizer/invalid-argument/bad-context.cpp +++ b/sycl/test-e2e/AddressSanitizer/invalid-argument/bad-device.cpp @@ -17,7 +17,7 @@ int main() { }); gpu_queue.wait(); // CHECK: DeviceSanitizer: invalid-argument on kernel - // CHECK: The {{[0-9]+}}th argument {{.*}} is allocated in other context + // CHECK: The {{[0-9]+}}th argument {{.*}} is allocated in other device // CHECK: {{.*}} is located inside of Device USM region sycl::free(data, cpu_queue);