Skip to content

[ARK] Optimize SYCL-TLA woqgemm_s8 on B70 - #2133

Open
Zhenzhong1 wants to merge 5 commits into
mainfrom
zhenzhong/woqgemm_s8_update
Open

[ARK] Optimize SYCL-TLA woqgemm_s8 on B70#2133
Zhenzhong1 wants to merge 5 commits into
mainfrom
zhenzhong/woqgemm_s8_update

Conversation

@Zhenzhong1

@Zhenzhong1 Zhenzhong1 commented Aug 7, 2026

Copy link
Copy Markdown
Contributor

Desc

  • Optimize SYCL-TLA woqgemm_s8.
  • ARK_RESCALE is enabled by default in setup.py.
  • Dectect machine (B60/B70) in setup.py

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/s

SYCL-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/s

Reference 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

Signed-off-by: Zhenzhong1 <zhenzhong.xu@intel.com>
Signed-off-by: Zhenzhong1 <zhenzhong.xu@intel.com>
@Zhenzhong1 Zhenzhong1 changed the title [ARK] Optimize SYCL-TLA woqgemm_s8 [ARK] Optimize SYCL-TLA woqgemm_s8 on B60&B70 Aug 10, 2026
@Zhenzhong1 Zhenzhong1 changed the title [ARK] Optimize SYCL-TLA woqgemm_s8 on B60&B70 [ARK] Optimize SYCL-TLA woqgemm_s8 on B70 Aug 11, 2026
Signed-off-by: Zhenzhong1 <zhenzhong.xu@intel.com>
@chensuyue chensuyue added this to the 0.15.0 milestone Aug 11, 2026
@Zhenzhong1
Zhenzhong1 marked this pull request as ready for review August 11, 2026 07:55

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 {
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants