[UR][HIP] Opt into the larger per-block shared-memory pool

[zjin-lcf]

Summary

The HIP adapter reported hipDeviceAttributeMaxSharedMemoryPerBlock as UR_DEVICE_INFO_LOCAL_MEM_SIZE and called hipFuncSetAttribute(hipFuncAttributeMaxDynamicSharedMemorySize) only when UR_HIP_MAX_LOCAL_MEM_SIZE / SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE was set. That matches neither the CUDA adapter nor HIP on NVIDIA.

On NVIDIA, HIP forwards these calls to CUDA. The default per-block limit stays 48 KiB, and the opt-in ceiling is larger (measured: T4 64 KiB, RTX A2000 99 KiB, H100 PCIe 227 KiB). A launch above 48 KiB has to opt in with cudaFuncSetAttribute / hipFuncSetAttribute. The CUDA adapter already does this; the HIP adapter did not, so a SYCL program on the HIP backend could not use that pool.

This change:

• Queries hipDeviceAttributeSharedMemPerBlockOptin and reports the larger of that value and the non-opt-in size as local_mem_size.
• On enqueue, sets hipFuncAttributeMaxDynamicSharedMemorySize when the kernel's dynamic local size exceeds the non-opt-in limit.
• Keeps UR_HIP_MAX_LOCAL_MEM_SIZE as a cap on that ceiling.

A failed query or a zero opt-in result keeps the non-opt-in size. ROCm 5.7 rejects the attribute on AMD, and ROCm 6.2 returns 0. ROCm 6.4+ reports the opt-in value, but it is the same 64 KiB LDS size as the non-opt-in limit for every ISA in the ROCclr table (gfx906, gfx90a, gfx942, and the RDNA targets). On those AMD devices this does not grow the pool; it also does not shrink local_mem_size to 0.

Fixes #23365.

Test plan

• MI210 (gfx90a, ROCm 7.1.1), both devices, with the patched libur_adapter_hip.so: local_mem_size is 65536; 1 KiB, 48 KiB, and 64 KiB local_accessor launches read back the first and last byte; 64 KiB + 1 KiB is rejected before launch.
• UR_HIP_MAX_LOCAL_MEM_SIZE=32768 lowers both the query and the launch limit. A cap above the device size is clamped to 65536.
• Native HIP attribute query on MI210 and MI50/MI60: opt-in equals the non-opt-in 64 KiB under ROCm 6.4/7.1, is 0 under ROCm 6.2, and the query is rejected under ROCm 5.7. A 64 KiB dynamic launch succeeds without hipFuncSetAttribute; 65 KiB fails with or without it.
• gfx942 (MI300A, ROCm 7.2.4), patched adapter, devices 0 and 3: opt-in equals the non-opt-in 65536, so sycl/test-e2e/Adapters/hip-local-mem-optin.cpp skips the default + 1 launch. local_mem_size is 65536; a 64 KiB launch reads back the first and last byte; 64 KiB + 1 is rejected. UR_HIP_MAX_LOCAL_MEM_SIZE=32768 lowers both the query and the launch limit.
• HIP program on an H100 PCIe (HIP_PLATFORM=nvidia, ROCm 7.1.1 HIP, CUDA 12.8). hipDeviceAttributeMaxSharedMemoryPerBlock is 49152 and hipDeviceAttributeSharedMemPerBlockOptin is 232448 (227 KiB). A dynamic launch of 48 KiB succeeds with no hipFuncSetAttribute. 49153 bytes fails with hipErrorInvalidValue unless hipFuncAttributeMaxDynamicSharedMemorySize is set to the launch size (or once, to 232448). The opted-in launch then succeeds through 232448 bytes and reads back the first and last byte; 232449 is rejected. This is a native HIP program, not a SYCL HIP-adapter launch.
• HIP program on MI210 (gfx90a, ROCm 7.1.1). Both attributes are 65536. A 64 KiB dynamic launch succeeds with no hipFuncSetAttribute. 65537 fails with or without it, and setting the attribute to 65537 is itself hipErrorInvalidValue.

Made with Cursor

[Copilot]

🟡 Changes recommended

The central HIP-on-NVIDIA opt-in launch path lacks automated coverage.

1 open finding

