diff --git a/gdb/elfread.c b/gdb/elfread.c
index c91b9ab7c21..db024c06344 100644
--- a/gdb/elfread.c
+++ b/gdb/elfread.c
@@ -1046,7 +1046,7 @@ elf_gnu_ifunc_resolver_return_stop (code_breakpoint *b)
 
   b->type = bp_breakpoint;
   update_breakpoint_locations (b, current_program_space,
-			       find_function_start_sal (resolved_pc, true),
+			       find_function_start_sal (resolved_pc, nullptr, true),
 			       {});
 }
 
diff --git a/gdb/linespec.c b/gdb/linespec.c
index b7ddd166c8a..4560459ad3c 100644
--- a/gdb/linespec.c
+++ b/gdb/linespec.c
@@ -4059,6 +4059,7 @@ minsym_found (struct linespec_state *self, struct objfile *objfile,
 
   CORE_ADDR func_addr;
   bool is_function = msymbol_is_function (objfile, msymbol, &func_addr);
+  obj_section *section = msymbol->obj_section (objfile);
 
   if (is_function)
     {
@@ -4066,7 +4067,15 @@ minsym_found (struct linespec_state *self, struct objfile *objfile,
 
       if (msymbol->type () == mst_text_gnu_ifunc
 	  || msymbol->type () == mst_data_gnu_ifunc)
-	want_start_sal = gnu_ifunc_resolve_name (msym_name, &func_addr);
+	{
+	  want_start_sal = gnu_ifunc_resolve_name (msym_name, &func_addr);
+
+	  /* We have found a different pc by resolving the ifunc. The
+	     section from the minsym may not be the same as the ifunc
+	     implementation.  Do not trust it.  */
+	  if (want_start_sal)
+	    section = nullptr;
+	}
       else
 	want_start_sal = true;
     }
@@ -4074,7 +4083,7 @@ minsym_found (struct linespec_state *self, struct objfile *objfile,
   symtab_and_line sal;
 
   if (is_function && want_start_sal)
-    sal = find_function_start_sal (func_addr, self->funfirstline);
+    sal = find_function_start_sal (func_addr, section, self->funfirstline);
   else
     {
       sal.objfile = objfile;
@@ -4086,14 +4095,15 @@ minsym_found (struct linespec_state *self, struct objfile *objfile,
       else
 	sal.pc = msymbol->value_address (objfile);
       sal.pspace = current_program_space;
-    }
 
-  /* Don't use the section from the msymbol, the code above might have
-     adjusted FUNC_ADDR, in which case the msymbol's section might not be
-     the section containing FUNC_ADDR.  It might not even be in the same
-     objfile.  As the section is primarily to assist with overlay
-     debugging, it should reflect the SAL's pc value.  */
-  sal.section = find_pc_overlay (sal.pc);
+      /* The minsym does not correspond to an ifunc that could be
+	 resolved.  The section from the minsym may thus be trusted,
+	 and cannot be nullptr (since the minsym is from an objfile).
+	 Ensure all resulting sals have a non-null section when
+	 possible.  */
+      gdb_assert (section != nullptr);
+      sal.section = section;
+    }
 
   if (self->maybe_add_address (objfile->pspace (), sal.pc))
     add_sal_to_sals (self, result, sal, msymbol->natural_name (), false);
diff --git a/gdb/symtab.c b/gdb/symtab.c
index bd3d55eecb0..3b0687c0750 100644
--- a/gdb/symtab.c
+++ b/gdb/symtab.c
@@ -3560,10 +3560,10 @@ find_function_start_sal_1 (CORE_ADDR func_addr, obj_section *section,
 /* See symtab.h.  */
 
 symtab_and_line
-find_function_start_sal (CORE_ADDR func_addr, bool funfirstline)
+find_function_start_sal (CORE_ADDR func_addr, obj_section *section, bool funfirstline)
 {
   symtab_and_line sal
-    = find_function_start_sal_1 (func_addr, nullptr, funfirstline);
+    = find_function_start_sal_1 (func_addr, section, funfirstline);
 
   /* find_function_start_sal_1 does a linetable search, so it finds
      the symtab and linenumber, but not a symbol.  Fill in the
diff --git a/gdb/symtab.h b/gdb/symtab.h
index 985843f76b6..e47033efd01 100644
--- a/gdb/symtab.h
+++ b/gdb/symtab.h
@@ -2516,6 +2516,7 @@ extern symtab_and_line find_function_start_sal (symbol *sym, bool
 
 /* Same, but start with a function address instead of a symbol.  */
 extern symtab_and_line find_function_start_sal (CORE_ADDR func_addr,
+						obj_section *section,
 						bool funfirstline);
 
 extern void skip_prologue_sal (struct symtab_and_line *);
diff --git a/gdb/testsuite/gdb.rocm/break-kernel-no-debug-info.cpp b/gdb/testsuite/gdb.rocm/break-kernel-no-debug-info.cpp
new file mode 100644
index 00000000000..f46a57c256b
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/break-kernel-no-debug-info.cpp
@@ -0,0 +1,30 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2025 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#include <hip/hip_runtime.h>
+
+__global__ void
+kern ()
+{
+}
+
+int
+main ()
+{
+  kern<<<1, 1>>> ();
+  return hipDeviceSynchronize () != hipSuccess;
+}
diff --git a/gdb/testsuite/gdb.rocm/break-kernel-no-debug-info.exp b/gdb/testsuite/gdb.rocm/break-kernel-no-debug-info.exp
new file mode 100644
index 00000000000..df65b7c064e
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/break-kernel-no-debug-info.exp
@@ -0,0 +1,53 @@
+# Copyright 2025 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+# Test setting a breakpoint on a kernel symbol without debug info,
+# relying on minimal symbols from the ELF.
+
+# A bug occured when GDB did not find the appropriate architecture for
+# breakpoints on minimal symbols. This had the effect that the
+# breakpoint would not be hit on the GPU when no debugging infos are
+# available.
+
+load_lib rocm.exp
+
+standard_testfile .cpp
+
+require allow_hipcc_tests
+
+# Build for hip, explicitly without debug infos
+if {[build_executable "failed to prepare" $testfile $srcfile {hip nodebug}]} {
+    return
+}
+
+clean_restart
+
+with_rocm_gpu_lock {
+    gdb_test "file $::binfile" ".*No debugging symbols.*" "load file"
+
+    if {![runto_main]} {
+	return
+    }
+
+    gdb_test "with breakpoint pending on -- break kern" \
+	"Breakpoint $::decimal \\(kern\\) pending."
+
+    gdb_test "continue" \
+	"Thread $::decimal hit Breakpoint $::decimal, $::hex in kern.*"
+
+    gdb_test "continue" \
+	"Inferior 1 .* exited normally.*" \
+	"continue to end"
+}