AMD General Thanks Alex!
> -----Original Message----- > From: Alex Deucher <[email protected]> > Sent: Monday, September 21, 2026 5:39 PM > To: Martin, Andrew <[email protected]> > Cc: [email protected]; Yang, Philip <[email protected]> > Subject: Re: [PATCH] drm/amdkfd: don't widen the validate GPU set past one > page owner > > Caution: This message originated from an External Source. Use proper caution > when opening attachments, clicking links, or responding. > > > On Thu, Sep 17, 2026 at 2:26 PM Martin, Andrew <[email protected]> > wrote: > > > > AMD General > > > > Greetings, sending peace > > > > Just a Ping to have a look at this. > > Looks reasonable to me. > Acked-by: Alex Deucher <[email protected]> > > > > > > -----Original Message----- > > > From: Martin, Andrew <[email protected]> > > > Sent: Tuesday, September 15, 2026 10:59 AM > > > To: [email protected] > > > Cc: Yang, Philip <[email protected]>; Martin, Andrew > > > <[email protected]> > > > Subject: [PATCH] drm/amdkfd: don't widen the validate GPU set past > > > one page owner > > > Importance: High > > > > > > Commit 6b2969831f7e ("drm/amdkfd: Unmap svm range from GPU set to > > > no-access") made svm_range_validate_and_map() OR bitmap_access into > > > ctx->bitmap unconditionally for a range carrying > > > KFD_IOCTL_SVM_FLAG_GPU_ALWAYS_MAPPED. Previously that only > happened > > > when > > > ctx->bitmap was still empty: > > > > > > if (bitmap_empty(ctx->bitmap, MAX_GPU_INSTANCE)) { > > > if (prange->mapped_to_gpu || > > > prange->flags & > > > KFD_IOCTL_SVM_FLAG_GPU_ALWAYS_MAPPED) > > > bitmap_copy(ctx->bitmap, prange->bitmap_access, > > > MAX_GPU_INSTANCE); > > > } > > > > > > The block just above already sets the bit for prange->actual_loc > > > when the range has been migrated to VRAM, so under the old code > > > ctx->bitmap was left holding that single GPU and the OR was skipped. > > > Without the guard it now also picks up every other GPU with the ACCESS > attribute. > > > > > > That matters because ctx->bitmap is not only the set of GPUs to map. > > > It also selects the HMM owner: owner is taken from the first set bit > > > and reset to NULL as soon as two set bits disagree. On a two GPU > > > system with ACCESS on both, owner becomes NULL, and > > > amdgpu_hmm_range_get_pages() then no longer recognises the process's > > > own device-private pages. HMM calls migrate_to_ram on each of them, > > > svm_migrate_to_ram() sees > > > p->svms.faulting_task == current and returns 0 without migrating, > > > p->and > > > hmm_range_fault() retries the same page forever. > > > > > > Which GPUs disagree depends on the interconnect. > > > SVM_ADEV_PGMAP_OWNER() is the XGMI hive when the device has one > and > > > the adev otherwise, so GPUs sharing a hive also share an owner and > > > never trip this. It takes two or more GPUs with ACCESS that are not > > > in a common hive, which is why this only shows up on multi-GPU > > > machines with no XGMI link between the cards. > > > > > > On a 2x Instinct MI350P (gfx950) this hangs > > > > > > HSA_XNACK=1 hipblaslt-test \ > > > > > > --gtest_filter=aux_handle_test.set_sm_count_target_default_is_zero > > > > > > The application never completes and burns 100% of a CPU in the kernel. > > > /proc/<pid>/stat shows utime frozen while stime climbs and minflt > > > advances at roughly 12 million faults per second, reaching billions; > > > majflt stays at 0 and the GPU sits at 0% busy. perf on the stuck > > > task is 24% svm_migrate_to_ram, 23% __handle_mm_fault, 5% > > > do_swap_page. The process cannot be killed and the machine has to be > rebooted. > > > > > > Restore the bitmap_empty() guard. The rest of that commit is kept: > > > the bitmap_mapped tracking, bitmap_needs_unmap, mapping_done and > the > > > unmap of ranges set to no-access are all unaffected, and > > > bitmap_mapped is still ORed in at the top of the branch, so a range > > > already mapped on a GPU still has its mapping updated there. > > > > > > The instrumentation, the hardware testing and this changelog were > > > done with the assistance of an AI coding agent; the resulting change > > > has been reviewed by hand. > > > > > > Fixes: 6b2969831f7e ("drm/amdkfd: Unmap svm range from GPU set to > > > no- > > > access") > > > Signed-off-by: Andrew Martin <[email protected]> > > > Assisted-by: Claude:claude-opus-5 > > > --- > > > drivers/gpu/drm/amd/amdkfd/kfd_svm.c | 9 +++++++-- > > > 1 file changed, 7 insertions(+), 2 deletions(-) > > > > > > diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_svm.c > > > b/drivers/gpu/drm/amd/amdkfd/kfd_svm.c > > > index c521ec07bf3c..75e25f075da3 100644 > > > --- a/drivers/gpu/drm/amd/amdkfd/kfd_svm.c > > > +++ b/drivers/gpu/drm/amd/amdkfd/kfd_svm.c > > > @@ -1758,9 +1758,14 @@ static int svm_range_validate_and_map(struct > > > mm_struct *mm, > > > > > > /* > > > * If prange with always mapped flag, update mapping > > > on GPUs with > > > - * ACCESS attribute > > > + * ACCESS attribute. Only widen the set when nothing else > > > has > > > + * selected a GPU: ctx->bitmap also picks the owner passed > > > to > > > + * amdgpu_hmm_range_get_pages() below, and spanning > > > GPUs with > > > + * different page owners forces that owner to NULL, > > > + which > > > makes > > > + * HMM fault on the process's own migrated pages. > > > */ > > > - if (prange->flags & > > > KFD_IOCTL_SVM_FLAG_GPU_ALWAYS_MAPPED) > > > + if (bitmap_empty(ctx->bitmap, MAX_GPU_INSTANCE) && > > > + prange->flags & > > > KFD_IOCTL_SVM_FLAG_GPU_ALWAYS_MAPPED) > > > bitmap_or(ctx->bitmap, ctx->bitmap, prange- > > > >bitmap_access, > > > MAX_GPU_INSTANCE); > > > } else { > > > -- > > > 2.43.0 > >
