2.1.3

clpeak 2.1.3

An Intel Arc release, plus new FP8 and INT8 sparsity coverage on CUDA. Several kernels were reporting the wrong rate on Arc, on CPU OpenCL runtimes, and for CUDA INT8 sparsity — some numbers move on those devices in both directions (see Measurement corrections).

Intel Arc (fixes #202)

Four separate faults were behind one report — an A380 that crashed under Vulkan, and a GUI that appeared to ignore the device selection made to work around it.

• Vulkan device features are now actually enabled. vkCreateDevice was handed feature structs nobody filled in — shaderFloat64 was the only VK_TRUE in the file, while Float16, Int8, BFloat16, VulkanMemoryModel, CooperativeMatrixKHR and Float8 were all requested as VK_FALSE with their extensions enabled. That is invalid usage; NVIDIA and AMD ran the shaders anyway, and Arc faulted inside the shader compiler on coopmat_int8.comp, the shader with the most unenabled capabilities.
• Cooperative-matrix stages are pinned to a 32-wide subgroup via VK_EXT_subgroup_size_control where the driver allows it. The shaders compute one tile per subgroup and store it at the work-group's offset, so a driver that splits the 32-thread group into 2 subgroups (Intel SIMD16) or 4 (SIMD8) recomputes the same tile into the same address and reports the rate divided by that factor. Falls back to an unpinned pipeline if the width is rejected.
• fp16/bf16/int8 coopmat rows are gated on the matching type feature, not just on the advertised tile — as the fp8 rows already were.
• The Run button in the GUI's Custom window is now a full-width block below the status line, instead of a small right-aligned pill parked next to the tab strip's bottom rule where it read as chrome and got skipped over.

Racing two MAD-chain shapes

Every compute test built on the shared x = x*x + c chain reported exactly half of Alchemist/Xe-HPG's real throughput: a three-source mad issues at full rate on Arc only when all three source operands are distinct registers, and x = x*x + c reads {x, x, c}. NVIDIA has the mirror-image constraint, so no single chain shape reaches peak on every vendor.

Vulkan, OpenCL, oneAPI and Metal now compile each compute kernel in two chain shapes, time both, and report the faster.

• Float families use an affine second shape (x = a*x + b). Integer families use a rotating shape (x_k = x_k * x_(k+1) + c) instead — an integer affine recurrence is legally foldable and Apple's OpenCL compiler does fold it, which came back as a 15.5x inflated int16 reading in testing. The rotating shape keeps three distinct registers while staying quadratic, so there is no closed form to fold to.
• Mixed-precision keeps the narrow type in the data path (x_k = narrow(m * wide(x_k) + b)) with exactly one narrowing conversion per mad — the same ratio the squaring kernels carry, so chains can be added without handing the alt build extra work. The narrowing also blocks any fold outright.
• A MAX_ALT_CHAIN_RATIO backstop (6x) rejects any alt reading that far outpaces the squaring one, in case a compiler finds a fold nobody predicted. Measured legitimate gains top out near 4x.
• CUDA and ROCm deliberately do not race, and common.h now records why: each targets one vendor, racing costs ~2x the compute-test budget, and on NVIDIA the squaring chain measures optimal (1.00 against 0.52 for a single affine chain, on both a 5060 and a 4060). AMD is unmeasured.

Verified with a standalone probe across an Arc A380, RTX 5060/4060, UHD 630, Adreno X1-45 and an Intel CPU OpenCL runtime: the racing pair lands within 3% of the best measured shape everywhere, against 50% off on Alchemist and 25% off on CPU runtimes before. Every OpenCL alt family is now syntax-checked with clang's OpenCL frontend at build time — overdue, since the hp, dp and mp alt kernels could not be compiled on the development machine at all.

New CUDA tests (fixes #203)

Test	Row	RTX 5060 (sm_120)
FP8(E4M3) mma.sync m16n8k32 +fp16	fp8_e4m3_f16acc	166.81 TFLOPS
FP8(E4M3) mma.sp 2:4 sparsity m16n8k64 +fp32	fp8_sparse	169.74 TFLOPS
FP8(E4M3) mma.sp 2:4 sparsity m16n8k64 +fp16	fp8_sparse_f16acc	326.15 TFLOPS

Together with the existing dense +fp32 row (85.13) these complete the FP8 grid and answer two questions: the fp32-accumulate half-rate cap that governs FP16 on GeForce governs FP8 too, and consumer Blackwell does accelerate FP8 2:4 sparsity at full rate. Sparsity and accumulator width are independent 2x multipliers and they compose — 3.83x the dense fp32-accumulate corner.

The dense FP8 rows are gated at sm_89+. fp8_sparse_f16acc is built only for sm_120a/121a: ptxas rejects the form for sm_89, and sm_90/sm_100 are untried, so the floor is provisional — Hopper currently reports Skipped rather than being tried.

Measurement corrections

• CUDA INT8 sparsity was measured at the wrong shape. int8_sparse was pinned to sparse m16n8k32, whose dense counterpart is m16n8k16 — not the k32 row it was printed next to. Against a same-K dense row a sparse kernel has no doubled K to turn into a 2x, whatever the hardware does, and the flat result was recorded as "GeForce gates the INT sparse data-path". It does not: at sparse m16n8k64 the RTX 5060 reads 327.30 TOPS, 1.99x dense k32 and 3.88x dense k16. The claim was corrected in the two FP4 sparse kernels as well.
• OpenCL mixed-precision was inflated 8x on work-item-vectorising CPU runtimes. compute_mp was the only family in any backend whose kernels derived every operand from the uniform kernel argument — CUDA, ROCm, Vulkan, Metal and oneAPI all seed from the thread or lane id. A fully uniform loop can be computed once and broadcast across the work-item vector, and that is what was happening: adding the per-work-item term to the fp32 accumulator seed dropped the readings by 7.96x, 7.99x, 7.96x and 8.02x at vector widths 2, 4, 8 and 16 — precisely the AVX2 fp32 lane count. On an Intel CPU runtime mp now reads 378 GFLOPS at best against 187 for pure half on the same device, where before it claimed 2077 — eleven times the fp16 rate it is built on. Scalar width 1 was inflated only 1.37x. GPUs are unaffected, and no saved baseline in results/ was: the mp rows there are Vulkan.
• Intel Arc and CPU OpenCL runtimes gain up to 2x across the compute families, from the chain racing above; Intel Vulkan coopmat rows gain the subgroup-split factor.
• oneAPI mp and bf16 go from one chain to four. They are as latency-bound as the rest and were simply not being given anything to overlap: the four-chain float path on an Intel CPU runtime went 396 → 1552 GFLOPS. Vulkan mp stays at 4/2/1 chains — a change to one chain in development cost an RTX 5060 10.6% on mp2 and was reverted.

Fixes

• clpeak --version no longer reports a stale hardcoded number (#198). Builds with no git (tarballs, no git binary) were pinned to whatever version was last hardcoded, silently misrepresenting the build as a specific release; they now report unknown. A git-describe that succeeds with empty output is guarded too.
• OpenCL int8_dp availability is detected properly. cl_khr_integer_dot_product advertises two independent input forms and a device may implement only the packed one, while the kernels need the unpacked 4x8-bit form — so gating on the extension string alone turned an honest "unsupported" into five clCreateKernel errors on Arc (-44) and Intel's CPU runtime (-46). The gate now reads CL_DEVICE_INTEGER_DOT_PRODUCT_CAPABILITIES_KHR and also accepts the OpenCL 3.0 optional feature __opencl_c_integer_dot_product_input_4x8bit. The KHR enums are defined locally behind #ifndef so the build still works against the OpenCL 2.2 headers ROCm ships.
• As a backstop, CL_INVALID_KERNEL_NAME and CL_INVALID_PROGRAM now map to a clean Unsupported skip: a kernel the compiler declined to build is a capability fact, not a failure worth five error rows.

Build and CI

• ROCm is discovered via HIP_PATH / ROCM_PATH, not just /opt/rocm. Discovery was hardcoded to the Linux default prefix in three places, so a Windows HIP SDK or a relocated Linux install was invisible no matter what was present. hipcc.bat joins the find_program names, since find_program only auto-appends .exe/.com. No behaviour change where ROCm lives at /opt/rocm.
• New release artifacts: clpeak-2.1.3-windows-AMD64-rocm.zip and clpeak-2.1.3-windows-AMD64-oneapi.zip. The AOT kernel design is what makes the ROCm one cheap — src/rocm links only hip::host and every .hip goes through hipcc --genco, so no host code touches the HIP device compiler. Note the oneAPI zips (Windows and Linux alike) are not self-contained: they need oneAPI installed for the SYCL runtime libraries. hipBLASLt is Linux-only and rocWMMA may not ship in the Windows HIP SDK; both degrade through existing feature gates.
• Both new rows are pinned to windows-2022: AMD publishes the HIP SDK only as a Server 2022 build, and MSVC's STL in the VS2026 image enforces P2652R2 through a clang-only attribute that icx honours and cl.exe does not, which Intel's SYCL headers trip over. The plain and CUDA Windows rows stay on windows-latest.
• Windows hipcc kernel compiles fixed by scrubbing INCLUDE/EXTERNAL_INCLUDE before invoking hipcc. MSBuild's custom-build step injects the full MSVC environment; clang honours INCLUDE when targeting the MSVC toolchain, and ROCm 6.4's bundled clang 20 choked on MSVC 14.44's vcruntime.h (unknown type name 'uintptr_t'). Applied to the configure-time arch probe as well, which fixes a latent bug where configuring from a VS developer prompt would have failed every probe and silently emitted stub kernels for the whole backend.
• A guard now fails the ROCm and oneAPI builds if the SDK did not resolve. Without it, an SDK CMake failed to find degraded to a green build shipping a plain clpeak under the -rocm artifact name.
• build.yml split into per-platform reusable workflows (desktop.yml, linux-gcc.yml, android.yml, ios.yml) plus composite setup actions under .github/actions/setup-*, taking the caller from ~440 lines to 79. Pure refactor — step bodies, matrix, triggers and release steps are byte-identical. Note for branch protection: reusable-workflow jobs report as <caller-job> / <job-name>, so required check names change (e.g. Linux x64 → desktop / Linux x64).

Other

• Refreshed reference results for RTX 5060, M1 Pro and Threadripper PRO 3955WX; the Threadripper file gains a oneAPI run.
• The Homebrew formula's comment on why it uses a git URL was updated for the unknown version fallback.

Full changelog: 2.1.2...2.1.3