[DeviceSanitizer] Check out-of-bounds on sycl::local_accessor #1532
Conversation
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
github-actions
bot
added
loader
|
Hi @oneapi-src/unified-runtime-maintain, please review. |
AllanZyne
commented
Apr 23, 2024
Comment on lines
+82
to
+101
| struct KernelInfo { | ||
| ur_kernel_handle_t Handle; | ||
|
|
||
| ur_shared_mutex Mutex; | ||
| // Need preserve the order of local arguments | ||
| std::map<uint32_t, LocalArgsInfo> LocalArgs; | ||
|
|
||
| explicit KernelInfo(ur_kernel_handle_t Kernel) : Handle(Kernel) { | ||
| [[maybe_unused]] auto Result = | ||
| context.urDdiTable.Kernel.pfnRetain(Kernel); | ||
| assert(Result == UR_RESULT_SUCCESS); | ||
| } | ||
|
|
||
| ~KernelInfo() { | ||
| [[maybe_unused]] auto Result = | ||
| context.urDdiTable.Kernel.pfnRelease(Handle); | ||
| assert(Result == UR_RESULT_SUCCESS); | ||
| } | ||
| }; | ||
|
|
There was a problem hiding this comment.
Just a note: KernelInfo related codes are duplicated with sycl::buffer PR #1533 (under review).
|
Kindly ping @oneapi-src/unified-runtime-maintain. Thanks very much! |
pbalcer
approved these changes
Apr 30, 2024
wenju-he
approved these changes
Apr 30, 2024
|
|
||
| struct LaunchInfo { | ||
| uintptr_t PrivateShadowOffset = | ||
| 0; // don't move this field, we use it in AddressSanitizerPass |
There was a problem hiding this comment.
to help me better understand the above comment, can we move other fields?
There was a problem hiding this comment.
Yes, you can reorder other fields except PrivateShadowOffset, because we assume PrivateShadowOffset is the first field in ASanPass.
steffenlarsen
pushed a commit
to intel/llvm
that referenced
this pull request
May 16, 2024
…ccessor (#13503) UR: oneapi-src/unified-runtime#1532 To check sycl::local_accessor(aka, dynamic local memory), we need to extend a new argument in spir kernel, this is because: - ASan needs to know some size information of local buffer, like its size and size with redzone, so that it can poison its shadow memory - By using this new argument, we can also pass some per-launch information (that is, it is different in each launch of kernel). One obvious example is SanitizerReport, which saves the error message, so that we can store and print multiple error reports for one kernel with different arguments. Another example is the shadow memory of local memory, this should be different per-launch as well, since one kernel can be launched multiple times and executed in parallel. I named this argument as "__asan_launch", which is a pointer pointed to "LaunchInfo" structure and allocated it in shared USM. To make this pointer can be used in spir_func w/o extending their argument, I created a global external local memory (external, so that it can be shared with other translation units, and its instance is defined in libdevice), and save the "__asan_launch" into this local memory immediately at the entry of kernel. UR can't check the name of kernel arguments, so it can't know if the kernel has "__asan_launch". So I assume the "__asan_launch" is always there, and added a check to prevent DAE pass from removing it. --------- Co-authored-by: Maosu Zhao <maosu.zhao@intel.com> Co-authored-by: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
LLVM: intel/llvm#13503