[PATCH] drm/amdkfd: don't widen the validate GPU set past one page owner

[PATCH] drm/amdkfd: don't widen the validate GPU set past one page owner

[Andrew Martin]

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, 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 <andrew.martin@amd.com>
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

RE: [PATCH] drm/amdkfd: don't widen the validate GPU set past one page owner

[Martin, Andrew]

AMD General

Greetings, sending peace

Just a Ping to have a look at this.

> -----Original Message-----
> From: Martin, Andrew <Andrew.Martin@amd.com>
> Sent: Tuesday, September 15, 2026 10:59 AM
> To: amd-gfx@lists.freedesktop.org
> Cc: Yang, Philip <Philip.Yang@amd.com>; Martin, Andrew
> <Andrew.Martin@amd.com>
> 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, 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 <andrew.martin@amd.com>
> 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

Re: [PATCH] drm/amdkfd: don't widen the validate GPU set past one page owner

[Alex Deucher]

On Thu, Sep 17, 2026 at 2:26 PM Martin, Andrew <Andrew.Martin@amd.com> wrote:
>
> AMD General
>
> Greetings, sending peace
>
> Just a Ping to have a look at this.

Looks reasonable to me.
Acked-by: Alex Deucher <alexander.deucher@amd.com>

>
> > -----Original Message-----
> > From: Martin, Andrew <Andrew.Martin@amd.com>
> > Sent: Tuesday, September 15, 2026 10:59 AM
> > To: amd-gfx@lists.freedesktop.org
> > Cc: Yang, Philip <Philip.Yang@amd.com>; Martin, Andrew
> > <Andrew.Martin@amd.com>
> > 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, 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 <andrew.martin@amd.com>
> > 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
>

RE: [PATCH] drm/amdkfd: don't widen the validate GPU set past one page owner

[Martin, Andrew]

AMD General

Thanks Alex!

> -----Original Message-----
> From: Alex Deucher <alexdeucher@gmail.com>
> Sent: Monday, September 21, 2026 5:39 PM
> To: Martin, Andrew <Andrew.Martin@amd.com>
> Cc: amd-gfx@lists.freedesktop.org; Yang, Philip <Philip.Yang@amd.com>
> 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 <Andrew.Martin@amd.com>
> wrote:
> >
> > AMD General
> >
> > Greetings, sending peace
> >
> > Just a Ping to have a look at this.
>
> Looks reasonable to me.
> Acked-by: Alex Deucher <alexander.deucher@amd.com>
>
> >
> > > -----Original Message-----
> > > From: Martin, Andrew <Andrew.Martin@amd.com>
> > > Sent: Tuesday, September 15, 2026 10:59 AM
> > > To: amd-gfx@lists.freedesktop.org
> > > Cc: Yang, Philip <Philip.Yang@amd.com>; Martin, Andrew
> > > <Andrew.Martin@amd.com>
> > > 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 <andrew.martin@amd.com>
> > > 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
> >

Re: [PATCH] drm/amdkfd: don't widen the validate GPU set past one page owner

[Philip Yang]

[-- Attachment #1: Type: text/plain, Size: 5834 bytes --]



On 2026-09-15 10:58, Andrew Martin wrote:
> 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, 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")
The fix tag is correct, suggest to change patch title and commit message

drm/amdkfd: Fix always mapped range mapping to all ACCESS GPUs

     CWSR ranges use the always mapped flag and only need a mapping on the
     GPU running the queue. However, svm_range_validate_and_map now updates
     the mapping on every GPU with the ACCESS attribute, which is all
     supported GPUs when XNACK is on. After the range is prefetched to VRAM,
     those GPUs may span different XGMI hives, so the page owner mismatches
     and is set to NULL. hmm_range_fault then tries to migrate the VRAM
     and is set to NULL. hmm_range_fault then tries to migrate the VRAM
     pages back to system memory, but svm_migrate_to_ram skips it because
     migrating pages back to system memory is not allowed while the GPU
     mapping is being updated, so the CPU page fault is never resolved.

     Only map to the ACCESS GPUs if the range is not mapped to any GPU and
     is not in VRAM.

> Signed-off-by: Andrew Martin<andrew.martin@amd.com>
> 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);
         /*
          * If prange with always mapped flag is not mapped or migrated
          * yet, update mapping on GPUs with ACCESS attribute. Don't add
          * them otherwise, ACCESS GPUs may be on different XGMI hives,
          * that would force hmm_range_fault to migrate VRAM back to ram.
          */
         if (bitmap_empty(ctx->bitmap, MAX_GPU_INSTANCE) &&
             prange->flags & KFD_IOCTL_SVM_FLAG_GPU_ALWAYS_MAPPED)
             bitmap_copy(ctx->bitmap, prange->bitmap_access,
                     MAX_GPU_INSTANCE);

Regards,
Philip
>   	} else {

[-- Attachment #2: Type: text/html, Size: 7211 bytes --]