Repository navigation
[SYCL][CUDA][HIP] Guidance for __launch_bounds__ in SYCL #8080
Description
Activity
__launch_bounds__()is seen in GPU programs written by researchers https://userweb.cs.txstate.edu/~burtscher/publications.html__launch_bounds__()is seen in GPU programs written by researchers https://userweb.cs.txstate.edu/~burtscher/publications.htmlThanks, Can you please point me the relevant publications among the above that makes use of launch_bounds() in SYCL
They make use of it in CUDA programs. Sorry for the confusion.
Reacted by Abhishek BagusettyUnfortunately this is not something we currently support in DPC++, we'll have to look into enabling this attribute for SYCL.
In the meantime, I haven't tested it carefully but it might be possible to tweak the register usage by using the
ptxasflag--maxrregcountinstead. You can pass it down throughclang++with the arguments:-Xcuda-ptxas --maxrregcount=<number>, for example:clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_60 cr.cpp -o cr \ -Xcuda-ptxas --maxrregcount=32Note that
ptxasis used during the linking stages in clang, so if you pass this when building with-cit will say the argument is unused.Reacted by Abhishek Bagusetty@npmiller Thanks for the pointers. We had used
-Xcuda-ptxas --maxrregcount=128(for A100) as a possible work around. We came across 2 issues:- (Major) It was still giving us the above error
- (Minor) The regcount is applied to all the kernels which might be detrimental. Since we have used
__launch_bounds__()with varied values for individual kernels.
Reacted by Nicolas MillerI hope there will be some links to the CUDA and SYCL programs in the future for studying the effect of "launch bounds" on the kernel characteristics.
Closing this.
Hi @abagusetty , this sounds like it may be a problem with index flipping. In CUDA the max dimensions of a Grid are for example:
Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535)In CUDA threads with globalIdx (0,0,0), (1,0,0) are adjacent.
Whereas since the SYCL API wants work items to be arranged in a row major fashion, meaning the work items with globalIdx (0,0,0), (0,0,1) are adjacent. Since SYCL performs this index flipping, the corresponding max global range in SYCL would be:
Max dimension size of a grid size (x,y,z): (65535, 65535, 2147483647)Sometimes when applications are ported from CUDA to SYCL the dpct does not perform this index flipping, which can cause kernel launch failures. Let me know if this is what is happening. We have some documentation about this here:
https://codeplay.com/portal/blogs/2019/11/18/computecpp-v1-1-6-changes-to-work-item-mapping-optimization.html@abagusetty you might be interested to know that it is possible to express CUDA's:
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor, maxBlocksPerCluster)with the use of following attributes:
intel::max_work_group_size,intel::min_work_groups_per_cu,intel::max_work_groups_per_mp.
For example:
cgh.single_task<class T3>( [=]() [[intel::max_work_group_size(1, 1, 256), intel::min_work_groups_per_cu(2), intel::max_work_groups_per_mp(8)]] { ... }
Subject to merging the PR: #11192
Is there a way to get/mention an equivalent functionality of
__launch_bounds__()in SYCL.Porting an optimized CUDA kernel to SYCL which preserves similar launch configuration (<<<....>>>) parameters, but without the functionality of
__launch_bounds__()in SYCL leads to the following error because ofthe kernel launch specifies too many threads for the kernel's register countAny suggestions: (a) One solution is to tweak with the global and local iteration space for the nd_range but wasn`t sure if this would be portable & performance approach when switching to other devices (i.e., PVC, MI250x, etc).
Error: