forked from llvm/llvm-project
-
Notifications
You must be signed in to change notification settings - Fork 0
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[Offload] Improve error reporting on memory faults (llvm#104254)
Since we can already track allocations, we can diagnose memory faults to some degree. If the fault happens in a prior allocation (use after free) or "close but outside" one, we can provide that information to the user. Note that the fault address might be page aligned, and not all accesses trigger a fault, especially for allocations that are backed by a MemoryManager. Still, if people disable the MemoryManager or the allocation is big enough, we can sometimes provide valueable feedback.
- Loading branch information
Showing
11 changed files
with
256 additions
and
18 deletions.
There are no files selected for viewing
This file contains 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 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 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 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 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 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 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 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,40 @@ | ||
| // clang-format off | ||
| // RUN: %libomptarget-compileopt-generic | ||
| // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,NTRCE | ||
| // RUN: %libomptarget-compileopt-generic | ||
| // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,TRACE | ||
| // clang-format on | ||
|
|
||
| // UNSUPPORTED: aarch64-unknown-linux-gnu | ||
| // UNSUPPORTED: aarch64-unknown-linux-gnu-LTO | ||
| // UNSUPPORTED: x86_64-pc-linux-gnu | ||
| // UNSUPPORTED: x86_64-pc-linux-gnu-LTO | ||
| // UNSUPPORTED: s390x-ibm-linux-gnu | ||
| // UNSUPPORTED: s390x-ibm-linux-gnu-LTO | ||
|
|
||
| #include <omp.h> | ||
|
|
||
| void *llvm_omp_target_alloc_host(size_t Size, int DeviceNum); | ||
| void llvm_omp_target_free_host(void *Ptr, int DeviceNum); | ||
|
|
||
| int main() { | ||
| int N = (1 << 30); | ||
| char *A = (char *)llvm_omp_target_alloc_host(N, omp_get_default_device()); | ||
| char *P; | ||
| #pragma omp target map(from : P) | ||
| { | ||
| P = &A[0]; | ||
| *P = 3; | ||
| } | ||
| // clang-format off | ||
| // CHECK: OFFLOAD ERROR: Memory access fault by GPU {{.*}} (agent 0x{{.*}}) at virtual address [[PTR:0x[0-9a-z]*]]. Reasons: {{.*}} | ||
| // NTRCE: Use 'OFFLOAD_TRACK_ALLOCATION_TRACES=true' to track device allocations | ||
| // TRACE: Device pointer [[PTR]] does not point into any (current or prior) host-issued allocation. | ||
| // TRACE: Closest host-issued allocation (distance 4096 bytes; might be by page): | ||
| // TRACE: Last allocation of size 1073741824 | ||
| // clang-format on | ||
| #pragma omp target | ||
| { P[-4] = 5; } | ||
|
|
||
| llvm_omp_target_free_host(A, omp_get_default_device()); | ||
| } |
This file contains 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,26 @@ | ||
| // clang-format off | ||
| // RUN: %libomptarget-compileopt-generic | ||
| // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK | ||
| // clang-format on | ||
|
|
||
| // UNSUPPORTED: aarch64-unknown-linux-gnu | ||
| // UNSUPPORTED: aarch64-unknown-linux-gnu-LTO | ||
| // UNSUPPORTED: x86_64-pc-linux-gnu | ||
| // UNSUPPORTED: x86_64-pc-linux-gnu-LTO | ||
| // UNSUPPORTED: s390x-ibm-linux-gnu | ||
| // UNSUPPORTED: s390x-ibm-linux-gnu-LTO | ||
|
|
||
| #include <omp.h> | ||
|
|
||
| int main() { | ||
| int N = (1 << 30); | ||
| char *A = (char *)malloc(N); | ||
| #pragma omp target map(A[ : N]) | ||
| { A[N] = 3; } | ||
| // clang-format off | ||
| // CHECK: OFFLOAD ERROR: Memory access fault by GPU {{.*}} (agent 0x{{.*}}) at virtual address [[PTR:0x[0-9a-z]*]]. Reasons: {{.*}} | ||
| // CHECK: Device pointer [[PTR]] does not point into any (current or prior) host-issued allocation. | ||
| // CHECK: Closest host-issued allocation (distance 1 byte; might be by page): | ||
| // CHECK: Last allocation of size 1073741824 | ||
| // clang-format on | ||
| } |
This file contains 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,39 @@ | ||
| // clang-format off | ||
| // RUN: %libomptarget-compileopt-generic | ||
| // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,NTRCE | ||
| // RUN: %libomptarget-compileopt-generic | ||
| // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,TRACE | ||
| // clang-format on | ||
|
|
||
| // UNSUPPORTED: aarch64-unknown-linux-gnu | ||
| // UNSUPPORTED: aarch64-unknown-linux-gnu-LTO | ||
| // UNSUPPORTED: x86_64-pc-linux-gnu | ||
| // UNSUPPORTED: x86_64-pc-linux-gnu-LTO | ||
| // UNSUPPORTED: s390x-ibm-linux-gnu | ||
| // UNSUPPORTED: s390x-ibm-linux-gnu-LTO | ||
|
|
||
| #include <omp.h> | ||
|
|
||
| void *llvm_omp_target_alloc_host(size_t Size, int DeviceNum); | ||
| void llvm_omp_target_free_host(void *Ptr, int DeviceNum); | ||
|
|
||
| int main() { | ||
| int N = (1 << 30); | ||
| char *A = (char *)llvm_omp_target_alloc_host(N, omp_get_default_device()); | ||
| char *P; | ||
| #pragma omp target map(from : P) | ||
| { | ||
| P = &A[N / 2]; | ||
| *P = 3; | ||
| } | ||
| llvm_omp_target_free_host(A, omp_get_default_device()); | ||
| // clang-format off | ||
| // CHECK: OFFLOAD ERROR: Memory access fault by GPU {{.*}} (agent 0x{{.*}}) at virtual address [[PTR:0x[0-9a-z]*]]. Reasons: {{.*}} | ||
| // NTRCE: Use 'OFFLOAD_TRACK_ALLOCATION_TRACES=true' to track device allocations | ||
| // TRACE: Device pointer [[PTR]] points into prior host-issued allocation: | ||
| // TRACE: Last deallocation: | ||
| // TRACE: Last allocation of size 1073741824 | ||
| // clang-format on | ||
| #pragma omp target | ||
| { *P = 5; } | ||
| } |
This file contains 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,32 @@ | ||
| // clang-format off | ||
| // RUN: %libomptarget-compileopt-generic | ||
| // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK | ||
| // clang-format on | ||
|
|
||
| // UNSUPPORTED: aarch64-unknown-linux-gnu | ||
| // UNSUPPORTED: aarch64-unknown-linux-gnu-LTO | ||
| // UNSUPPORTED: x86_64-pc-linux-gnu | ||
| // UNSUPPORTED: x86_64-pc-linux-gnu-LTO | ||
| // UNSUPPORTED: s390x-ibm-linux-gnu | ||
| // UNSUPPORTED: s390x-ibm-linux-gnu-LTO | ||
|
|
||
| #include <omp.h> | ||
|
|
||
| int main() { | ||
| int N = (1 << 30); | ||
| char *A = (char *)malloc(N); | ||
| char *P; | ||
| #pragma omp target map(A[ : N]) map(from : P) | ||
| { | ||
| P = &A[N / 2]; | ||
| *P = 3; | ||
| } | ||
| // clang-format off | ||
| // CHECK: OFFLOAD ERROR: Memory access fault by GPU {{.*}} (agent 0x{{.*}}) at virtual address [[PTR:0x[0-9a-z]*]]. Reasons: {{.*}} | ||
| // CHECK: Device pointer [[PTR]] points into prior host-issued allocation: | ||
| // CHECK: Last deallocation: | ||
| // CHECK: Last allocation of size 1073741824 | ||
| // clang-format on | ||
| #pragma omp target | ||
| { *P = 5; } | ||
| } |