Add OpenAI SwiGLU benchmarks - #334
Open
zjin-lcf wants to merge 2 commits into
Open
Conversation
Add CUDA, HIP, OpenMP, and SYCL implementations of the fused SwiGLU activation with low-precision input and output modes.
There was a problem hiding this comment.
🟡 Changes recommended
CMake discovery, SYCL subgroup correctness, toolchain compatibility, and benchmark failure propagation have unresolved defects.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Pull request overview
Adds fused OpenAI/gpt-oss SwiGLU benchmarks across CUDA, HIP, SYCL, and OpenMP.
Changes:
- Adds low-precision input, output quantization, and bias paths.
- Adds shared reference and validation logic.
- Adds standalone and CMake build registrations.
File summaries
| File | Description |
|---|---|
src/swiglu-oai-sycl/Makefile |
Builds and runs SYCL cases. |
src/swiglu-oai-sycl/main.cpp |
Implements the SYCL benchmark. |
src/swiglu-oai-sycl/CMakeLists.txt |
Registers the SYCL target. |
src/swiglu-oai-omp/Makefile.nvc |
Adds NVIDIA OpenMP build support. |
src/swiglu-oai-omp/Makefile.aomp |
Adds AOMP build support. |
src/swiglu-oai-omp/Makefile |
Adds Intel OpenMP build support. |
src/swiglu-oai-omp/main.cpp |
Implements the OpenMP benchmark. |
src/swiglu-oai-omp/CMakeLists.txt |
Registers the OpenMP target. |
src/swiglu-oai-hip/reference.h |
Adds HIP-specific FP8 reference behavior. |
src/swiglu-oai-hip/Makefile |
Builds and runs HIP cases. |
src/swiglu-oai-hip/main.cu |
Implements the HIP benchmark. |
src/swiglu-oai-hip/CMakeLists.txt |
Registers the HIP target. |
src/swiglu-oai-cuda/reference.h |
Defines shared formats and reference logic. |
src/swiglu-oai-cuda/Makefile |
Builds and runs CUDA cases. |
src/swiglu-oai-cuda/main.cu |
Implements the CUDA benchmark. |
src/swiglu-oai-cuda/CMakeLists.txt |
Registers the CUDA target. |
Review details
- Files reviewed: 16/16 changed files
- Comments generated: 7
- Review effort level: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+3
to
+7
| add_hecbench_benchmark( | ||
| NAME swiglu-oai | ||
| MODEL cuda | ||
| SOURCES main.cu | ||
| CATEGORIES algorithms |
Comment on lines
+79
to
+88
| # Standard LLM inference: BF16 activation, no bias, no output quantization. \ | ||
| run_case $(ROWS) $(DIM) $(REPEAT) bf16 none none; \ | ||
| # Bias-bearing FP16 model. \ | ||
| run_case $(ROWS) $(DIM) $(REPEAT) fp16 none fp16; \ | ||
| # vLLM/SGLang-style fused SwiGLU followed by FP8 activation quantization. \ | ||
| run_case $(ROWS) $(DIM) $(REPEAT) bf16 fp8 none; \ | ||
| # Blackwell-style fused SwiGLU followed by MXFP4 activation quantization. \ | ||
| run_case $(ROWS) $(DIM) $(REPEAT) bf16 mxfp4 none; \ | ||
| # Low-precision MoE path with MXFP8 input, FP8 output, and BF16 bias. \ | ||
| run_case $(ROWS) $(DIM) $(REPEAT) mxfp8 fp8 bf16 |
| free(Y); | ||
| free(Y_ref); | ||
|
|
||
| return 0; |
Comment on lines
+223
to
+227
| const sycl::sub_group sg = item.get_sub_group(); | ||
| float amax = active ? sycl::fabs(y) : 0.f; | ||
| for (int mask = 1; mask < MX_BLOCK; mask <<= 1) | ||
| amax = sycl::fmax( | ||
| amax, sycl::permute_group_by_xor(sg, amax, mask)); |
Comment on lines
+14
to
+16
| #elif (__CUDACC_VER_MAJOR__ < 12) || \ | ||
| (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ < 8) | ||
| #error "swiglu-oai-cuda requires CUDA Toolkit 12.8 or newer for FP4 support" |
Comment on lines
+28
to
+31
| #elif HIP_VERSION_MAJOR < 7 | ||
| #error "swiglu-oai-hip requires ROCm 7.0 or newer for hip_fp4.h" | ||
| #else | ||
| #include <hip/hip_fp4.h> |
| float* Ydata) | ||
| { | ||
| #pragma omp target teams distribute parallel for num_threads(block_size) | ||
| for (int index = 0; index < rows * dim; index++) { |
- register swiglu-oai in the master CMake benchmark list so the four model directories are actually configured - move the CUDA `make run` case annotations out of the backslash-continued recipe, where the first comment swallowed every following run_case - build on CUDA toolkits older than 12.8 and ROCm older than 7.0 by falling back to software E2M1/E8M0 codecs instead of #error - SYCL: request a sub-group of MX_BLOCK lanes for the MXFP4 epilogue where the device offers it, keep the plain kernel for wider-only sub-groups (wave64), and skip devices that may pick a narrower one - OpenMP: iterate the fused kernel over a 64-bit element count so large shapes no longer overflow rows * dim as int Co-authored-by: Cursor <cursoragent@cursor.com>
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
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