Introduce CUDA Device Assertions Infrastructure#84609
Closed
r-barnes wants to merge 1 commit intopytorch:masterfrom
Closed
Introduce CUDA Device Assertions Infrastructure#84609r-barnes wants to merge 1 commit intopytorch:masterfrom
r-barnes wants to merge 1 commit intopytorch:masterfrom
Conversation
Contributor
🔗 Helpful links
❌ 6 New Failures, 2 Base FailuresAs of commit 0c1508614c (more details on the Dr. CI page): Expand to see more
🕵️ 6 new failures recognized by patternsThe following CI failures do not appear to be due to upstream breakages
|
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
0c15086 to
73f43b1
Compare
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
🔗 Helpful Links🧪 See artifacts and rendered test results at hud.pytorch.org/pr/84609
Note: Links to docs will display an error until the docs builds have been completed. ✅ No FailuresAs of commit 19934bd: This comment was automatically generated by Dr. CI and updates every 15 minutes. |
Collaborator
|
this is pretty cool! |
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
73f43b1 to
57d647a
Compare
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
57d647a to
2afacff
Compare
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
2afacff to
466d1cd
Compare
466d1cd to
c217d64
Compare
c217d64 to
b9763d7
Compare
b9763d7 to
be5b129
Compare
be5b129 to
e0b4ee2
Compare
e0b4ee2 to
19642d3
Compare
Contributor
|
@r-barnes has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
1 similar comment
Contributor
|
@r-barnes has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
malfet
reviewed
Nov 8, 2022
c10/cuda/CUDADeviceAssertion.h
Outdated
|
|
||
| #ifdef TORCH_USE_CUDA_DSA | ||
| // Copy string from `src` to `dst` | ||
| static __device__ void dstrcpy(char* const dst, const char* const src) { |
Contributor
There was a problem hiding this comment.
Not sure if it supported across all cuda versions, but isn't it just strncpy? (and if it is not, perhaps it should be called strncpy
Also, not sure if static keyword has any effect on __device__ only functions (as they do not have global visibility to beging with)
Contributor
Author
There was a problem hiding this comment.
Leaving the static since my recollection is someone else requested it earlier in review and it seems otherwise benign.
I don't think strncpy works in CUDA and am not sure if we can adequately test for that, so I'd prefer to keep this for now.
This function's signature differs from strncpy, so I'd prefer to keep the name as-is unless you have other suggestions.
tl;dr I haven't made a change here.
c10/cuda/CUDAException.cpp
Outdated
| C10_CUDA_DSA_ASSERTION_COUNT) { | ||
| oss << "But at least " << assertion_data_for_device.assertion_count | ||
| << " assertion failures occurred on the device" << std::endl; | ||
| oss << "Adjust `C10_CUDA_DSA_ASSERTION_COUNT` if you need more assertion failure info" |
Contributor
There was a problem hiding this comment.
Though kinds of messages often get out of date (as one edits the environment variable but ignores the rest
Contributor
Author
There was a problem hiding this comment.
Do you have a proposed solution?
malfet
reviewed
Nov 8, 2022
Contributor
|
@r-barnes has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
3 similar comments
Contributor
|
@r-barnes has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
Contributor
|
@r-barnes has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
Contributor
|
@r-barnes has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
malfet
approved these changes
Nov 11, 2022
Contributor
|
@r-barnes has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
Contributor
|
@r-barnes has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
2 similar comments
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
Summary: This diff introduces a set of changes that makes it possible for the host to get assertions from CUDA devices. This includes the introduction of **`CUDA_KERNEL_ASSERT2`** A preprocessor macro to be used within a CUDA kernel that, upon an assertion failure, writes the assertion message, file, line number, and possibly other information to UVM (Managed memory). Once this is done, the original assertion is triggered, which places the GPU in a Bad State requiring recovery. In my tests, data written to UVM appears there before the GPU reaches the Bad State and is still accessible from the host after the GPU is in this state. Messages are written to a multi-message buffer which can, in theory, hold many assertion failures. I've done this as a precaution in case there are several, but I don't actually know whether that is possible and a simpler design which holds only a single message may well be all that is necessary. **`TORCH_DSA_KERNEL_ARGS`** This preprocess macro is added as an _argument_ to a kernel function's signature. It expands to supply the standardized names of all the arguments needed by `C10_CUDA_COMMUNICATING_KERNEL_ASSERTION` to handle device-side assertions. This includes, eg, the name of the pointer to the UVM memory the assertion would be written to. This macro abstracts the arguments so there is a single point of change if the system needs to be modified. **`c10::cuda::get_global_cuda_kernel_launch_registry()`** This host-side function returns a singleton object that manages the host's part of the device-side assertions. Upon allocation, the singleton allocates sufficient UVM (Managed) memory to hold information about several device-side assertion failures. The singleton also provides methods for getting the current traceback (used to identify when a kernel was launched). To avoid consuming all the host's memory the singleton stores launches in a circular buffer; a unique "generation number" is used to ensure that kernel launch failures map to their actual launch points (in the case that the circular buffer wraps before the failure is detected). **`TORCH_DSA_KERNEL_LAUNCH`** This host-side preprocessor macro replaces the standard ``` kernel_name<<<blocks, threads, shmem, stream>>>(args) ``` invocation with ``` TORCH_DSA_KERNEL_LAUNCH(blocks, threads, shmem, stream, args); ``` Internally, it fetches the UVM (Managed) pointer and generation number from the singleton and append these to the standard argument list. It also checks to ensure the kernel launches correctly. This abstraction on kernel launches can be modified to provide additional safety/logging. **`c10::cuda::c10_retrieve_device_side_assertion_info`** This host-side function checks, when called, that no kernel assertions have occurred. If one has. It then raises an exception with: 1. Information (file, line number) of what kernel was launched. 2. Information (file, line number, message) about the device-side assertion 3. Information (file, line number) about where the failure was detected. **Checking for device-side assertions** Device-side assertions are most likely to be noticed by the host when a CUDA API call such as `cudaDeviceSynchronize` is made and fails with a `cudaError_t` indicating > CUDA error: device-side assert triggered CUDA kernel errors Therefore, we rewrite `C10_CUDA_CHECK()` to include a call to `c10_retrieve_device_side_assertion_info()`. To make the code cleaner, most of the logic of `C10_CUDA_CHECK()` is now contained within a new function `c10_cuda_check_implementation()` to which `C10_CUDA_CHECK` passes the preprocessor information about filenames, function names, and line numbers. (In C++20 we can use `std::source_location` to eliminate macros entirely!) # Notes on special cases * Multiple assertions from the same block are recorded * Multiple assertions from different blocks are recorded * Launching kernels from many threads on many streams seems to be handled correctly * If two process are using the same GPU and one of the processes fails with a device-side assertion the other process continues without issue * X Multiple assertions from separate kernels on different streams seem to be recorded, but we can't reproduce the test condition * X Multiple assertions from separate devices should be all be shown upon exit, but we've been unable to generate a test that produces this condition Pull Request resolved: pytorch#84609 Reviewed By: ezyang Differential Revision: D37621532 Pulled By: r-barnes fbshipit-source-id: efdfc57d6a1fa6dadfe30e693157e4cc040c7191
Contributor
|
This pull request was exported from Phabricator. Differential Revision: D37621532 |
Contributor
|
@r-barnes has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
Contributor
Author
|
@pytorchbot merge -f "Internal changes is incorrect." |
Collaborator
Merge startedYour change will be merged immediately since you used the force (-f) flag, bypassing any CI checks (ETA: 1-5 minutes). Learn more about merging in the wiki. Questions? Feedback? Please reach out to the PyTorch DevX Team |
Collaborator
Merge failedReason: This PR has internal changes and must be landed via Phabricator Details for Dev Infra teamRaised by workflow job |
Contributor
Author
|
@pytorchbot merge -f "No internal changes" |
Collaborator
Merge startedYour change will be merged immediately since you used the force (-f) flag, bypassing any CI checks (ETA: 1-5 minutes). Learn more about merging in the wiki. Questions? Feedback? Please reach out to the PyTorch DevX Team |
Collaborator
Merge failedReason: This PR has internal changes and must be landed via Phabricator Details for Dev Infra teamRaised by workflow job |
Contributor
Author
|
@pytorchbot merge |
Collaborator
Merge startedYour change will be merged once all checks pass (ETA 0-4 Hours). Learn more about merging in the wiki. Questions? Feedback? Please reach out to the PyTorch DevX Team |
Closed
Collaborator
|
Why is this feature behind a compile-time flag The release not enabling this flag by default means that most users are not benefiting from this great feature. |
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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
11 participants
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.
Summary:
This diff introduces a set of changes that makes it possible for the host to get assertions from CUDA devices. This includes the introduction of
CUDA_KERNEL_ASSERT2A preprocessor macro to be used within a CUDA kernel that, upon an assertion failure, writes the assertion message, file, line number, and possibly other information to UVM (Managed memory). Once this is done, the original assertion is triggered, which places the GPU in a Bad State requiring recovery. In my tests, data written to UVM appears there before the GPU reaches the Bad State and is still accessible from the host after the GPU is in this state.
Messages are written to a multi-message buffer which can, in theory, hold many assertion failures. I've done this as a precaution in case there are several, but I don't actually know whether that is possible and a simpler design which holds only a single message may well be all that is necessary.
TORCH_DSA_KERNEL_ARGSThis preprocess macro is added as an argument to a kernel function's signature. It expands to supply the standardized names of all the arguments needed by
C10_CUDA_COMMUNICATING_KERNEL_ASSERTIONto handle device-side assertions. This includes, eg, the name of the pointer to the UVM memory the assertion would be written to. This macro abstracts the arguments so there is a single point of change if the system needs to be modified.c10::cuda::get_global_cuda_kernel_launch_registry()This host-side function returns a singleton object that manages the host's part of the device-side assertions. Upon allocation, the singleton allocates sufficient UVM (Managed) memory to hold information about several device-side assertion failures. The singleton also provides methods for getting the current traceback (used to identify when a kernel was launched). To avoid consuming all the host's memory the singleton stores launches in a circular buffer; a unique "generation number" is used to ensure that kernel launch failures map to their actual launch points (in the case that the circular buffer wraps before the failure is detected).
TORCH_DSA_KERNEL_LAUNCHThis host-side preprocessor macro replaces the standard
invocation with
Internally, it fetches the UVM (Managed) pointer and generation number from the singleton and append these to the standard argument list. It also checks to ensure the kernel launches correctly. This abstraction on kernel launches can be modified to provide additional safety/logging.
c10::cuda::c10_retrieve_device_side_assertion_infoThis host-side function checks, when called, that no kernel assertions have occurred. If one has. It then raises an exception with:
Checking for device-side assertions
Device-side assertions are most likely to be noticed by the host when a CUDA API call such as
cudaDeviceSynchronizeis made and fails with acudaError_tindicatingTherefore, we rewrite
C10_CUDA_CHECK()to include a call toc10_retrieve_device_side_assertion_info(). To make the code cleaner, most of the logic ofC10_CUDA_CHECK()is now contained within a new functionc10_cuda_check_implementation()to whichC10_CUDA_CHECKpasses the preprocessor information about filenames, function names, and line numbers. (In C++20 we can usestd::source_locationto eliminate macros entirely!)Notes on special cases
Differential Revision: D37621532