• Add HIP test coverage for opt-in local memory branch · New

What changed in this PR

Aligns the HIP adapter with CUDA’s opt-in shared-memory behavior.

Changes:

• Reports the larger opt-in local-memory ceiling with fallback.
• Opts kernels into larger dynamic shared memory when needed.
• Preserves environment-variable caps.

File	Description
hip/​enqueue.cpp	Enables per-kernel opt-in shared memory.
hip/​device.hpp	Queries and stores default and opt-in limits.
hip/​device.cpp	Documents the reported local-memory ceiling.

🧠 Review effort: Balanced

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[Copilot]

Warning

Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.

Copilot review overview

6 open findings

• This test implicitly relies on the default SYCL device selection to pick the HIP backend device… · New
• This test implicitly relies on the default SYCL device selection to pick the HIP backend device… · New
• This path sets hipFuncAttributeMaxDynamicSharedMemorySize to LocalSize without guarding against… · New
• Add HIP test coverage for opt-in local memory branch
• The comment about later hipDeviceGetAttribute calls being 'poisoned' is potentially misleading:… · New
• Calling hipFuncSetAttribute on every enqueue can add measurable overhead for workloads with many… · New

🧠 Review effort: Lite

[Copilot]

🟡 Changes recommended

Launch correctness issues remain in NVIDIA API usage, shared-function synchronization, and static-memory accounting, alongside missing test SDK gating.

10 open findings

• HIP test must require the hip_dev_kit feature · New
• Use the NVIDIA driver API for module function attributes · New
• Race updating shared function limits across concurrent launches · New
• Include static shared memory when checking launch limits · New
• This path sets hipFuncAttributeMaxDynamicSharedMemorySize to LocalSize without guarding against…
• This test implicitly relies on the default SYCL device selection to pick the HIP backend device…
• This test implicitly relies on the default SYCL device selection to pick the HIP backend device…
• Add HIP test coverage for opt-in local memory branch
• Calling hipFuncSetAttribute on every enqueue can add measurable overhead for workloads with many…
• The comment about later hipDeviceGetAttribute calls being 'poisoned' is potentially misleading:…

🧠 Review effort: Balanced

[ldorau]

Please fix failing CI builds

[zjin-lcf]

Addressed the latest review in 475f087:

• select the HIP backend explicitly and require hip_dev_kit in the E2E test;
• validate static + dynamic shared memory before opting in;
• configure one stable, capacity-bounded dynamic limit per module function, shared across UR kernel handles;
• use cuFuncSetAttribute for NVIDIA module functions and retain hipFuncSetAttribute for AMD;
• cache the static-memory query/configuration and clarify the HIP last-error comment.

Hardware validation on gilgamesh:

• MI210 (gfx90a, ROCm 7.1.1): the updated AMD HIP adapter builds successfully. The E2E test passes uncapped (65536-byte launch succeeds, 65537 is rejected) and with UR_HIP_MAX_LOCAL_MEM_SIZE=32768 (32768 succeeds, 32769 is rejected).
• H100 PCIe (CUDA 12.8, ROCm 7.1.1 HIP headers): a module-function probe using hipModuleGetFunction + cuFuncSetAttribute reports default 49152 and opt-in 232448. A 232448-byte dynamic launch succeeds. A second module kernel reports 16384 bytes of static shared memory; setting the stable dynamic maximum to 232448 - 16384 makes a 40 KiB dynamic launch succeed, covering the static+dynamic threshold case.

The changed NVIDIA program.cpp compiles with __HIP_PLATFORM_NVIDIA__. A complete local HIP-on-NVIDIA adapter link was not available because this checkout has unrelated existing build failures in image/command-buffer code against the installed HIP-on-NVIDIA headers; the H100 validation therefore exercised the exact driver/module API path directly rather than through a linked UR adapter.

[zjin-lcf]

The actionable Windows CI failure was the new HIP test being selected by the build-only Windows configuration: %{build} targeted SPIR-V while %hip_options added --rocm-path, which became an unused-argument error under -Werror. Commit cc63e0a marks this HIP hardware test unsupported on Windows. The separate BMG job ran all 1,359 selected tests successfully; it failed during cleanup because an unrelated bindless_images/aligned_alloc_device_2d process remained hung and kept the build directory busy.