Script 'mail_helper' called by obssrc Hello community, here is the log from the commit of package clpeak for openSUSE:Factory checked in at 2026-08-17 16:59:08 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Comparing /work/SRC/openSUSE:Factory/clpeak (Old) and /work/SRC/openSUSE:Factory/.clpeak.new.1258 (New) ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
Package is "clpeak" Mon Aug 17 16:59:08 2026 rev:16 rq:1371494 version:2.0.19 Changes: -------- --- /work/SRC/openSUSE:Factory/clpeak/clpeak.changes 2026-07-13 14:28:31.056341514 +0200 +++ /work/SRC/openSUSE:Factory/.clpeak.new.1258/clpeak.changes 2026-08-17 17:02:34.226616428 +0200 @@ -1,0 +2,26 @@ +Mon Aug 17 05:38:00 UTC 2026 - Martin Pluskal <[email protected]> + +- Update to 2.0.19: + * OpenCL: replace the mad() builtin with a contracted (y*x)+y + expression in the fp32/fp64/fp16 compute kernels. The spec only + permits, but does not require, mad() to fuse into a hardware + FMA, and Apple's CL-on-Metal compiler never lowered it, halving + the reported fp32 throughput there; other platforms are + unaffected because they already fused it. + * CUDA/ROCm: mark the global-bandwidth kernel in/out pointers + __restrict__, so the compiler may batch the loads ahead of the + trailing store and NVIDIA can take the read-only ld.global.nc + path. + * CPU backend: fix the GCC build failure on the AArch64 crypto ISA + extensions (gh#krrishnarraj/clpeak#189). GCC exposes the + AES/SHA-256 NEON intrinsics only under the umbrella +crypto + target attribute, so a translation unit compiled with just +aes + or +sha2 failed to inline them; GCC now gets +crypto + (+crypto+sha3 for the SHA-512 unit) while clang keeps the + granular flags. +- Correct License to Apache-2.0: upstream relicensed away from the + Unlicense in December 2023 and the bundled LICENSE has shipped the + Apache License 2.0 ever since, but the spec still declared + LicenseRef-SUSE-Public-Domain. + +------------------------------------------------------------------- Old: ---- clpeak-2.0.18.tar.gz New: ---- clpeak-2.0.19.tar.gz ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ Other differences: ------------------ ++++++ clpeak.spec ++++++ --- /var/tmp/diff_new_pack.yHxWuP/_old 2026-08-17 17:02:34.784636144 +0200 +++ /var/tmp/diff_new_pack.yHxWuP/_new 2026-08-17 17:02:34.788636286 +0200 @@ -18,10 +18,14 @@ Name: clpeak -Version: 2.0.18 +Version: 2.0.19 Release: 0 Summary: Find peak OpenCL capacities like bandwidth & compute -License: LicenseRef-SUSE-Public-Domain +# Legal-Review-Notice: upstream relicensed from the Unlicense to Apache-2.0 +# in commit 65e57245 (2023-12-22); LICENSE is the sole licence file in the +# tree, there is no bundled third-party code and no source carries a +# divergent SPDX header. Fedora declares Apache-2.0 as well. +License: Apache-2.0 URL: https://github.com/krrishnarraj/clpeak Source: https://github.com/krrishnarraj/clpeak/archive/refs/tags/%{version}.tar.gz#/%{name}-%{version}.tar.gz BuildRequires: cmake ++++++ clpeak-2.0.18.tar.gz -> clpeak-2.0.19.tar.gz ++++++ diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/clpeak-2.0.18/.github/workflows/build.yml new/clpeak-2.0.19/.github/workflows/build.yml --- old/clpeak-2.0.18/.github/workflows/build.yml 2026-07-11 18:02:40.000000000 +0200 +++ new/clpeak-2.0.19/.github/workflows/build.yml 2026-08-15 09:13:05.000000000 +0200 @@ -223,6 +223,52 @@ path: build/clpeak-*.zip # --------------------------------------------------------------------------- + # Linux x64 GCC — build-error smoke test only. + # + # The shipped Linux x64 artifact is built with clang (the root CMake prefers it + # for the CPU backend's FMA-chain scheduling), so GCC's stricter always_inline + # target-attribute checks — e.g. the ARM NEON crypto intrinsics, and GCC-only + # diagnostics generally — never get exercised in CI otherwise. This job forces + # a GCC/g++ build to catch those. It is deliberately NOT wired into `release` + # (release depends on build-desktop only): no packaging, no artifact upload. + # clang is purged so the run genuinely mirrors a GCC-only toolchain, and g++ is + # also pinned via -DCMAKE_*_COMPILER so the build can't silently fall back. + # --------------------------------------------------------------------------- + build-linux-gcc: + name: Linux x64 GCC + runs-on: ubuntu-latest + + steps: + - name: Checkout repository + uses: actions/checkout@v4 + with: + fetch-depth: 0 # full history for git-describe versioning + + - name: Install dependencies (GCC only, no clang) + run: | + sudo apt-get update + # Purge any clang the runner image ships with, so the build can only use + # GCC (the whole point of this job). '|| true' keeps it green if clang + # was never installed. + sudo apt-get remove -y --purge 'clang*' || true + sudo apt-get install -y --no-install-recommends build-essential g++ + + - name: Install Vulkan SDK + uses: jakoch/install-vulkan-sdk-action@v1 + with: + install_runtime: true + cache: true + + - name: Configure CMake & Build (GCC) + run: | + cmake -B build -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_C_COMPILER=gcc -DCMAKE_CXX_COMPILER=g++ + cmake --build build --config Release -j4 + + - name: Run + run: ./build/clpeak + + # --------------------------------------------------------------------------- # Android # --------------------------------------------------------------------------- build-android: diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/clpeak-2.0.18/src/cpu/CMakeLists.txt new/clpeak-2.0.19/src/cpu/CMakeLists.txt --- old/clpeak-2.0.18/src/cpu/CMakeLists.txt 2026-07-11 18:02:40.000000000 +0200 +++ new/clpeak-2.0.19/src/cpu/CMakeLists.txt 2026-08-15 09:13:05.000000000 +0200 @@ -414,17 +414,38 @@ # Crypto TUs: +aes = the AES instructions, +sha2 = SHA-256, +sha3 = the # SHA-512 (and SHA-3) instructions, +crc = CRC32/CRC32C. All fixed- # width NEON/GPR types, so they build on Windows clang-cl too. - check_cxx_compiler_flag("${_clpeak_gnuflag}-march=armv8-a+aes" _clpeak_f_a_aes) + # + # GCC's <arm_neon.h> declares every AES/SHA-256 intrinsic under one + # `+crypto` target attribute (aes AND sha2 together, on the v8.0 crypto + # base), and its always_inline rule needs the intrinsic's feature set to + # be a SUBSET of the caller's. A TU built with only the granular +aes + # (missing sha2) or only +sha2 (missing aes) therefore can't inline the + # intrinsic -> "target specific option mismatch" hard error; the sha512 + # TU hits it too, on the SHA-256 kernel it also carries (+sha3 implies + # sha2 but not aes). Give GCC the umbrella +crypto so the caller is + # always a superset. Clang declares each intrinsic with its own granular + # feature, so keep the split flags there -- +crypto is deprecated on + # clang and would only warn. + if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU") + set(_aes_ext "+crypto") + set(_sha2_ext "+crypto") + set(_sha3_ext "+crypto+sha3") + else() + set(_aes_ext "+aes") + set(_sha2_ext "+sha2") + set(_sha3_ext "+sha3") + endif() + check_cxx_compiler_flag("${_clpeak_gnuflag}-march=armv8-a${_aes_ext}" _clpeak_f_a_aes) if(_clpeak_f_a_aes) - clpeak_add_isa_tu(aes ${_clpeak_gnuflag}-march=armv8-a+aes) + clpeak_add_isa_tu(aes ${_clpeak_gnuflag}-march=armv8-a${_aes_ext}) endif() - check_cxx_compiler_flag("${_clpeak_gnuflag}-march=armv8-a+sha2" _clpeak_f_a_sha2) + check_cxx_compiler_flag("${_clpeak_gnuflag}-march=armv8-a${_sha2_ext}" _clpeak_f_a_sha2) if(_clpeak_f_a_sha2) - clpeak_add_isa_tu(sha ${_clpeak_gnuflag}-march=armv8-a+sha2) + clpeak_add_isa_tu(sha ${_clpeak_gnuflag}-march=armv8-a${_sha2_ext}) endif() - check_cxx_compiler_flag("${_clpeak_gnuflag}-march=armv8.2-a+sha3" _clpeak_f_a_sha3) + check_cxx_compiler_flag("${_clpeak_gnuflag}-march=armv8.2-a${_sha3_ext}" _clpeak_f_a_sha3) if(_clpeak_f_a_sha3) - clpeak_add_isa_tu(sha512 ${_clpeak_gnuflag}-march=armv8.2-a+sha3) + clpeak_add_isa_tu(sha512 ${_clpeak_gnuflag}-march=armv8.2-a${_sha3_ext}) endif() check_cxx_compiler_flag("${_clpeak_gnuflag}-march=armv8-a+crc" _clpeak_f_a_crc) if(_clpeak_f_a_crc) diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/clpeak-2.0.18/src/cuda/cuda_kernels/global_bandwidth.cu new/clpeak-2.0.19/src/cuda/cuda_kernels/global_bandwidth.cu --- old/clpeak-2.0.18/src/cuda/cuda_kernels/global_bandwidth.cu 2026-07-11 18:02:40.000000000 +0200 +++ new/clpeak-2.0.19/src/cuda/cuda_kernels/global_bandwidth.cu 2026-08-15 09:13:05.000000000 +0200 @@ -4,10 +4,15 @@ // measure VRAM rather than L2. CUDA's widest native load is LDG.128 // (float4), so float8/float16 would lower to multiple float4 loads with // no new HW path -- we stop at float4. +// +// `in`/`out` are __restrict__ so nvcc can prove they do not alias and issue +// the fetches as read-only ld.global.nc; without it the trailing store forces +// the conservative coherent load path. #define FETCH_PER_WI 16 -extern "C" __global__ void global_bandwidth_v1(const float *in, float *out) +extern "C" __global__ void global_bandwidth_v1(const float * __restrict__ in, + float * __restrict__ out) { unsigned int gid = blockIdx.x * blockDim.x + threadIdx.x; unsigned int lid = threadIdx.x; @@ -25,7 +30,8 @@ out[gid] = sum; } -extern "C" __global__ void global_bandwidth_v2(const float2 *in, float *out) +extern "C" __global__ void global_bandwidth_v2(const float2 * __restrict__ in, + float * __restrict__ out) { unsigned int gid = blockIdx.x * blockDim.x + threadIdx.x; unsigned int lid = threadIdx.x; @@ -44,7 +50,8 @@ out[gid] = sum.x + sum.y; } -extern "C" __global__ void global_bandwidth_v4(const float4 *in, float *out) +extern "C" __global__ void global_bandwidth_v4(const float4 * __restrict__ in, + float * __restrict__ out) { unsigned int gid = blockIdx.x * blockDim.x + threadIdx.x; unsigned int lid = threadIdx.x; diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/clpeak-2.0.18/src/opencl/kernels/compute_dp_kernels.cl new/clpeak-2.0.19/src/opencl/kernels/compute_dp_kernels.cl --- old/clpeak-2.0.18/src/opencl/kernels/compute_dp_kernels.cl 2026-07-11 18:02:40.000000000 +0200 +++ new/clpeak-2.0.19/src/opencl/kernels/compute_dp_kernels.cl 2026-08-15 09:13:05.000000000 +0200 @@ -10,11 +10,13 @@ \n #define DOUBLE_AVAILABLE \n#endif +// Plain contracted expression instead of mad() -- see compute_sp_kernels.cl + \n#undef MAD_4 \n#undef MAD_16 \n#undef MAD_64 \n -\n#define MAD_4(x, y) x = mad(y, x, y); y = mad(x, y, x); x = mad(y, x, y); y = mad(x, y, x); +\n#define MAD_4(x, y) x = (y*x) + y; y = (x*y) + x; x = (y*x) + y; y = (x*y) + x; \n#define MAD_16(x, y) MAD_4(x, y); MAD_4(x, y); MAD_4(x, y); MAD_4(x, y); \n#define MAD_64(x, y) MAD_16(x, y); MAD_16(x, y); MAD_16(x, y); MAD_16(x, y); \n diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/clpeak-2.0.18/src/opencl/kernels/compute_hp_kernels.cl new/clpeak-2.0.19/src/opencl/kernels/compute_hp_kernels.cl --- old/clpeak-2.0.18/src/opencl/kernels/compute_hp_kernels.cl 2026-07-11 18:02:40.000000000 +0200 +++ new/clpeak-2.0.19/src/opencl/kernels/compute_hp_kernels.cl 2026-08-15 09:13:05.000000000 +0200 @@ -7,11 +7,13 @@ \n #define HALF_AVAILABLE \n#endif +// Plain contracted expression instead of mad() -- see compute_sp_kernels.cl + \n#undef MAD_4 \n#undef MAD_16 \n#undef MAD_64 \n -\n#define MAD_4(x, y) x = mad(y, x, y); y = mad(x, y, x); x = mad(y, x, y); y = mad(x, y, x); +\n#define MAD_4(x, y) x = (y*x) + y; y = (x*y) + x; x = (y*x) + y; y = (x*y) + x; \n#define MAD_16(x, y) MAD_4(x, y); MAD_4(x, y); MAD_4(x, y); MAD_4(x, y); \n#define MAD_64(x, y) MAD_16(x, y); MAD_16(x, y); MAD_16(x, y); MAD_16(x, y); \n diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/clpeak-2.0.18/src/opencl/kernels/compute_mp_kernels.cl new/clpeak-2.0.19/src/opencl/kernels/compute_mp_kernels.cl --- old/clpeak-2.0.18/src/opencl/kernels/compute_mp_kernels.cl 2026-07-11 18:02:40.000000000 +0200 +++ new/clpeak-2.0.19/src/opencl/kernels/compute_mp_kernels.cl 2026-08-15 09:13:05.000000000 +0200 @@ -4,12 +4,15 @@ // Dominant arithmetic path in LLM training/prefill -- distinct from // compute_hp (fp16 accumulator). // -// Each MAD_4 issues 4 mixed-precision fma's: -// a = fma(convert_floatN(x), convert_floatN(y), a); +// Each MAD_4 issues 4 mixed-precision multiply-accumulates: +// a = convert_floatN(x) * convert_floatN(y) + a; // and writes a back into x (resp. y) via an fp16 downcast, so the // compiler can't CSE (float)x*(float)y across iterations. The downcast // is ~1 cycle on every vendor we care about and does not distort the // FMA measurement meaningfully. +// +// The MAC is a plain contracted expression rather than fma() for the same +// reason compute_sp_kernels.cl avoids mad() -- see the comment there. \n#if defined(cl_khr_fp16) \n #pragma OPENCL EXTENSION cl_khr_fp16 : enable @@ -20,10 +23,10 @@ \n#undef MAD_16 \n \n#define MAD_4(HT, FT, x, y, a) \ - a = fma(convert_##FT(x), convert_##FT(y), a); x = convert_##HT(a); \ - a = fma(convert_##FT(y), convert_##FT(x), a); y = convert_##HT(a); \ - a = fma(convert_##FT(x), convert_##FT(y), a); x = convert_##HT(a); \ - a = fma(convert_##FT(y), convert_##FT(x), a); y = convert_##HT(a); + a = convert_##FT(x) * convert_##FT(y) + a; x = convert_##HT(a); \ + a = convert_##FT(y) * convert_##FT(x) + a; y = convert_##HT(a); \ + a = convert_##FT(x) * convert_##FT(y) + a; x = convert_##HT(a); \ + a = convert_##FT(y) * convert_##FT(x) + a; y = convert_##HT(a); \n#define MAD_16(HT, FT, x, y, a) MAD_4(HT,FT,x,y,a); MAD_4(HT,FT,x,y,a); MAD_4(HT,FT,x,y,a); MAD_4(HT,FT,x,y,a); \n diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/clpeak-2.0.18/src/opencl/kernels/compute_sp_kernels.cl new/clpeak-2.0.19/src/opencl/kernels/compute_sp_kernels.cl --- old/clpeak-2.0.18/src/opencl/kernels/compute_sp_kernels.cl 2026-07-11 18:02:40.000000000 +0200 +++ new/clpeak-2.0.19/src/opencl/kernels/compute_sp_kernels.cl 2026-08-15 09:13:05.000000000 +0200 @@ -2,11 +2,20 @@ // Avoiding auto-vectorize by using vector-width locked dependent code +// The MAD is written as a plain contracted expression rather than the mad() +// builtin. Both request the same thing under the -cl-mad-enable we always +// build with (mad() carries no accuracy guarantee either -- the spec permits +// it to be a multiply followed by an add), but some frontends treat the +// builtin as a slow precise path and never lower it to the hardware FMA: +// Apple's deprecated CL-on-Metal compiler halves peak fp32 that way. This +// also matches the integer/char/short kernels here and the explicit fma +// intrinsics every other clpeak backend uses. Do not "restore" mad(). + \n#undef MAD_4 \n#undef MAD_16 \n#undef MAD_64 \n -\n#define MAD_4(x, y) x = mad(y, x, y); y = mad(x, y, x); x = mad(y, x, y); y = mad(x, y, x); +\n#define MAD_4(x, y) x = (y*x) + y; y = (x*y) + x; x = (y*x) + y; y = (x*y) + x; \n#define MAD_16(x, y) MAD_4(x, y); MAD_4(x, y); MAD_4(x, y); MAD_4(x, y); \n#define MAD_64(x, y) MAD_16(x, y); MAD_16(x, y); MAD_16(x, y); MAD_16(x, y); \n diff -urN '--exclude=CVS' '--exclude=.cvsignore' '--exclude=.svn' '--exclude=.svnignore' old/clpeak-2.0.18/src/rocm/rocm_kernels/global_bandwidth.hip new/clpeak-2.0.19/src/rocm/rocm_kernels/global_bandwidth.hip --- old/clpeak-2.0.18/src/rocm/rocm_kernels/global_bandwidth.hip 2026-07-11 18:02:40.000000000 +0200 +++ new/clpeak-2.0.19/src/rocm/rocm_kernels/global_bandwidth.hip 2026-08-15 09:13:05.000000000 +0200 @@ -1,8 +1,14 @@ #include <hip/hip_runtime.h> +// `in`/`out` are __restrict__ so the compiler can prove they do not alias and +// issue the FETCH_PER_WI loads as one batch ahead of the trailing store; +// without it the store forces a conservative, serialized load schedule. +// Mirrors cuda_kernels/global_bandwidth.cu. + #define FETCH_PER_WI 16 -extern "C" __global__ void global_bandwidth_v1(const float *in, float *out) +extern "C" __global__ void global_bandwidth_v1(const float * __restrict__ in, + float * __restrict__ out) { unsigned int gid = blockIdx.x * blockDim.x + threadIdx.x; unsigned int lid = threadIdx.x; @@ -20,7 +26,8 @@ out[gid] = sum; } -extern "C" __global__ void global_bandwidth_v2(const float2 *in, float *out) +extern "C" __global__ void global_bandwidth_v2(const float2 * __restrict__ in, + float * __restrict__ out) { unsigned int gid = blockIdx.x * blockDim.x + threadIdx.x; unsigned int lid = threadIdx.x; @@ -39,7 +46,8 @@ out[gid] = sum.x + sum.y; } -extern "C" __global__ void global_bandwidth_v4(const float4 *in, float *out) +extern "C" __global__ void global_bandwidth_v4(const float4 * __restrict__ in, + float * __restrict__ out) { unsigned int gid = blockIdx.x * blockDim.x + threadIdx.x; unsigned int lid = threadIdx.x;
