Skip to content

2.0.19

Choose a tag to compare

@krrishnarraj krrishnarraj released this 16 Aug 05:06
· 1 commit to 2.0.x since this release

Fixes

  • OpenCL: fix ~2x fp32 throughput undercount on Apple. The compute
    kernels used the mad() builtin, which the OpenCL spec permits (but does
    not require) to fuse into hardware FMA. Apple's deprecated CL-on-Metal 1.2
    compiler never lowers it, silently halving measured fp32 throughput
    (2.2 → 4.5 TFLOPS on M1 Pro). Rewritten as a plain contracted expression
    ((y*x)+y), which forces the compiler to choose between fusing or
    violating -cl-mad-enable — all vendors honor this. OpenCL fp32 now
    matches Metal on Apple (within 1–2%); other platforms are unaffected
    (mad() already fused there).
  • CUDA/ROCm: mark global-bandwidth kernel pointers __restrict__. The
    in/out pointer pair had no aliasing annotation, so the compiler
    couldn't batch loads ahead of the trailing store and NVIDIA couldn't take
    the read-only ld.global.nc path. These were the only kernels in either
    backend with two pointer parameters, so the fix is contained to them.
  • CPU: fix GCC build failure on the crypto ISA extensions (#189).
    GCC declares AES/SHA-256 NEON intrinsics under a single umbrella
    +crypto target attribute, so a translation unit built with only the
    granular +aes or +sha2 flag failed to inline them ("target specific
    option mismatch"). GCC now gets +crypto (+crypto+sha3 for the SHA-512
    TU); Clang keeps the granular flags, since +crypto is deprecated there.

CI

  • Added a GCC-only build job to catch this class of compiler-specific
    regression going forward.

Full Changelog: 2.0.18...2.0.19