Declare indirect SHARED access on Level Zero and tag managed memory as host USM on OpenCL - #1515
Merged
Merged
Conversation
…on Level Zero A kernel whose only path to a hipMallocManaged buffer is a pointer stored in other memory (a member of a by-value struct argument, or an entry of a device array of pointers) loses its write on Intel Data Center GPU Max under Level Zero: the host reads the old value after hipDeviceSynchronize and after a hipMemcpy readback. Managed memory there is shared USM since e7d1f6e, and launchImpl only declares DEVICE and HOST indirect access, so the driver never migrates the shared allocation for the launch. Kokkos hip.mixed_then_host_device_nodes (4 instead of 6) and hip.interact_with_hip_node (0 instead of 1) are this bug: a View inside a by-value functor. The test prints the runtime's HasNoIGBAs verdict for the module and fails if the IGBA detector did not flag it, so a pass is known to have gone through the zeKernelSetIndirectAccess path.
…al buffer accesses launchImpl declared ZE_KERNEL_INDIRECT_ACCESS_FLAG_DEVICE | HOST for a kernel whose module the IGBA detector flagged, but never SHARED. Since e7d1f6e hipMallocManaged is backed by zeMemAllocShared on Intel Data Center GPU Max, and the driver only migrates shared allocations for a launch when the kernel declares indirect shared access (KernelImp::setIndirectAccess -> indirectSharedAllocationsAllowed -> CommandList::migrateSharedAllocations). A managed buffer reached only through a pointer stored in a by-value struct argument or another buffer was therefore updated on a stale copy and the write lost: Kokkos hip.mixed_then_host_device_nodes got 4 instead of 6 and hip.interact_with_hip_node read 0 instead of 1. Add SHARED unconditionally, mirroring the OpenCL backend's CL_KERNEL_EXEC_INFO_INDIRECT_{HOST,DEVICE,SHARED}_ACCESS_INTEL; it is a no-op where no shared allocations exist, and on devices whose managed memory stays host USM nothing changes. Log the call once per kernel at info level so a run shows the flags were applied. Level Zero spec, scripts/core/module.yml, ze_kernel_indirect_access_flags_t: SHARED "Indicates that the kernel accesses shared allocations indirectly"; zeKernelSetIndirectAccess: "The application should specify which allocations will be indirectly accessed by the kernel to allow driver to optimize which allocations are made resident".
…ccess flags MemoryManager::allocate() tagged hipMemoryTypeManaged and hipMemoryTypeUnified allocations as shared USM, but allocateUSM() backs them with clHostMemAllocINTEL. annotateIndirectPointers() then set CL_KERNEL_EXEC_INFO_INDIRECT_SHARED_ACCESS_INTEL and, unless a hipHostMalloc allocation also existed, left INDIRECT_HOST_ACCESS unset. The Intel CPU OpenCL runtime (2026.20.1.0.12) rejects such a launch with CL_INVALID_OPERATION at clEnqueueNDRangeKernel: a USM pointer listed in CL_KERNEL_EXEC_INFO_USM_PTRS_INTEL must be accompanied by the indirect access flag of its own kind (intel/compute-runtime issue 748, the same rule chipStar already follows for host and device allocations since f06b694). Any module the IGBA detector flags that also touches a hipMallocManaged buffer failed to launch there; TestFixManagedIndirectAccess is such a module. Record managed and unified allocations as host allocations, which is what they are.
… launch TestFixManagedIndirectAccess read the hipMallocManaged counter directly from the host and reset it with a host write. On rusticl managed memory is coarse-grained SVM: host reads after hipDeviceSynchronize return stale data and host writes do not reach the device, so the test failed there (struct path host read 0, table path readback 2) although both indirect writes landed. A control launch with the managed pointer as a direct kernel argument now decides whether the host is coherent on the device under test; the host read checks and host-side resets apply only then, and every path is still checked through a hipMemcpy readback, which #1503 shows stale as well.
…ot a probe TestFixManagedIndirectAccess decided whether to write and read the managed counter from the host by reading it once after a control launch. On Mali-G52 (OpenCL, coarse-grained SVM only) that read returned 1 by chance, the test then reset the counter with an unmapped host write which the next kernel never saw, and both the host read and the hipMemcpy readback of the struct and table launches came back 0; the same binary passed in the translator lane and failed in the native lane, and both fail when run by hand. Host access to coarse-grained SVM outside clEnqueueSVMMap/Unmap is undefined, so no empirical probe can decide it. A pure OpenCL run of the same sequence on Mali loses the write on every repetition with an unmapped host write and keeps it on every repetition with a clEnqueueSVMMemcpy reset, so the indirect access path itself works there in both SPIR-V lanes. Touch managed memory from the host only where the device reports directManagedMemAccessFromHost or concurrentManagedAccess; elsewhere reset and read the counter through hipMemcpy, which is the check that showed the stale device copy in the original Level Zero report. Log which mode applies and print "skipped" instead of a meaningless host read.
Collaborator
Author
|
/run-aurora-ci |
2 similar comments
Collaborator
Author
|
/run-aurora-ci |
Collaborator
Author
|
/run-aurora-ci |
Collaborator
Author
|
Aurora PVC manual run: PASS. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
A kernel whose only path to a
hipMallocManagedbuffer is a pointer stored in other memory (a member of a by-value struct argument, or an entry of a device array of pointers) lost its write on Aurora PVC under Level Zero:launchImpldeclaredZE_KERNEL_INDIRECT_ACCESS_FLAG_DEVICE | HOSTfor IGBA-flagged kernels but never SHARED, so once managed memory is shared USM the driver did not migrate it for the launch. Kokkoship.mixed_then_host_device_nodes(4 instead of 6) andhip.interact_with_hip_node(0 instead of 1) are this bug. Separately, the OpenCL backend tagged managed and unified allocations as shared USM althoughallocateUSM()backs them withclHostMemAllocINTEL, soannotateIndirectPointersset the wrong indirect access flag and the Intel CPU OpenCL runtime rejected the launch withCL_INVALID_OPERATION.Level Zero now adds
ZE_KERNEL_INDIRECT_ACCESS_FLAG_SHAREDunconditionally, mirroring the OpenCL backend, and logs the call once per kernel. The OpenCL memory manager records managed and unified allocations as host allocations.TestFixManagedIndirectAccesscovers both indirect paths and checks the module was actually flagged by the IGBA detector.Fixes #1503
Fixes #1504