Modified CUDA/HIP reduction algorithm#62

Open

vikaskurapati wants to merge 3 commits intomasterfrom

vikas/cudareduction

Contributor

vikaskurapati commented Feb 25, 2026 •

edited

Loading

Modified CUDA reduction algorithm based on #56. Tested, and benchmarked comparing it with the current implementation on Hopper using Vista's GPU; the average runtime is around 1.5x-6x faster just for the reduction on Hopper. I did not test it with SeisSol, and as reduction is only a small component of SeisSol, I do not anticipate any significant speed-up in SeisSol runs due to this.

The WorkGroupSize, and ItemsPerWorkItem are heuristically picked from benchmarking after running it with different configurations. The results of best configurations are here depending on vector sizes, and data types. I chose currently 1024, 4 on basis of 1e5-1e6 for float -- which are typical time cluster sizes for big problems we use. But if there is a better idea for choosing these, please let me know.

=== Benchmark Comparison for Type: float (Time Metric) ===
Vector Size | Old Time (ms) | New Time (ms) | Speedup | Best Config (WG, Items)

1e+00 | 0.0030 | 0.0041 | 0.72x | (512, 8)
1e+01 | 0.0030 | 0.0041 | 0.74x | (128, 4)
1e+02 | 0.0029 | 0.0042 | 0.69x | (128, 4)
1e+03 | 0.0029 | 0.0043 | 0.67x | (256, 4)
1e+04 | 0.0054 | 0.0044 | 1.24x | (256, 8)
1e+05 | 0.0304 | 0.0044 | 6.86x | (1024, 4)
1e+06 | 0.2783 | 0.0048 | 57.49x | (1024, 4)
1e+07 | 3.8529 | 0.0111 | 348.49x | (512, 16)
1e+08 | 46.2797 | 0.1087 | 425.86x | (256, 16)
1e+09 | 461.8960 | 1.0327 | 447.28x | (256, 16)

=== Benchmark Comparison for Type: int (Time Metric) ===
Vector Size | Old Time (ms) | New Time (ms) | Speedup | Best Config (WG, Items)

1e+00 | 0.0030 | 0.0041 | 0.73x | (256, 8)
1e+01 | 0.0029 | 0.0042 | 0.69x | (256, 8)
1e+02 | 0.0029 | 0.0041 | 0.71x | (128, 8)
1e+03 | 0.0030 | 0.0040 | 0.74x | (128, 4)
1e+04 | 0.0054 | 0.0043 | 1.27x | (128, 8)
1e+05 | 0.0307 | 0.0043 | 7.07x | (512, 4)
1e+06 | 0.2807 | 0.0048 | 58.37x | (512, 8)
1e+07 | 3.8434 | 0.0107 | 358.10x | (512, 8)
1e+08 | 46.2018 | 0.1086 | 425.34x | (256, 16)
1e+09 | 461.2680 | 1.0323 | 446.82x | (512, 16)

=== Benchmark Comparison for Type: double (Time Metric) ===
Vector Size | Old Time (ms) | New Time (ms) | Speedup | Best Config (WG, Items)

1e+00 | 0.0030 | 0.0042 | 0.73x | (128, 8)
1e+01 | 0.0029 | 0.0040 | 0.73x | (256, 8)
1e+02 | 0.0030 | 0.0041 | 0.72x | (256, 8)
1e+03 | 0.0030 | 0.0043 | 0.70x | (512, 4)
1e+04 | 0.0061 | 0.0043 | 1.41x | (512, 4)
1e+05 | 0.0335 | 0.0045 | 7.44x | (256, 4)
1e+06 | 0.3053 | 0.0056 | 54.97x | (512, 8)
1e+07 | 4.9312 | 0.0262 | 188.43x | (256, 8)
1e+08 | 49.1560 | 0.2110 | 232.95x | (256, 8)
1e+09 | 491.2810 | 2.0570 | 238.83x | (128, 16)


          Reduction for CUDA

bc193c3

bug fix
clang-format
modified workgroup sizes

vikaskurapati requested a review from davschneller

February 25, 2026 15:31

davschneller approved these changes

View reviewed changes

Contributor

davschneller left a comment •

edited

Loading

LGTM.

Large workgroup sizes are a good idea—just to keep enough data in flight in case some other wave/warp has to wait.

algorithms/cudahip/Reduction.cpp Outdated

-                dim3 grid(1, 1, 1);
-                dim3 block(1024, 1, 1);
+                size_t totalItems = WorkGroupSize * ItemsPerWorkItem;

Contributor

davschneller Feb 25, 2026

Actually more an AMD name :) (NVIDIA would call it a "block")
Also make const; same for the block count

algorithms/cudahip/Reduction.cpp Show resolved Hide resolved

algorithms/cudahip/Reduction.cpp

-                for (int offset = 1; offset < warpSize; offset *= 2) {
+                for (int offset = warpSize / 2; offset > 0; offset /= 2) {
                   value = operation(value, shuffledown(value, offset));

Contributor

davschneller Feb 25, 2026 •

edited

Loading

Just as a side note; you could use cub in the CUDA case here. (there are sometimes, e.g. for min/max more advanced reduction functions starting with Ampere) That should come with CUDA included.

For the AMD case, there should be an equivalent library (which might require explicit linking, however) — or you could use the OCKL functions (as long as they exist). See here for a forward as it's used: https://github.com/ROCm/clr/blob/b90c29358c694e66ea78cb1e3957edad09f35cbf/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h#L57-L92 (you can almost just take a copy of that declaration code for the relevant functions—or try to use the HIP functions below ... EDIT: found the docs: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_cpp_language_extensions.html )

... though I'm not sure if there's much speedup to gain in either case.
And it could also be put in some future PR.

Contributor Author

vikaskurapati Mar 2, 2026

I am not sure if I understand it correctly, and in https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_cpp_language_extensions.html -- I found the __reduce functions for warps only for int or unsigned int types. So, I added these if conditions, I am unsure how benificial they would be. If you suggest removing this, and keeping the manual ones, I will remove them.

Contributor Author

vikaskurapati Mar 3, 2026

I am not sure if I understand it correctly, and in https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_cpp_language_extensions.html -- I found the __reduce functions for warps only for int or unsigned int types. So, I added these if conditions, I am unsure how benificial they would be. If you suggest removing this, and keeping the manual ones, I will remove them.

Surprisingly, these functions are not recognized by HIP in the CI, and they fail, while the NVIDIA tests pass silently. I am a bit confused, but for now, I removed this for HIP too. If there is a particular way to get this done, please let me know. :)

algorithms/cudahip/Reduction.cpp Outdated

+                const int warpId = threadIdx.x / warpSize;
+                val = warpReduce(val, operation);
+                if (laneId == 0)

Contributor

davschneller Feb 25, 2026

Brackets

algorithms/cudahip/Reduction.cpp Outdated

    
                  auto value = operation.defaultValue;

                  auto lastAcc = operation.defaultValue;

                if (warpId == 0)

Contributor

davschneller Feb 25, 2026

Brackets (it's too easy to slip in another statement between those two at some point otherwise)

algorithms/cudahip/Reduction.cpp

+                val = warpReduce(val, operation);
+                if (laneId == 0)
+                  shmem[warpId] = val;

Contributor

davschneller Feb 25, 2026 •

edited

Loading

In principle, you could already use the atomics here.
That's most likely how e.g. SYCL did it internally (or the respective OCKL reductions would do it, or maybe cub has something as well—though I'm not sure how public those are).

(though IIRC there was a performance regression for FP32 on the MI250X in LDS, and only there with that)

Contributor Author

vikaskurapati Mar 2, 2026

Thanks for pointing it out. However, given that exact LDS FP32 regression you mentioned on the MI250X, I think sticking to the two-pass shuffle (Warp 0 sweeping the shared memory) seems like a safer bet for now. But if you insist on implementing it that way, I can try in a different PR, but I will need to benchmark the differences, and I don't have access to LUMI as of now.

davschneller changed the title ~~Modified CUDA reduction algorithm~~ Modified CUDA/HIP reduction algorithm


          Address review comments: rename to NVIDIA convention, add brackets, m…

18a86c0

…ake const

- Rename WorkGroupSize -> BlockSize and ItemsPerWorkItem -> ItemsPerThread
  (WorkGroup/WorkItem are AMD/SYCL terminology; block/thread are NVIDIA terms)
- Add curly brackets to if (laneId == 0) and if (warpId == 0) bodies in blockReduce
- Make totalItems and numBlocks const in reduceVector

vikaskurapati requested a review from davschneller

March 2, 2026 09:20

vikaskurapati force-pushed the vikas/cudareduction branch 2 times, most recently from 398d6b7 to 0192755 Compare

March 3, 2026 10:51


          vendor warp reduce functions

ffc30c4

Format and fix
compilation fix
remove hip for warp reduce functions

vikaskurapati force-pushed the vikas/cudareduction branch from 0192755 to ffc30c4 Compare

March 3, 2026 10:58

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet