[ARK] Optimize SYCL-TLA woqgemm_s8 on B70 - #2133
Open
Zhenzhong1 wants to merge 5 commits into
Open
Conversation
Signed-off-by: Zhenzhong1 <zhenzhong.xu@intel.com>
Signed-off-by: Zhenzhong1 <zhenzhong.xu@intel.com>
Signed-off-by: Zhenzhong1 <zhenzhong.xu@intel.com>
Zhenzhong1
marked this pull request as ready for review
August 11, 2026 07:55
Contributor
There was a problem hiding this comment.
Pull request overview
Optimizes ARK’s SYCL-TLA int8 weight-only GEMM for B70 and adds device-specific build targeting.
Changes:
- Adds B60/B70/PVC target detection and enables rescaling.
- Fuses K-block accumulation, dequantization, and output finalization.
- Adds G21/G31 compiler target mappings.
Reviewed changes
Copilot reviewed 4 out of 4 changed files in this pull request and generated 3 comments.
| File | Description |
|---|---|
auto_round_extension/ark/setup.py |
Detects SYCL targets and configures rescaling. |
auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_s8_gemm.hpp |
Implements optimized fused K-block GEMM. |
auto_round_extension/ark/auto_round_kernel/CMakeLists.txt |
Maps GPU targets for the primary kernel build. |
auto_round_extension/ark/auto_round_kernel/bestla/CMakeLists.txt |
Extends benchmark target mappings. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+110
to
+111
| int k_tiles_per_block = blocksize / k_tile_size; | ||
| int k_tile_count = blks * k_tiles_per_block; |
Comment on lines
+116
to
120
| if(DPCPP_SYCL_TARGET STREQUAL "intel_gpu_bmg_g21") | ||
| set(SYCL_DEVICE_NAME "bmg-g21") | ||
| elseif(DPCPP_SYCL_TARGET STREQUAL "intel_gpu_bmg_g31") | ||
| set(SYCL_DEVICE_NAME "bmg-g31") | ||
| elseif(DPCPP_SYCL_TARGET STREQUAL "intel_gpu_pvc" OR DPCPP_SYCL_TARGET STREQUAL "pvc") |
| launch_igemm_kblock_tile<false, ElementOut, 128, 128, MediumTileSG>( | ||
| q, m, n, k, a, b, c, scale_a, scale_b, bias, blocksize, blks); | ||
| } | ||
| } else { |
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.
Desc
Perf
SYCL-TLA on B70
=== W4A8 / A8W4 === m=4096, n=4096, k=4096, blocksize=4096, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.085938, Mean Diff: 0.007950, ref mean:7.507812, OUT mean:7.507812 [Performance] Time: 0.7353 ms GFLOPS: 186918.01 Bandwidth: 102.69 GB/s m=4096, n=4096, k=4096, blocksize=128, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.066406, Mean Diff: 0.008209, ref mean:7.546875, OUT mean:7.546875 [Performance] Time: 0.9017 ms GFLOPS: 152419.01 Bandwidth: 84.02 GB/s m=4096, n=4096, k=4096, blocksize=64, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.062500, Mean Diff: 0.008217, ref mean:7.546875, OUT mean:7.546875 [Performance] Time: 1.0235 ms GFLOPS: 134284.96 Bandwidth: 74.28 GB/s m=4096, n=4096, k=4096, blocksize=32, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.007812, Mean Diff: 0.001299, ref mean:7.507812, OUT mean:7.507812 [Performance] Time: 1.2077 ms GFLOPS: 113804.86 Bandwidth: 63.38 GB/sSYCL-TLA on B60
=== W4A8=== m=4096, n=4096, k=4096, blocksize=4096, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.078125, Mean Diff: 0.007988, ref mean:7.476562, OUT mean:7.476562 [Performance] Time: 1.0585 ms GFLOPS: 129838.74 Bandwidth: 71.33 GB/s m=4096, n=4096, k=4096, blocksize=128, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.062500, Mean Diff: 0.008240, ref mean:7.500000, OUT mean:7.500000 [Performance] Time: 1.4031 ms GFLOPS: 97953.18 Bandwidth: 53.99 GB/s m=4096, n=4096, k=4096, blocksize=64, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.062500, Mean Diff: 0.008240, ref mean:7.480469, OUT mean:7.480469 [Performance] Time: 1.5445 ms GFLOPS: 88987.05 Bandwidth: 49.22 GB/s m=4096, n=4096, k=4096, blocksize=32, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.007812, Mean Diff: 0.001297, ref mean:7.511719, OUT mean:7.511719 [Performance] Time: 1.7683 ms GFLOPS: 77723.89 Bandwidth: 43.29 GB/sReference Joint_matrix on B60
=== W4A8 / A8W4 === m=4096, n=4096, k=4096, blocksize=4096, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.078125, Mean Diff: 0.007988, ref mean:7.476562, OUT mean:7.476562 [Performance] Time: 1.1099 ms GFLOPS: 123827.48 Bandwidth: 68.03 GB/s m=4096, n=4096, k=4096, blocksize=128, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.062500, Mean Diff: 0.008240, ref mean:7.500000, OUT mean:7.500000 [Performance] Time: 1.5408 ms GFLOPS: 89199.99 Bandwidth: 49.17 GB/s m=4096, n=4096, k=4096, blocksize=64, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.062500, Mean Diff: 0.008240, ref mean:7.480469, OUT mean:7.480469 [Performance] Time: 1.9249 ms GFLOPS: 71400.29 Bandwidth: 39.49 GB/s m=4096, n=4096, k=4096, blocksize=32, compute_type=int8, weight_type=int4, scale_type=fp16, asym=False, device=xpu Max Diff: 0.007812, Mean Diff: 0.001297, ref mean:7.511719, OUT mean:7.511719 [Performance] Time: 1.5898 ms GFLOPS: 86448.29 Bandwidth: 48.15 GB/s