Add support for enabling/querying local memory exceptions. #105
There are no files selected for viewing
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
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -208,11 +208,15 @@ struct amd_dbgapi_inferior_info | |
| { | ||
| explicit amd_dbgapi_inferior_info | ||
| (inferior *inf, bool precise_memory_requested = false, | ||
| bool precise_alu_exceptions_requested = false, | ||
| bool local_memory_out_of_addr_range_exception_requested = false) | ||
|
|
||
| : inf (inf) | ||
| { | ||
| precise_memory.requested = precise_memory_requested; | ||
| precise_alu_exceptions.requested = precise_alu_exceptions_requested; | ||
| local_memory_out_of_addr_range_exception.requested = | ||
| local_memory_out_of_addr_range_exception_requested; | ||
| } | ||
|
|
||
| /* Backlink to inferior. */ | ||
|
|
@@ -252,6 +256,10 @@ struct amd_dbgapi_inferior_info | |
| enablement. */ | ||
| struct dbgapi_feature_request precise_alu_exceptions; | ||
|
|
||
| /* Track the status of local memory out-of-address-range exception reporting | ||
| request and enablement. */ | ||
| struct dbgapi_feature_request local_memory_out_of_addr_range_exception; | ||
|
|
||
| gdb::unordered_map<decltype (amd_dbgapi_breakpoint_id_t::handle), | ||
| struct breakpoint *> | ||
| breakpoint_map; | ||
|
|
@@ -2495,6 +2503,33 @@ set_process_alu_exceptions_precision (amd_dbgapi_inferior_info &info) | |
| get_status_string (status)); | ||
| } | ||
|
|
||
| /* Set the process' local memory out-of-address-range exception reporting mode. | ||
|
|
||
| Warn if the requested mode is not supported by at least one of the agents | ||
| in the process. */ | ||
|
|
||
| static void | ||
| set_process_local_memory_out_of_addr_range_exception (amd_dbgapi_inferior_info &info) | ||
| { | ||
| auto mode = (info.local_memory_out_of_addr_range_exception.requested | ||
| ? AMD_DBGAPI_GROUP_SEGMENT_EXCEPTIONS_OUT_OF_ADDR_RANGE | ||
| : AMD_DBGAPI_GROUP_SEGMENT_EXCEPTIONS_NONE); | ||
|
|
||
| amd_dbgapi_status_t status | ||
| = amd_dbgapi_set_group_segment_out_of_addr_range_exception | ||
| (info.process_id, mode); | ||
|
|
||
| if (status == AMD_DBGAPI_STATUS_SUCCESS) | ||
| info.local_memory_out_of_addr_range_exception.enabled = | ||
| info.local_memory_out_of_addr_range_exception.requested; | ||
| else if (status == AMD_DBGAPI_STATUS_ERROR_NOT_SUPPORTED) | ||
| warning (_("AMDGPU local memory out-of-address-range exception reporting could not " | ||
|
Contributor
There was a problem hiding this comment. could not be enabled -> is not supported. |
||
| "be enabled.")); | ||
| else if (status != AMD_DBGAPI_STATUS_SUCCESS) | ||
| error (_("amd_dbgapi_set_local_memory_out_of_addr_range_exception failed (%s)"), | ||
| get_status_string (status)); | ||
| } | ||
|
|
||
| /* Handle extra initialisation after we have attached to a AMDGPU corefile. */ | ||
|
|
||
| static void | ||
|
|
@@ -2617,6 +2652,7 @@ attach_amd_dbgapi (inferior *inf) | |
|
|
||
| set_process_memory_precision (info); | ||
| set_process_alu_exceptions_precision (info); | ||
| set_process_local_memory_out_of_addr_range_exception (info); | ||
|
|
||
| /* If GDB is attaching to a process that has the runtime loaded, there will | ||
| already be a "runtime loaded" event available. Consume it and push the | ||
|
|
@@ -2669,7 +2705,8 @@ detach_amd_dbgapi (inferior *inf) | |
| /* Reset the amd_dbgapi_inferior_info, except for precise_memory_mode and | ||
| precise_alu_exceptions. */ | ||
| info = amd_dbgapi_inferior_info (inf, info.precise_memory.requested, | ||
| info.precise_alu_exceptions.requested, | ||
| info.local_memory_out_of_addr_range_exception.requested); | ||
|
|
||
| maybe_reset_amd_dbgapi (); | ||
| } | ||
|
|
@@ -3074,12 +3111,14 @@ amd_dbgapi_target_inferior_cloned (inferior *original_inferior, | |
| = get_amd_dbgapi_inferior_info (new_inferior); | ||
|
|
||
| /* At this point, the process is not started. Therefore it is sufficient to | ||
| copy the precise memory and local memory request, it will be applied when | ||
| the process starts. */ | ||
|
Collaborator
There was a problem hiding this comment. We now have two things, so some words must be adjusted for plural: request => requests |
||
| gdb_assert (new_info.process_id == AMD_DBGAPI_PROCESS_NONE); | ||
| new_info.precise_memory.requested = orig_info.precise_memory.requested; | ||
| new_info.precise_alu_exceptions.requested | ||
| = orig_info.precise_alu_exceptions.requested; | ||
| new_info.local_memory_out_of_addr_range_exception.requested | ||
| = orig_info.local_memory_out_of_addr_range_exception.requested; | ||
| } | ||
|
|
||
| /* inferior_execd observer. */ | ||
|
|
@@ -3092,14 +3131,17 @@ amd_dbgapi_inferior_execd (inferior *exec_inf, inferior *follow_inf) | |
| the new process image. */ | ||
| detach_amd_dbgapi (exec_inf); | ||
|
|
||
| /* If using "follow-exec-mode new", carry over the precise-memory and local-memory | ||
| settings to the new inferior (otherwise, FOLLOW_INF and ORIG_INF point to | ||
| the same inferior, so this is a no-op). */ | ||
| get_amd_dbgapi_inferior_info (follow_inf).precise_memory.requested | ||
| = get_amd_dbgapi_inferior_info (exec_inf).precise_memory.requested; | ||
| get_amd_dbgapi_inferior_info (follow_inf).precise_alu_exceptions.requested | ||
| = get_amd_dbgapi_inferior_info (exec_inf) | ||
| .precise_alu_exceptions.requested; | ||
| get_amd_dbgapi_inferior_info (follow_inf).local_memory_out_of_addr_range_exception | ||
| .requested = get_amd_dbgapi_inferior_info (exec_inf) | ||
| .local_memory_out_of_addr_range_exception.requested; | ||
|
|
||
| attach_amd_dbgapi (follow_inf); | ||
| } | ||
|
|
@@ -3112,7 +3154,7 @@ amd_dbgapi_inferior_forked (inferior *parent_inf, inferior *child_inf, | |
| { | ||
| if (child_inf != nullptr) | ||
| { | ||
| /* Copy precise-memory and local-memory requested values from parent to child. */ | ||
| const amd_dbgapi_inferior_info &parent_info | ||
| = get_amd_dbgapi_inferior_info (parent_inf); | ||
| amd_dbgapi_inferior_info &child_info | ||
|
|
@@ -3121,6 +3163,8 @@ amd_dbgapi_inferior_forked (inferior *parent_inf, inferior *child_inf, | |
| = parent_info.precise_memory.requested; | ||
| child_info.precise_alu_exceptions.requested | ||
| = parent_info.precise_alu_exceptions.requested; | ||
| child_info.local_memory_out_of_addr_range_exception.requested | ||
| = parent_info.local_memory_out_of_addr_range_exception.requested; | ||
|
|
||
| if (fork_kind != TARGET_WAITKIND_VFORKED) | ||
| { | ||
|
|
@@ -3350,6 +3394,13 @@ may not be accurate. See \"show amdgpu precise-alu-exceptions\".\n")); | |
| gdb_printf (_("\ | ||
| Warning: precise memory violation signal reporting is not enabled, reported\n\ | ||
| location may not be accurate. See \"show amdgpu precise-memory\".\n")); | ||
| } | ||
| if (!info.local_memory_out_of_addr_range_exception.enabled) | ||
| { | ||
| gdb_printf (_("\ | ||
| Warning: local memory out-of-address-range exception reporting is not enabled.\n\ | ||
| See \"show amdgpu lds-addr-range-exception\".\n")); | ||
|
|
||
| } | ||
| return; | ||
| default: | ||
|
|
@@ -3693,6 +3744,58 @@ get_effective_precise_alu_exception_mode () | |
| return info.precise_alu_exceptions.enabled; | ||
| } | ||
|
|
||
| /* Callback for "show amdgpu lds-addr-range-exception". */ | ||
|
|
||
| static void | ||
| show_lds_addr_range_excp_mode (struct ui_file *file, int from_tty, | ||
| struct cmd_list_element *c, const char *value) | ||
|
Contributor
There was a problem hiding this comment. The alignment of the text is off. I'm going to suppress other GNU coding standard violations at this point. Please find and fix them. |
||
| { | ||
| amd_dbgapi_inferior_info &info | ||
| = get_amd_dbgapi_inferior_info (current_inferior ()); | ||
|
|
||
| gdb_printf (file, | ||
| _("AMDGPU local memory out-of-address-range exception reporting is %s " | ||
| "(currently %s).\n"), | ||
| info.local_memory_out_of_addr_range_exception.requested ? "on" : "off", | ||
| info.local_memory_out_of_addr_range_exception.enabled ? | ||
| "enabled" : "disabled"); | ||
| } | ||
|
|
||
| /* Callback for "set amdgpu lds-addr-range-exception". */ | ||
|
|
||
| static void | ||
| set_lds_addr_range_excp_mode (bool value) | ||
| { | ||
| amd_dbgapi_inferior_info &info | ||
| = get_amd_dbgapi_inferior_info (current_inferior ()); | ||
|
|
||
| info.local_memory_out_of_addr_range_exception.requested = value; | ||
|
|
||
| if (info.process_id != AMD_DBGAPI_PROCESS_NONE) | ||
| set_process_local_memory_out_of_addr_range_exception (info); | ||
| } | ||
|
|
||
| /* Get the local memory out-of-address-range exception reporting requested mode. */ | ||
|
|
||
| static bool | ||
| get_lds_addr_range_excp_mode () | ||
| { | ||
| amd_dbgapi_inferior_info &info | ||
| = get_amd_dbgapi_inferior_info (current_inferior ()); | ||
| return info.local_memory_out_of_addr_range_exception.requested; | ||
| } | ||
|
|
||
| /* Get the local memory out-of-address-range exception reporting effective mode. */ | ||
|
|
||
| static bool | ||
| get_effective_local_memory_out_of_addr_range_exception_mode () | ||
| { | ||
| amd_dbgapi_inferior_info &info | ||
| = get_amd_dbgapi_inferior_info (current_inferior ()); | ||
| return info.local_memory_out_of_addr_range_exception.enabled; | ||
| } | ||
|
|
||
|
|
||
| static const char * | ||
| get_dbgapi_library_file_path () | ||
| { | ||
|
|
@@ -5035,6 +5138,20 @@ running. If off (default), precise ALU exceptions reporting is disabled."), | |
| cmds.show->var->set_effective_value_getter<bool> | ||
| (get_effective_precise_alu_exception_mode); | ||
|
|
||
| cmds = add_setshow_boolean_cmd ("lds-addr-range-exception", no_class, | ||
| _("Set lds-addr-range-exception mode."), | ||
| _("Show lds-addr-range-exception mode."), _("\ | ||
| If on, local memory out-of-address-range exception reporting is enabled if/when\n\ | ||
|
Contributor
There was a problem hiding this comment. looks like you copy/pasted this from the code block above, nonetheless I think "if/when" is not adding any value. a mere "when" would suffice. |
||
| the inferior is running. If off (default), local memory out-of-address-range\n\ | ||
| exception reporting is disabled."), | ||
| set_lds_addr_range_excp_mode, | ||
| get_lds_addr_range_excp_mode, | ||
| show_lds_addr_range_excp_mode, | ||
| &set_amdgpu_list, &show_amdgpu_list); | ||
|
Collaborator
There was a problem hiding this comment. This setting / feature must be documented in the GDB manual. |
||
|
|
||
| cmds.show->var->set_effective_value_getter<bool> | ||
| (get_effective_local_memory_out_of_addr_range_exception_mode); | ||
|
|
||
| add_cmd ("version", no_set_class, show_dbgapi_version, | ||
| _("Show the ROCdbgapi library version and build information."), | ||
| &show_amdgpu_list); | ||
|
|
||
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
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
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
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,68 @@ | ||
| /* Copyright (C) 2022-2025 Advanced Micro Devices, Inc. All rights reserved. | ||
|
|
||
|
|
||
| This file is part of GDB. | ||
|
|
||
| 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> | ||
| #include <stdio.h> | ||
| #include <stdlib.h> | ||
| #include <cassert> | ||
| #include "rocm-test-utils.h" | ||
|
|
||
| __global__ void | ||
| kern (size_t dyn_alloc, size_t lds_size) | ||
| { | ||
| extern __shared__ char arr[]; | ||
|
|
||
| /* Access up to dyn_alloc should always work. */ | ||
| for (size_t i = 0; i < dyn_alloc; i++) | ||
| arr[i] = i; | ||
|
|
||
| __syncthreads (); | ||
|
|
||
| for (size_t i = 0; i < dyn_alloc; i++) | ||
| assert (arr[i] == i); | ||
|
|
||
| /* Use the last four bytes of the LDS */ | ||
| size_t idx = lds_size - 4; | ||
| if (threadIdx.x == 0) | ||
| { | ||
| assert (idx >= dyn_alloc); | ||
| arr[idx] = 8; | ||
| } | ||
| __syncthreads (); | ||
|
|
||
| /* This is expected to fail. | ||
| One could run this kernel once and expect to fail at the assert, then run | ||
| it again with the LDS reporting on, and check you receive the memviol. */ | ||
| if (threadIdx.x == 0) | ||
| assert (arr[idx] == 8); | ||
| __syncthreads (); | ||
| } | ||
|
|
||
| int | ||
| main (int argc, char* argv[]) | ||
| { | ||
| hipDeviceProp_t props; | ||
| int deviceId; | ||
| CHECK (hipGetDevice (&deviceId)); | ||
| CHECK (hipGetDeviceProperties(&props, deviceId)); | ||
|
|
||
| constexpr size_t shared_mem_alloc = 64; | ||
| size_t lds_size = props.sharedMemPerBlock; | ||
| kern<<<1, 128, shared_mem_alloc, 0>>> (shared_mem_alloc, lds_size); | ||
| CHECK (hipDeviceSynchronize ()); | ||
| return 0; | ||
| } | ||
Oops, something went wrong.
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.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
these new added lines are duplicate and should be removed.
The new entries should also go at the end of the "Added" list.