From ec23f3605785a7b632d8fa31e722076b0004f011 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sat, 18 Jul 2026 07:51:14 -0700 Subject: [PATCH 1/3] [UR][HIP] Implement urKernelSuggestMaxCooperativeGroupCount Implement the cooperative-group-count occupancy query for the HIP (AMD) adapter, which previously returned UR_RESULT_ERROR_UNSUPPORTED_FEATURE and forced the SYCL runtime onto a coarse 0/1 host-side fallback for kernel::ext_oneapi_get_info. As noted in intel/llvm#21803, a direct port of the CUDA implementation does not work: hipOccupancyMaxActiveBlocksPerMultiprocessor expects a host function pointer, whereas our kernel handle (hKernel->get()) is a module-loaded hipFunction_t. Use the module-occupancy entry point hipModuleOccupancyMaxActiveBlocksPerMultiprocessor instead, mirroring the existing hipModuleOccupancyMaxPotentialBlockSize usage in the adapter. The total suggested group count is the per-CU occupancy multiplied by the device multiprocessor (compute-unit) count. When the per-CU occupancy is 0, fall back to reporting 1 group if the requested work-group size, dynamic local memory, and per-block register usage all fit within device limits, matching the CUDA adapter's semantics. Closes: https://github.com/intel/llvm/issues/21803 Co-authored-by: Cursor --- .../source/adapters/hip/kernel.cpp | 74 ++++++++++++++++++- 1 file changed, 70 insertions(+), 4 deletions(-) diff --git a/unified-runtime/source/adapters/hip/kernel.cpp b/unified-runtime/source/adapters/hip/kernel.cpp index de20e31699a66..604aa6193abd6 100644 --- a/unified-runtime/source/adapters/hip/kernel.cpp +++ b/unified-runtime/source/adapters/hip/kernel.cpp @@ -156,10 +156,76 @@ urKernelGetNativeHandle(ur_kernel_handle_t, ur_native_handle_t *) { } UR_APIEXPORT ur_result_t UR_APICALL urKernelSuggestMaxCooperativeGroupCount( - ur_kernel_handle_t /*hKernel*/, ur_device_handle_t /*hDevice*/, - uint32_t /*workDim*/, const size_t * /*pLocalWorkSize*/, - size_t /*dynamicSharedMemorySize*/, uint32_t * /*pGroupCountRet*/) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + ur_kernel_handle_t hKernel, ur_device_handle_t /*hDevice*/, + uint32_t workDim, const size_t *pLocalWorkSize, + size_t dynamicSharedMemorySize, uint32_t *pGroupCountRet) { + UR_ASSERT(hKernel, UR_RESULT_ERROR_INVALID_KERNEL); + + size_t localWorkSize = pLocalWorkSize[0]; + localWorkSize *= (workDim >= 2 ? pLocalWorkSize[1] : 1); + localWorkSize *= (workDim == 3 ? pLocalWorkSize[2] : 1); + + // We need to set the active current device for this kernel explicitly here, + // because the occupancy querying API does not take a device parameter. + ur_device_handle_t Device = hKernel->getProgram()->getDevice(); + ScopedDevice Active(Device); + try { + // We need to calculate max num of work-groups using per-device semantics. + + // Unlike CUDA's cuOccupancyMaxActiveBlocksPerMultiprocessor, the + // hipOccupancy* variant expects a host function pointer; our kernel is a + // module-loaded hipFunction_t (hKernel->get()), so we must use the + // hipModule* occupancy entry point instead (see issue #21803). + int MaxNumActiveGroupsPerCU{0}; + UR_CHECK_ERROR(hipModuleOccupancyMaxActiveBlocksPerMultiprocessor( + &MaxNumActiveGroupsPerCU, hKernel->get(), + static_cast(localWorkSize), dynamicSharedMemorySize)); + assert(MaxNumActiveGroupsPerCU >= 0); + // Handle the case where we can't have all CUs active with at least 1 group + // per CU. In that case, the device is still able to run 1 work-group, hence + // we will manually check if it is possible with the available HW resources. + if (MaxNumActiveGroupsPerCU == 0) { + size_t MaxWorkGroupSize{}; + urKernelGetGroupInfo( + hKernel, Device, UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE, + sizeof(MaxWorkGroupSize), &MaxWorkGroupSize, nullptr); + size_t MaxLocalSizeBytes{}; + urDeviceGetInfo(Device, UR_DEVICE_INFO_LOCAL_MEM_SIZE, + sizeof(MaxLocalSizeBytes), &MaxLocalSizeBytes, nullptr); + // Check whether the kernel's per-thread register usage would exceed the + // device's per-block register budget for the requested work-group size. + int NumRegs{0}; + UR_CHECK_ERROR(hipFuncGetAttribute(&NumRegs, HIP_FUNC_ATTRIBUTE_NUM_REGS, + hKernel->get())); + int MaxRegsPerBlock{0}; + UR_CHECK_ERROR(hipDeviceGetAttribute( + &MaxRegsPerBlock, hipDeviceAttributeMaxRegistersPerBlock, + Device->get())); + const bool HasExceededMaxRegistersPerBlock = + localWorkSize * static_cast(NumRegs) > + static_cast(MaxRegsPerBlock); + if (localWorkSize > MaxWorkGroupSize || + dynamicSharedMemorySize > MaxLocalSizeBytes || + HasExceededMaxRegistersPerBlock) + *pGroupCountRet = 0; + else + *pGroupCountRet = 1; + } else { + // Multiply by the number of CUs (compute units = streaming + // multiprocessors) on the device in order to retrieve the total number + // of groups/blocks that can be launched. + int NumComputeUnits{0}; + UR_CHECK_ERROR(hipDeviceGetAttribute( + &NumComputeUnits, hipDeviceAttributeMultiprocessorCount, + Device->get())); + assert(NumComputeUnits >= 0); + *pGroupCountRet = + static_cast(NumComputeUnits) * MaxNumActiveGroupsPerCU; + } + } catch (ur_result_t Err) { + return Err; + } + return UR_RESULT_SUCCESS; } UR_APIEXPORT ur_result_t UR_APICALL urKernelGetInfo(ur_kernel_handle_t hKernel, From 3ce7d533f300161f0fb0fc626472010c76fe1203 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sat, 18 Jul 2026 08:25:30 -0700 Subject: [PATCH 2/3] [UR][HIP][libclc] Support cooperative kernel launch and root-group barrier The HIP adapter previously ignored the UR_KERNEL_LAUNCH_FLAG_COOPERATIVE launch property and always issued an ordinary launch, so the work-groups of a root-group kernel were not guaranteed to be co-resident. In addition, the amdgpu libspirv __spirv_ControlBarrier always emitted a work-group scope barrier regardless of the requested execution scope, so a Device/CrossDevice (root-group) barrier only synchronized within a work-group. Issue a cooperative launch via hipModuleLaunchCooperativeKernel when the COOPERATIVE flag is set, and lower a Device/CrossDevice execution scope barrier to the ROCm device library grid synchronization (__ockl_grid_sync), which is valid because the launch is now cooperative. This makes root-group synchronization work end-to-end on AMD GPUs. Co-authored-by: Cursor --- .../lib/amdgpu/synchronization/barrier.cl | 13 +++++++++- .../source/adapters/hip/enqueue.cpp | 25 +++++++++++++------ 2 files changed, 29 insertions(+), 9 deletions(-) diff --git a/libclc/libspirv/lib/amdgpu/synchronization/barrier.cl b/libclc/libspirv/lib/amdgpu/synchronization/barrier.cl index 49cd4a3926577..db25e48a7e75a 100644 --- a/libclc/libspirv/lib/amdgpu/synchronization/barrier.cl +++ b/libclc/libspirv/lib/amdgpu/synchronization/barrier.cl @@ -57,11 +57,22 @@ _CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(int scope_memory, __mem_fence(scope_memory, semantics); } +// Grid-wide synchronization provided by the ROCm device library (ockl). It is +// only valid when the kernel is launched cooperatively (all work-groups are +// co-resident), which the HIP UR adapter guarantees for root-group barriers. +extern _CLC_CONVERGENT void __ockl_grid_sync(void); + _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT void __spirv_ControlBarrier(int scope_execution, int scope_memory, int semantics) { if (semantics) { __mem_fence(scope_memory, semantics); } - __builtin_amdgcn_s_barrier(); + // A Device/CrossDevice execution scope corresponds to a root-group barrier + // that must synchronize every work-item in the grid, not just the work-group. + if (scope_execution == Device || scope_execution == CrossDevice) { + __ockl_grid_sync(); + } else { + __builtin_amdgcn_s_barrier(); + } } diff --git a/unified-runtime/source/adapters/hip/enqueue.cpp b/unified-runtime/source/adapters/hip/enqueue.cpp index d84586f53ee74..e08923fa22c92 100644 --- a/unified-runtime/source/adapters/hip/enqueue.cpp +++ b/unified-runtime/source/adapters/hip/enqueue.cpp @@ -260,10 +260,10 @@ static ur_result_t urEnqueueKernelLaunch( ur_kernel_launch_ext_properties_t *_launchPropList = const_cast(launchPropList); - // Adapters that don't support cooperative kernels are currently expected - // to ignore COOPERATIVE launch properties. Ideally we should avoid passing - // these at the SYCL RT level instead, see - // https://github.com/intel/llvm/issues/18421 + // HIP supports cooperative kernel launches through + // hipModuleLaunchCooperativeKernel. Any other launch property flag is + // unsupported, see https://github.com/intel/llvm/issues/18421 + bool UseCooperativeLaunch = false; if (_launchPropList && _launchPropList->flags & ~UR_KERNEL_LAUNCH_FLAG_COOPERATIVE) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; @@ -274,6 +274,8 @@ static ur_result_t urEnqueueKernelLaunch( as_stype()) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } + if (_launchPropList->flags & UR_KERNEL_LAUNCH_FLAG_COOPERATIVE) + UseCooperativeLaunch = true; _launchPropList = static_cast( _launchPropList->pNext); } @@ -331,10 +333,17 @@ static ur_result_t urEnqueueKernelLaunch( UR_CHECK_ERROR(RetImplEvent->start()); } - UR_CHECK_ERROR(hipModuleLaunchKernel( - HIPFunc, BlocksPerGrid[0], BlocksPerGrid[1], BlocksPerGrid[2], - ThreadsPerBlock[0], ThreadsPerBlock[1], ThreadsPerBlock[2], - hKernel->getLocalSize(), HIPStream, ArgPointers.data(), nullptr)); + if (UseCooperativeLaunch) { + UR_CHECK_ERROR(hipModuleLaunchCooperativeKernel( + HIPFunc, BlocksPerGrid[0], BlocksPerGrid[1], BlocksPerGrid[2], + ThreadsPerBlock[0], ThreadsPerBlock[1], ThreadsPerBlock[2], + hKernel->getLocalSize(), HIPStream, ArgPointers.data())); + } else { + UR_CHECK_ERROR(hipModuleLaunchKernel( + HIPFunc, BlocksPerGrid[0], BlocksPerGrid[1], BlocksPerGrid[2], + ThreadsPerBlock[0], ThreadsPerBlock[1], ThreadsPerBlock[2], + hKernel->getLocalSize(), HIPStream, ArgPointers.data(), nullptr)); + } if (phEvent) { UR_CHECK_ERROR(RetImplEvent->record()); From bceb53b35e97738dbed3b3f4a69960b309f337a0 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sat, 18 Jul 2026 08:26:09 -0700 Subject: [PATCH 3/3] [UR][HIP] Support work_group_scratch_size launch property Free-function kernels request dynamic work-group local memory through a UR_STRUCTURE_TYPE_KERNEL_LAUNCH_WORKGROUP_PROPERTY launch property node. The HIP adapter previously rejected any launch property node other than the cooperative flag with UR_RESULT_ERROR_UNSUPPORTED_FEATURE, so such launches failed. Parse the work-group launch property and add the requested amount to the dynamic shared memory passed to the kernel launch, matching the CUDA adapter. Co-authored-by: Cursor --- .../source/adapters/hip/enqueue.cpp | 35 +++++++++++++------ 1 file changed, 24 insertions(+), 11 deletions(-) diff --git a/unified-runtime/source/adapters/hip/enqueue.cpp b/unified-runtime/source/adapters/hip/enqueue.cpp index e08923fa22c92..f51bb094ae44a 100644 --- a/unified-runtime/source/adapters/hip/enqueue.cpp +++ b/unified-runtime/source/adapters/hip/enqueue.cpp @@ -261,21 +261,31 @@ static ur_result_t urEnqueueKernelLaunch( ur_kernel_launch_ext_properties_t *_launchPropList = const_cast(launchPropList); // HIP supports cooperative kernel launches through - // hipModuleLaunchCooperativeKernel. Any other launch property flag is + // hipModuleLaunchCooperativeKernel and dynamic work-group local memory + // requested via a work-group launch property. Any other launch property is // unsupported, see https://github.com/intel/llvm/issues/18421 bool UseCooperativeLaunch = false; - if (_launchPropList && - _launchPropList->flags & ~UR_KERNEL_LAUNCH_FLAG_COOPERATIVE) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; - } + size_t WorkGroupMemory = 0; while (_launchPropList != nullptr) { - if (_launchPropList->stype != - as_stype()) { + switch (_launchPropList->stype) { + case UR_STRUCTURE_TYPE_KERNEL_LAUNCH_EXT_PROPERTIES: { + if (_launchPropList->flags & ~UR_KERNEL_LAUNCH_FLAG_COOPERATIVE) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + if (_launchPropList->flags & UR_KERNEL_LAUNCH_FLAG_COOPERATIVE) + UseCooperativeLaunch = true; + break; + } + case UR_STRUCTURE_TYPE_KERNEL_LAUNCH_WORKGROUP_PROPERTY: { + WorkGroupMemory = + reinterpret_cast( + _launchPropList) + ->workgroup_mem_size; + break; + } + default: return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } - if (_launchPropList->flags & UR_KERNEL_LAUNCH_FLAG_COOPERATIVE) - UseCooperativeLaunch = true; _launchPropList = static_cast( _launchPropList->pNext); } @@ -333,16 +343,19 @@ static ur_result_t urEnqueueKernelLaunch( UR_CHECK_ERROR(RetImplEvent->start()); } + // Dynamic work-group local memory requested through a launch property is + // added on top of the static local memory of the kernel arguments. + uint32_t SharedMemBytes = hKernel->getLocalSize() + WorkGroupMemory; if (UseCooperativeLaunch) { UR_CHECK_ERROR(hipModuleLaunchCooperativeKernel( HIPFunc, BlocksPerGrid[0], BlocksPerGrid[1], BlocksPerGrid[2], ThreadsPerBlock[0], ThreadsPerBlock[1], ThreadsPerBlock[2], - hKernel->getLocalSize(), HIPStream, ArgPointers.data())); + SharedMemBytes, HIPStream, ArgPointers.data())); } else { UR_CHECK_ERROR(hipModuleLaunchKernel( HIPFunc, BlocksPerGrid[0], BlocksPerGrid[1], BlocksPerGrid[2], ThreadsPerBlock[0], ThreadsPerBlock[1], ThreadsPerBlock[2], - hKernel->getLocalSize(), HIPStream, ArgPointers.data(), nullptr)); + SharedMemBytes, HIPStream, ArgPointers.data(), nullptr)); } if (phEvent) {