Add LTimes tuning to represent Kripke LTimes - #684
Conversation
… inner loop. Grid/block size changes to match Kripke Ltimessdom
michaelmckinsey1
changed the title
|
Yes you would have to make LTIMES and LTIMES_NO_VIEW match. |
Done |
| RAJA::statement::CudaKernelAsync< | ||
| RAJA::statement::For<1, RAJA::cuda_block_x_loop, // z | ||
| RAJA::statement::For<2, RAJA::cuda_block_y_loop, // g | ||
| RAJA::statement::For<3, RAJA::cuda_thread_x_loop, // m |
There was a problem hiding this comment.
I assume this is a non-size loop policy because it is in ltimes. Here we know the block size at compile time, is that also true in kripke?
There was a problem hiding this comment.
In Kripke, the block sizes are determined by parameters passed in at runtime. For example, like in this version of LTimes, it will be blocked on zones and groups. The exact parameters which will be blocked are not always the same for each Kripke run. For instance, if we use the DZG layout at runtime, then the loops will be blocked with directions and zones (while groups are threaded).
There was a problem hiding this comment.
The size policies tell the implementation the compile time dimensions of the block and/or grid, not the extents of the iteration space.
| RAJA::statement::For<3, RAJA::cuda_thread_x_loop, // m | |
| RAJA::statement::For<3, RAJA::cuda_thread_size_x_loop<block_size>, // m |
There was a problem hiding this comment.
Wait... I reverted f52fcfe because we want the block size to be m for the kripke tuning. This change sets the compile time block size of 25 always.
| static const size_t default_gpu_block_size = 256; | ||
| using gpu_block_sizes_type = integer::make_gpu_block_size_list_type<default_gpu_block_size, | ||
| integer::MultipleOf<32>>; | ||
| static const size_t default_gpu_block_size = 25; |
There was a problem hiding this comment.
How was a block size of 25 chosen? Would it be better for both GPU platforms to set this to 32?
There was a problem hiding this comment.
It is the square of legendre order + 1. In kripke this is also the case, the default legendre is 4, so (4+1)^2 means the kernel in kripke will be (5,5,1)
In RAJAPerf, we set m=25 directly (and this is the default). So m=36 would then be equivalent to legendre=5 in kripke.
|
Do we want to only have this kripke conforming tuning, or should does it make sense to keep the current tuning as well? |
|
@michaelmckinsey1 take a look here: https://github.com/llnl/RAJA/blob/develop/benchmark/ltimes.cpp, it would be cool to also have a GPU shared memory version as a tuning! |
9fd34d7 to
dfb8ab5
Compare
| RAJA::LaunchParams(RAJA::Teams(*num_z, *num_g, 1), | ||
| RAJA::Threads(block_size, 1, 1)), | ||
| [=] RAJA_HOST_DEVICE(RAJA::LaunchContext ctx) { | ||
|
|
There was a problem hiding this comment.
It might be worth looking into the CachePolicy for the launch_context:
https://github.com/llnl/RAJAPerf/blob/develop/src/apps/MASS3DEA-Cuda.cpp#L206-L209
We got some speedups here, this might help the nested loops
michaelmckinsey1
changed the title
I have now Implemented the Kripke implementation as a separate tuning, since we see many cases where the kripke tuning is slower. |
This reverts commit f52fcfe.
4 participants
Matching parameters like zones, groups, moments, and directions, the runtime of RAJAPerf LTimes is equivalent (<1%) to Kripke LTimes on CPU. But I have noticed runtime differences on GPU by ~15% on CUDA and ~60% on ROCm. After these changes CUDA runtime is within 1% and ROCm within 6% (CUDA is now faster, ROCm actually (~60%) slower with block 25). This is because AMD wavefront size 64 on block 25 wastes significantly more threads than NVIDIA warp size 32.
Summary
elllayout fromdcontiguous tomcontiguous to make accesses coalesced. This matches Kripke ZGD, whereField_Ellis layout orderingmomentas stride-1. Doesn't apply to psi because does not depend on m (threads access same value), phi because no d (already contiguous inm).block_m/kernel_mwhere the default block size to beminstead of 256 and remapped the kernel so z -> blockIdx.x, g -> blockIdx.y, m -> threadIdx.x. This changes the launch togrid=(num_z, num_g, 1), block=(m, 1, 1), which matches the kripke launch. However, running at different legendre orders (m) will result in different block sizes.Would need to make similar changes fordoneLTIMES-NOVIEW?kernel_m_{blocksize}kernel_zgm_{blocksize}