From 502395c4499e07af992fd34890c0de358e35d963 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 8 Oct 2026 14:00:17 -0700 Subject: [PATCH 1/5] [UR][HIP] Opt into the larger per-block shared-memory pool The HIP adapter reported only the non-opt-in shared-memory size and set hipFuncAttributeMaxDynamicSharedMemorySize only when UR_HIP_MAX_LOCAL_MEM_SIZE was set. On NVIDIA that hides the opt-in pool above 48 KiB. Query the opt-in ceiling and set the attribute when a launch exceeds the default limit, matching the CUDA adapter. A missing or zero opt-in result keeps the non-opt-in size. Co-authored-by: Cursor --- .../source/adapters/hip/device.cpp | 4 ++- .../source/adapters/hip/device.hpp | 32 ++++++++++++++++++- .../source/adapters/hip/enqueue.cpp | 8 +++++ 3 files changed, 42 insertions(+), 2 deletions(-) diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp index 8fb95f00c9a40..02fb0132ab47c 100644 --- a/unified-runtime/source/adapters/hip/device.cpp +++ b/unified-runtime/source/adapters/hip/device.cpp @@ -402,7 +402,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_LOCAL_MEM_SIZE: { // OpenCL's "local memory" maps most closely to HIP's "shared memory". // HIP has its own definition of "local memory", which maps to OpenCL's - // "private memory". + // "private memory". The size is the opt-in per-block ceiling + // (hipDeviceAttributeSharedMemPerBlockOptin, or the non-opt-in size when + // that attribute is absent). UR_HIP_MAX_LOCAL_MEM_SIZE caps it. if (hDevice->getMaxChosenLocalMem()) { return ReturnValue( static_cast(hDevice->getMaxChosenLocalMem())); diff --git a/unified-runtime/source/adapters/hip/device.hpp b/unified-runtime/source/adapters/hip/device.hpp index 639167ce6ce78..a23ecb811e53a 100644 --- a/unified-runtime/source/adapters/hip/device.hpp +++ b/unified-runtime/source/adapters/hip/device.hpp @@ -13,6 +13,8 @@ #include +#include + /// UR device mapping to a hipDevice_t. /// Includes an observer pointer to the platform, /// and implements the reference counting semantics since @@ -28,7 +30,11 @@ struct ur_device_handle_t_ : ur::hip::handle_base { int MaxWorkGroupSize{0}; size_t MaxBlockDim[3]; + // Opt-in per-block shared memory ceiling reported as local_mem_size. int MaxCapacityLocalMem{0}; + // Non-opt-in per-block shared memory. Dynamic local memory above this + // requires hipFuncAttributeMaxDynamicSharedMemorySize. + int MaxDefaultLocalMem{0}; int MaxChosenLocalMem{0}; int ManagedMemSupport{0}; int ConcurrentManagedAccess{0}; @@ -54,9 +60,31 @@ struct ur_device_handle_t_ : ur::hip::handle_base { &MaxDim, hipDeviceAttributeMaxBlockDimZ, HIPDevice)); MaxBlockDim[2] = size_t(MaxDim); + int MaxSharedMemPerBlock = 0; UR_CHECK_ERROR(hipDeviceGetAttribute( - &MaxCapacityLocalMem, hipDeviceAttributeMaxSharedMemoryPerBlock, + &MaxSharedMemPerBlock, hipDeviceAttributeMaxSharedMemoryPerBlock, HIPDevice)); + // Default dynamic shared-memory limit. Kernels that need more must opt in + // via hipFuncAttributeMaxDynamicSharedMemorySize, up to the opt-in ceiling. + MaxDefaultLocalMem = MaxSharedMemPerBlock; + + int MaxSharedMemOptin = 0; + hipError_t OptinStatus = hipDeviceGetAttribute( + &MaxSharedMemOptin, hipDeviceAttributeSharedMemPerBlockOptin, + HIPDevice); + if (OptinStatus != hipSuccess) { + // ROCm 5.7 rejects this attribute on AMD (hipErrorInvalidValue). Clear + // the sticky error so later hipDeviceGetAttribute calls are not poisoned. + (void)hipGetLastError(); + MaxSharedMemOptin = 0; + } + // ROCm 6.2 accepts the query on gfx90a but returns 0. CDNA parts that do + // report an opt-in value (ROCm 6.4+) report the same 64KiB LDS size as + // the non-opt-in limit, not a larger pool. Keep the larger of the two so + // a missing or zero opt-in cannot hide memory the device already exposes. + if (MaxSharedMemOptin <= 0) + MaxSharedMemOptin = MaxSharedMemPerBlock; + MaxCapacityLocalMem = std::max(MaxSharedMemPerBlock, MaxSharedMemOptin); UR_CHECK_ERROR(hipDeviceGetAttribute( &ManagedMemSupport, hipDeviceAttributeManagedMemory, HIPDevice)); UR_CHECK_ERROR(hipDeviceGetAttribute( @@ -114,6 +142,8 @@ struct ur_device_handle_t_ : ur::hip::handle_base { int getMaxCapacityLocalMem() const noexcept { return MaxCapacityLocalMem; }; + int getMaxDefaultLocalMem() const noexcept { return MaxDefaultLocalMem; }; + int getMaxChosenLocalMem() const noexcept { return MaxChosenLocalMem; }; int getManagedMemSupport() const noexcept { return ManagedMemSupport; }; diff --git a/unified-runtime/source/adapters/hip/enqueue.cpp b/unified-runtime/source/adapters/hip/enqueue.cpp index d84586f53ee74..ef7ad36c56105 100644 --- a/unified-runtime/source/adapters/hip/enqueue.cpp +++ b/unified-runtime/source/adapters/hip/enqueue.cpp @@ -1878,6 +1878,14 @@ setKernelParams(const ur_device_handle_t Device, const uint32_t WorkDim, UR_CHECK_ERROR(hipFuncSetAttribute( HIPFunc, hipFuncAttributeMaxDynamicSharedMemorySize, MaxLocalMem)); + } else if (LocalSize > + static_cast(Device->getMaxDefaultLocalMem())) { + // Match the CUDA adapter: dynamic shared memory above the default + // per-block limit must be opted in. UR_HIP_MAX_LOCAL_MEM_SIZE is only a + // cap on that ceiling; it is not required to reach the opt-in pool. + UR_CHECK_ERROR(hipFuncSetAttribute( + HIPFunc, hipFuncAttributeMaxDynamicSharedMemorySize, + static_cast(LocalSize))); } } catch (ur_result_t Err) { return Err; From 19db92819dd0ae61a5f9bb704d322aa23755e909 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Wed, 14 Oct 2026 07:26:41 -0700 Subject: [PATCH 2/5] [SYCL][E2E] Test HIP opt-in local memory launches The opt-in attribute path is only reachable when the per-block ceiling exceeds the non-opt-in limit. Skip that launch when the two are equal, and still check the reported size, a launch at that size, and rejection of one extra byte, including a UR_HIP_MAX_LOCAL_MEM_SIZE cap. --- .../test-e2e/Adapters/hip-local-mem-optin.cpp | 177 ++++++++++++++++++ 1 file changed, 177 insertions(+) create mode 100644 sycl/test-e2e/Adapters/hip-local-mem-optin.cpp diff --git a/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp b/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp new file mode 100644 index 0000000000000..0f1ea25de94b0 --- /dev/null +++ b/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp @@ -0,0 +1,177 @@ +// REQUIRES: hip + +// RUN: %{build} -o %t.out %hip_options +// The opt-in launch is skipped when the opt-in ceiling equals the non-opt-in +// limit. Otherwise it uses one byte more than that limit. +// RUN: %{run} %t.out +// A cap below the device ceiling lowers the query and rejects a larger launch. +// RUN: %{run} UR_HIP_MAX_LOCAL_MEM_SIZE=32768 %t.out + +//==------------------------ hip-local-mem-optin.cpp ----------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// Dynamic local memory above the non-opt-in per-block limit has to opt in with +// hipFuncAttributeMaxDynamicSharedMemorySize. AMD devices report the same +// ceiling for both attributes (64KiB LDS on gfx942), so that launch is skipped +// there. The test still checks the reported ceiling, a launch of that size, +// and rejection of one byte more. UR_HIP_MAX_LOCAL_MEM_SIZE caps both. + +#include +#include + +#include +#include +#include +#include + +#if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__) +#define __HIP_PLATFORM_AMD__ +#endif +#include + +namespace { + +int queryAttribute(int Device, hipDeviceAttribute_t Attr, bool *Ok) { + int Value = 0; + hipError_t Err = hipDeviceGetAttribute(&Value, Attr, Device); + if (Err != hipSuccess) { + (void)hipGetLastError(); + *Ok = false; + return 0; + } + *Ok = true; + return Value; +} + +bool launchAndCheck(sycl::queue &Q, size_t Bytes) { + sycl::buffer Out{sycl::range<1>(2)}; + Q.submit([&](sycl::handler &CGH) { + sycl::accessor OutAcc{Out, CGH, sycl::write_only}; + sycl::local_accessor Local{Bytes, CGH}; + CGH.parallel_for(sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1>) { + Local[0] = 0xA5; + Local[Bytes - 1] = 0x5A; + OutAcc[0] = Local[0]; + OutAcc[1] = Local[Bytes - 1]; + }); + }).wait(); + sycl::host_accessor Host{Out, sycl::read_only}; + if (Host[0] != 0xA5 || Host[1] != 0x5A) { + std::printf("launch %zu read back %u %u\n", Bytes, Host[0], Host[1]); + return false; + } + std::printf("launch %zu OK\n", Bytes); + return true; +} + +bool expectReject(sycl::queue &Q, size_t Bytes, const char *Expected) { + try { + launchAndCheck(Q, Bytes); + } catch (const std::exception &E) { + if (std::string(E.what()).find(Expected) == std::string::npos) { + std::printf("launch %zu rejected with unexpected message: %s\n", Bytes, + E.what()); + return false; + } + std::printf("launch %zu rejected as expected\n", Bytes); + return true; + } + std::printf("launch %zu was expected to fail\n", Bytes); + return false; +} + +} // namespace + +int main() { + sycl::queue Q; + const std::string Pci = + Q.get_device().get_info(); + int DeviceCount = 0; + if (hipGetDeviceCount(&DeviceCount) != hipSuccess || DeviceCount <= 0) { + std::printf("failed to enumerate HIP devices\n"); + return 1; + } + int HipDevice = -1; + for (int Index = 0; Index < DeviceCount; ++Index) { + char BusId[32] = {}; + if (hipDeviceGetPCIBusId(BusId, sizeof(BusId), Index) == hipSuccess && + Pci == BusId) { + HipDevice = Index; + break; + } + } + if (HipDevice < 0) { + std::printf("no HIP device matches PCI address %s\n", Pci.c_str()); + return 1; + } + + bool DefaultOk = false; + const int Default = queryAttribute( + HipDevice, hipDeviceAttributeMaxSharedMemoryPerBlock, &DefaultOk); + if (!DefaultOk || Default <= 0) { + std::printf("failed to query non-opt-in shared memory on HIP device %d\n", + HipDevice); + return 1; + } + + bool OptinOk = false; + int Optin = queryAttribute( + HipDevice, hipDeviceAttributeSharedMemPerBlockOptin, &OptinOk); + // Match the adapter: a rejected or non-positive opt-in query keeps the + // non-opt-in size. + if (!OptinOk || Optin <= 0) + Optin = Default; + const int Ceiling = std::max(Default, Optin); + + const auto Reported = static_cast( + Q.get_device().get_info()); + + const char *CapEnv = std::getenv("UR_HIP_MAX_LOCAL_MEM_SIZE"); + size_t Expected = static_cast(Ceiling); + if (CapEnv) { + const int Cap = std::atoi(CapEnv); + if (Cap <= 0) { + std::printf("UR_HIP_MAX_LOCAL_MEM_SIZE is not a positive size\n"); + return 1; + } + Expected = static_cast(std::min(Cap, Ceiling)); + } + + std::printf("device %d default %d opt-in %d ceiling %d reported %zu\n", + HipDevice, Default, Optin, Ceiling, Reported); + if (Reported != Expected) { + std::printf("local_mem_size %zu, expected %zu\n", Reported, Expected); + return 1; + } + + if (Reported == 0 || !launchAndCheck(Q, Reported)) + return 1; + + const char *Reject = + (Reported + 1 > static_cast(Ceiling)) + ? "Excessive allocation of local memory on the device" + : "UR_HIP_MAX_LOCAL_MEM_SIZE"; + if (!expectReject(Q, Reported + 1, Reject)) + return 1; + + // One byte past the non-opt-in limit is the hipFuncSetAttribute path. It is + // unreachable when the ceiling equals that limit (gfx942 and the other AMD + // ISAs) or when the env cap keeps the launch inside the default pool. + const size_t OptInBytes = static_cast(Default) + 1; + if (Ceiling > Default && OptInBytes <= Reported && OptInBytes != Reported) { + if (!launchAndCheck(Q, OptInBytes)) + return 1; + std::printf("opt-in launch %zu OK\n", OptInBytes); + } else if (Ceiling > Default && OptInBytes == Reported) { + std::printf("opt-in launch covered by the %zu-byte launch\n", Reported); + } else { + std::printf("skipped opt-in launch: ceiling %d default %d reported %zu\n", + Ceiling, Default, Reported); + } + return 0; +} From 475f0879eed9ad1b0a0c330b18b584673c9e4fe5 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 15 Oct 2026 07:33:27 -0700 Subject: [PATCH 3/5] [UR][HIP] Address opt-in local memory review Co-authored-by: Cursor --- .../test-e2e/Adapters/hip-local-mem-optin.cpp | 6 +-- .../source/adapters/hip/device.hpp | 2 +- .../source/adapters/hip/enqueue.cpp | 31 +---------- .../source/adapters/hip/program.cpp | 52 +++++++++++++++++++ .../source/adapters/hip/program.hpp | 9 ++++ 5 files changed, 67 insertions(+), 33 deletions(-) diff --git a/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp b/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp index 0f1ea25de94b0..b3976a5f92430 100644 --- a/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp +++ b/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp @@ -1,11 +1,11 @@ -// REQUIRES: hip +// REQUIRES: hip, hip_dev_kit // RUN: %{build} -o %t.out %hip_options // The opt-in launch is skipped when the opt-in ceiling equals the non-opt-in // limit. Otherwise it uses one byte more than that limit. -// RUN: %{run} %t.out +// RUN: env SYCL_DEVICE_FILTER=hip %{run} %t.out // A cap below the device ceiling lowers the query and rejects a larger launch. -// RUN: %{run} UR_HIP_MAX_LOCAL_MEM_SIZE=32768 %t.out +// RUN: env SYCL_DEVICE_FILTER=hip UR_HIP_MAX_LOCAL_MEM_SIZE=32768 %{run} %t.out //==------------------------ hip-local-mem-optin.cpp ----------------------===// // diff --git a/unified-runtime/source/adapters/hip/device.hpp b/unified-runtime/source/adapters/hip/device.hpp index a23ecb811e53a..65b86aa882405 100644 --- a/unified-runtime/source/adapters/hip/device.hpp +++ b/unified-runtime/source/adapters/hip/device.hpp @@ -74,7 +74,7 @@ struct ur_device_handle_t_ : ur::hip::handle_base { HIPDevice); if (OptinStatus != hipSuccess) { // ROCm 5.7 rejects this attribute on AMD (hipErrorInvalidValue). Clear - // the sticky error so later hipDeviceGetAttribute calls are not poisoned. + // HIP's last-error state so it does not contaminate a later error check. (void)hipGetLastError(); MaxSharedMemOptin = 0; } diff --git a/unified-runtime/source/adapters/hip/enqueue.cpp b/unified-runtime/source/adapters/hip/enqueue.cpp index ef7ad36c56105..0d6f81a0bbfd8 100644 --- a/unified-runtime/source/adapters/hip/enqueue.cpp +++ b/unified-runtime/source/adapters/hip/enqueue.cpp @@ -1858,35 +1858,8 @@ setKernelParams(const ur_device_handle_t Device, const uint32_t WorkDim, Kernel->setImplicitOffsetArg(sizeof(ImplicitOffset), ImplicitOffset); } - uint32_t LocalSize = Kernel->getLocalSize(); - if (LocalSize > static_cast(Device->getMaxCapacityLocalMem())) { - setErrorMessage("Excessive allocation of local memory on the device", - UR_RESULT_ERROR_OUT_OF_RESOURCES); - return UR_RESULT_ERROR_ADAPTER_SPECIFIC; - } - - if (int MaxLocalMem = Device->getMaxChosenLocalMem()) { - if (LocalSize > static_cast(MaxLocalMem)) { - setErrorMessage( - "Local memory for kernel exceeds the amount requested using " - "UR_HIP_MAX_LOCAL_MEM_SIZE (or deprecated " - "SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE). Try increasing the maximum " - "local memory.", - UR_RESULT_ERROR_OUT_OF_RESOURCES); - return UR_RESULT_ERROR_ADAPTER_SPECIFIC; - } - - UR_CHECK_ERROR(hipFuncSetAttribute( - HIPFunc, hipFuncAttributeMaxDynamicSharedMemorySize, MaxLocalMem)); - } else if (LocalSize > - static_cast(Device->getMaxDefaultLocalMem())) { - // Match the CUDA adapter: dynamic shared memory above the default - // per-block limit must be opted in. UR_HIP_MAX_LOCAL_MEM_SIZE is only a - // cap on that ceiling; it is not required to reach the opt-in pool. - UR_CHECK_ERROR(hipFuncSetAttribute( - HIPFunc, hipFuncAttributeMaxDynamicSharedMemorySize, - static_cast(LocalSize))); - } + Kernel->getProgram()->prepareFunctionLocalMemory(HIPFunc, + Kernel->getLocalSize()); } catch (ur_result_t Err) { return Err; } diff --git a/unified-runtime/source/adapters/hip/program.cpp b/unified-runtime/source/adapters/hip/program.cpp index d63175baa0aa6..07cc12d8fe874 100644 --- a/unified-runtime/source/adapters/hip/program.cpp +++ b/unified-runtime/source/adapters/hip/program.cpp @@ -10,6 +10,58 @@ #include "program.hpp" #include "ur_util.hpp" +void ur_program_handle_t_::prepareFunctionLocalMemory(hipFunction_t Function, + uint32_t DynamicSize) { + std::lock_guard Lock(FunctionLocalMemMutex); + auto &Info = FunctionLocalMem[Function]; + + if (Info.StaticSize < 0) { + UR_CHECK_ERROR(hipFuncGetAttribute( + &Info.StaticSize, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, Function)); + } + + const uint64_t TotalSize = + static_cast(Info.StaticSize) + DynamicSize; + const int Capacity = Device->getMaxCapacityLocalMem(); + if (TotalSize > static_cast(Capacity)) { + setErrorMessage("Excessive allocation of local memory on the device", + UR_RESULT_ERROR_OUT_OF_RESOURCES); + throw UR_RESULT_ERROR_ADAPTER_SPECIFIC; + } + + const int ChosenSize = Device->getMaxChosenLocalMem(); + if (ChosenSize && DynamicSize > static_cast(ChosenSize)) { + setErrorMessage( + "Local memory for kernel exceeds the amount requested using " + "UR_HIP_MAX_LOCAL_MEM_SIZE (or deprecated " + "SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE). Try increasing the maximum " + "local memory.", + UR_RESULT_ERROR_OUT_OF_RESOURCES); + throw UR_RESULT_ERROR_ADAPTER_SPECIFIC; + } + + if (TotalSize <= static_cast(Device->getMaxDefaultLocalMem()) || + Info.MaxDynamicSizeSet) + return; + + // Set one stable limit for this module function. Multiple UR kernel handles + // can share Function, so changing the limit to each launch's requested size + // would race with launches through the other handles. + const int MaxDynamicSize = + std::min(ChosenSize ? ChosenSize : Capacity, Capacity - Info.StaticSize); +#if defined(__HIP_PLATFORM_NVIDIA__) + // hipModuleGetFunction returns a driver function. HIP's NVIDIA runtime + // hipFuncSetAttribute wrapper expects a runtime-registered kernel symbol. + UR_CHECK_ERROR(hipCUResultTohipError(cuFuncSetAttribute( + Function, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, + MaxDynamicSize))); +#else + UR_CHECK_ERROR(hipFuncSetAttribute( + Function, hipFuncAttributeMaxDynamicSharedMemorySize, MaxDynamicSize)); +#endif + Info.MaxDynamicSizeSet = true; +} + #ifdef SYCL_ENABLE_KERNEL_FUSION #ifdef UR_COMGR_VERSION4_INCLUDE #include diff --git a/unified-runtime/source/adapters/hip/program.hpp b/unified-runtime/source/adapters/hip/program.hpp index 0fd9c33b591b7..829521c039e98 100644 --- a/unified-runtime/source/adapters/hip/program.hpp +++ b/unified-runtime/source/adapters/hip/program.hpp @@ -11,6 +11,7 @@ #include #include +#include #include #include "common/ur_ref_count.hpp" @@ -41,6 +42,13 @@ struct ur_program_handle_t_ : ur::hip::handle_base { KernelReqdWorkGroupSizeMD; std::unordered_map KernelReqdSubGroupSizeMD; + struct FunctionLocalMemInfo { + int StaticSize{-1}; + bool MaxDynamicSizeSet{false}; + }; + std::mutex FunctionLocalMemMutex; + std::unordered_map FunctionLocalMem; + constexpr static size_t MAX_LOG_SIZE = 8192u; char ErrorLog[MAX_LOG_SIZE], InfoLog[MAX_LOG_SIZE]; @@ -66,6 +74,7 @@ struct ur_program_handle_t_ : ur::hip::handle_base { ur_result_t buildProgram(const char *BuildOptions); ur_result_t finalizeRelocatable(); + void prepareFunctionLocalMemory(hipFunction_t Function, uint32_t DynamicSize); ur_context_handle_t getContext() const { return Context; }; ur_device_handle_t getDevice() const { return Device; }; From cc63e0ac110bc210706d4b2dfd8f2d2135d9e06c Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 15 Oct 2026 07:35:29 -0700 Subject: [PATCH 4/5] [SYCL][E2E] Skip HIP local memory test on Windows Co-authored-by: Cursor --- sycl/test-e2e/Adapters/hip-local-mem-optin.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp b/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp index b3976a5f92430..c96be6b0e73b4 100644 --- a/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp +++ b/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp @@ -1,4 +1,5 @@ // REQUIRES: hip, hip_dev_kit +// UNSUPPORTED: windows // RUN: %{build} -o %t.out %hip_options // The opt-in launch is skipped when the opt-in ceiling equals the non-opt-in From f0f665b4291aeb0793fa1dca0a50151e6324a003 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Mon, 19 Oct 2026 16:16:17 -0700 Subject: [PATCH 5/5] [SYCL][E2E] Explain why the HIP local-memory test skips Windows The Windows build compiles this test for spir64, where --rocm-path from %hip_options is unused and fails under -Werror. Co-authored-by: Cursor --- sycl/test-e2e/Adapters/hip-local-mem-optin.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp b/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp index c96be6b0e73b4..e9337785167fd 100644 --- a/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp +++ b/sycl/test-e2e/Adapters/hip-local-mem-optin.cpp @@ -1,5 +1,7 @@ // REQUIRES: hip, hip_dev_kit // UNSUPPORTED: windows +// UNSUPPORTED-INTENDED: %{build} targets spir64, so --rocm-path from +// %hip_options is an unused argument and fails under -Werror. // RUN: %{build} -o %t.out %hip_options // The opt-in launch is skipped when the opt-in ceiling equals the non-opt-in