[GSD-13401] OpenCL JIT causes deterministic host heap corruption on Arc B580 after clBuildProgram

[SINF-KEN]

Pre-submission Checklist

• I am using the latest GPU driver version (releases)
• I have searched for similar issues and found none

GPU Hardware

Intel Arc B580 (Battlemage G21, ASRock), PCI 0000:03:00.0

DRI Devices Information

$ ls -ls /dev/dri/*

0 crw-rw----+ 1 root video  226,   0  5 sep 14:51 /dev/dri/card0
0 crw-rw----+ 1 root video  226,   1  5 sep 13:58 /dev/dri/card1
0 crw-rw-rw-  1 root render 226, 128  4 sep 13:03 /dev/dri/renderD128
0 crw-rw-rw-  1 root render 226, 129  4 sep 13:03 /dev/dri/renderD129
 

$ ls -la /dev/dri/by-path/

total 0
0 lrwxrwxrwx 1 root root  8  4 sep 13:03 pci-0000:03:00.0-card -> ../card0
0 lrwxrwxrwx 1 root root 13  4 sep 13:03 pci-0000:03:00.0-render -> ../renderD128
0 lrwxrwxrwx 1 root root  8  4 sep 13:03 pci-0000:0e:00.0-card -> ../card1
0 lrwxrwxrwx 1 root root 13  4 sep 13:03 pci-0000:0e:00.0-render -> ../renderD129

GPU Detailed Information (lspci output)

$ lspci -vvv -k -s 0000:03:00.0
03:00.0 VGA compatible controller: Intel Corporation Battlemage G21 [Arc B580] (prog-if 00 [VGA controller])
	Subsystem: ASRock Incorporation Device 6021
	Control: I/O- Mem+ BusMaster+ SpecCycle- MemWINV- VGASnoop- ParErr- Stepping- SERR- FastB2B- DisINTx-
	Status: Cap+ 66MHz- UDF- FastB2B- ParErr- DEVSEL=fast >TAbort- <TAbort- <MAbort- >SERR- <PERR- INTx-
	Latency: 0, Cache Line Size: 64 bytes
	Interrupts: unknown pin routed to IRQ 89, MSI(X) routed to IRQ 89
	IOMMU group: 15
	Region 0: Memory at f4000000 (64-bit, non-prefetchable) [size=16M]
	Region 2: Memory at f800000000 (64-bit, prefetchable) [size=16G]
	Expansion ROM at f5000000 [disabled] [size=2M]
	Capabilities: <access denied>
	Kernel driver in use: xe
	Kernel modules: xe

Driver Version

26.31.39395 (package intel-compute-runtime 26.31.39395.13-1.1; clinfo: Driver Version 26.31.39395, OpenCL 3.0 NEO)

Installed GPU Driver Packages

$ pacman -Q | grep -iE "igc|gmm|opencl|level-zero|ocloc|compute|clinfo|libze"
clinfo 3.0.25.02.14-1.1
intel-compute-runtime 26.31.39395.13-1.1
intel-graphics-compiler 1:2.40.13-1.1
intel-gmmlib 22.10.0-1.1
level-zero-headers 1.32.0-1
level-zero-loader 1.32.0-1
ocl-icd 2.3.5-1.1
opencl-headers 2:2026.05.29-1
opencl-mesa 3:26.2.2-2

Driver Installation Details

• Installation method: Distribution repository (CachyOS, Arch-based)
• sudo pacman -S intel-compute-runtime intel-graphics-compiler (pulled
  in intel-gmmlib; ocl-icd as the ICD loader)
• No custom kernel parameters, no out-of-tree driver builds, no PPA/Intel repo

Linux Distribution

Other (please specify below)

Other Linux Distribution

CachyOS (Arch-based, rolling), fully updated as of 2026-09-05

Kernel Version & Boot Parameters

$ uname -r
7.2.3-1-cachyos
 
$ cat /proc/cmdline
quiet nowatchdog splash rw rootflags=subvol=/@ root=UUID=939efe19-f8c2-4c1a-8425-33008d10c447
 
$ grep -E '^(xe|i915) ' /proc/modules
xe 4677632 49 - Live 0x0000000000000000

Actual Behavior

Actual Behavior (*)

A program that builds a 10-kernel OpenCL program in-process
(clCreateProgramWithSource + clBuildProgram) and then executes ONE
trivial kernel aborts in the first kernel-run block, at the read-back /
clReleaseMemObject / clReleaseKernel calls, or the first heap operation
right after them. with glibc heap-corruption errors:

munmap_chunk(): invalid pointer (SIGABRT, exit 134)

Every OpenCL call before the abort returns CL_SUCCESS and the kernel's
read-back values are 100% correct; this is host heap corruption, not
wrong code. Under AddressSanitizer the same block is a SEGV
("unknown-crash"): an 8-byte READ of a stack local whose shadow bytes
are marked addressable, i.e. process state was corrupted earlier by
uninstrumented driver code, most plausibly during the build path.

Fault is on the Intel compute-runtime path: with the Mesa rusticl ICD
excluded from the loader (OCL_ICD_VENDORS=/etc/OpenCL/vendors/intel.icd)
the reproducer still aborts 5/5.

The corruption is heap-layout sensitive: minimized variants of the same
program run clean (0/11) ; the FULL attached source must be used
verbatim; a sibling program building the same kernel shapes from a
different host allocation pattern runs clean (0/10); removing the u64
div/rem kernels keeps the 100% abort rate (5/5); a 2-kernel program with
the same host structure is clean over dozens of runs.

Expected Behavior

The program should run the full builtin sweep, print
"ALL 64-BIT BUILTINS CORRECT (0 mismatches)", release all objects, and
exit 0. No OpenCL error is returned at any step before the abort (every
call returns CL_SUCCESS; the blocking read-back completes with correct
data).

Reproduction Rate

Always reproduces - 100%

Steps to Reproduce

1. Save the ATTACHED builtin64.cpp verbatim (exact source required — the
   bug is heap-layout sensitive and minimized variants do NOT reproduce)
2. g++ -O2 -g -std=c++17 builtin64.cpp -o builtin64 -lOpenCL
3. ./builtin64
   → prints "build OK", runs the first kernel and reads it back
   successfully (values verified correct), then aborts with
   "munmap_chunk(): invalid pointer" (SIGABRT, exit 134)
4. Re-running the same binary aborts every time
5. Rebuild with -fsanitize=address: SIGSEGV "unknown-crash" in the same
   block (8-byte read of an addressable stack local)

Observed output (step 3):

build OK
S1 start
S1 buf ok
  run2 k_clz: create kernel
  run2 k_clz: create ob
  run2 k_clz: set args
  run2 k_clz: enqueue
  run2 k_clz: read
  run2 k_clz: release
  run2 k_clz: done
munmap_chunk(): invalid pointer

builtin64.cpp

Is this a regression?

• Yes, this is a regression - functionality that previously worked is now broken

Last Known Working Driver Version

No response

First Known Failing Driver Version

No response

API Call Logs

No response

strace Logs

$ gdb ./builtin64
(gdb) run
...
Thread 1 "builtin64" received signal SIGABRT, Aborted.
#0  pthread_kill () from /usr/lib/libc.so.6
#1  raise () from /usr/lib/libc.so.6
#2  abort () from /usr/lib/libc.so.6
#3  malloc_printerr () from /usr/lib/libc.so.6     # "munmap_chunk(): invalid pointer"
#4-#6 glibc free/munmap_chunk internals
#7+ frames above are corrupted garbage — consistent with the heap
    corruption glibc is reporting

All driver worker threads idle at abort time, e.g.:
Thread 5 "b:clctxworker0" ... pthread_cond_wait () from libigdrcl.so
(no NEO thread is executing during the abort — the corruption happened
earlier, most plausibly during clBuildProgram)

Under -fsanitize=address (separate build):
==41796==ERROR: AddressSanitizer: unknown-crash ... READ of size 8 ...
Address is located in stack of thread T0 ... 'ob' ...
shadow byte 00 (addressable) — i.e. a valid-looking stack object read
through a corrupted pointer/path.

System Logs / dmesg Output

❯ sudo dmesg | grep -iE 'xe|drm'
[21293.706580] amdgpu 0000:0e:00.0: [drm] PCIE GART of 1024M enabled (table at 0x000000F41FC00000).
[21293.800100] amdgpu 0000:0e:00.0: [drm] DMUB hardware initialized: version=0x05003300
[21293.805852] amdgpu 0000:0e:00.0: [drm] Cannot find any crtc or sizes
[21293.805861] amdgpu 0000:0e:00.0: [drm] Cannot find any crtc or sizes
[21297.732055] mei_lb xe.mei-gscfi.768-e2c2afa2-3817-4d19-9d95-06b16b588a5d: bound 0000:03:00.0 (ops xe_late_bind_component_ops [xe])
[22944.761050] amdgpu 0000:0e:00.0: [drm] PCIE GART of 1024M enabled (table at 0x000000F41FC00000).
[22944.853488] amdgpu 0000:0e:00.0: [drm] DMUB hardware initialized: version=0x05003300
[22944.858105] amdgpu 0000:0e:00.0: [drm] Cannot find any crtc or sizes
[22944.858116] amdgpu 0000:0e:00.0: [drm] Cannot find any crtc or sizes
[22948.783747] mei_lb xe.mei-gscfi.768-e2c2afa2-3817-4d19-9d95-06b16b588a5d: bound 0000:03:00.0 (ops xe_late_bind_component_ops [xe])
[57172.131426] xe 0000:03:00.0: [drm] Tile0: GT0: log buffer overflow
[57172.131470] xe 0000:03:00.0: [drm] Tile0: GT0: 
[57172.131475] xe 0000:03:00.0: [drm] Tile0: GT0: Fault response: Unsuccessful -ENOENT
[57716.597913] xe 0000:03:00.0: [drm] Tile0: GT0: 
[57716.597920] xe 0000:03:00.0: [drm] Tile0: GT0: Fault response: Unsuccessful -ENOENT

Backtrace (if crash or hang occurred)

$ gdb ./builtin64
(gdb) run
...
Thread 1 "builtin64" received signal SIGABRT, Aborted.
#0  pthread_kill () from /usr/lib/libc.so.6
#1  raise () from /usr/lib/libc.so.6
#2  abort () from /usr/lib/libc.so.6
#3  malloc_printerr () from /usr/lib/libc.so.6     # "munmap_chunk(): invalid pointer"
#4-#6 glibc free/munmap_chunk internals
#7+ frames above are corrupted garbage — consistent with the heap
    corruption glibc is reporting

All driver worker threads idle at abort time, e.g.:
Thread 5 "b:clctxworker0" ... pthread_cond_wait () from libigdrcl.so
(no NEO thread is executing during the abort — the corruption happened
earlier, most plausibly during clBuildProgram)

Under -fsanitize=address (separate build):
==41796==ERROR: AddressSanitizer: unknown-crash ... READ of size 8 ...
Address is located in stack of thread T0 ... 'ob' ...
shadow byte 00 (addressable) — i.e. a valid-looking stack object read
through a corrupted pointer/path.

Source Code / Reproducer

• Attached: builtin64.cpp — FULL source, use verbatim (10 kernels:
  clz/ctz/popcount/mul_hi/mad_hi/u64 div/rem/shl/shr/sshr; runs k_clz
  first; aborts deterministically). Minimized variants do NOT reproduce.
• Compilation command: g++ -O2 -g -std=c++17 builtin64.cpp -o builtin64 -lOpenCL

Command Line / Application Details

./builtin64

(no options; GPU idle during all runs, nothing else running)

oneAPI Version (if applicable)

NA

Screenshots / Video

NA

Additional Notes

NA

[SINF-KEN]

ps i love my B580 its an amazing tool
thanks intel <3

[kgibala]

Hi @SINF-KEN,

Thank you for the submission - this is a genuinely well-prepared report. The attached reproducer, the gdb and Valgrind output, the dmesg excerpt, the ICD-isolation experiment and the note that minimized variants behave differently all gave us a lot to work with.

We built and ran builtin64.cpp locally on Battlemage hardware on our side. The crash reproduces for us as well. However, after tracing it we believe the failure comes from the reproducer itself rather than from the driver, and we would like to walk you through what we found so you can check it independently.

---

The main thing to look at:

run2 (line 80) and run3 (line 97) are lambdas with no declared return type, and your CK macro (line 43) expands to a return 1 statement. Because those return statements exist, the compiler deduces the lambdas return type as int. On the success path no CK fires, so control reaches the closing brace at line 96 (and 107) without returning a value. In C++ that is undefined behaviour, even when the caller discards the result - which it does, at line 117.

The quickest way to confirm this on your side is to add one flag: -Wreturn-type 

...
builtin64.cpp: In lambda function:
builtin64.cpp:96:5: error: control reaches end of non-void function [-Werror=return-type]
   96 |     };
      |     ^
builtin64.cpp: In lambda function:
builtin64.cpp:107:5: error: control reaches end of non-void function [-Werror=return-type]
  107 |     };
      |     ^
cc1plus: some warnings being treated as errors

That should stop the build and point at lines 96 and 107 directly. -Wreturn-type is on by default in C++, so if you scroll back through your original build output you will probably find the warning was already there.

---

Other things worth a look while you are in there:

Once the crash was out of the way, the sweep ran to the end and reported mismatches. Those turned out to trace back to the reproducer too, so these are worth checking before reading anything into them:

1. k_mad_hi takes four arguments -a, b, c, out (line 23) - but run3 only sets three (lines 100-102). Argument 3 is never set, so the enqueue fails with CL_INVALID_KERNEL_ARGS (-52). Your comment at line 136 suggests you intended to bind c = b, which run3 has no slot for.
2. A failed run is not detected. run3 returns an error code that the call sites discard, and out.resize(n) is a no-op when the size is unchanged. So after the failure above, the verification loop compares the previous kernel's output. The four mad_hi mismatches we first saw were bit-identical to the preceding k_mul_hi results.
3. k_sshr reads out of bounds. It is dispatched over N3 = M * S = 960 work-items (line 186), but argument 0 is built from a 15-element vector (lines 178-180), so it is a 120-byte buffer being read as up to 7680 bytes. Worth comparing against k_shl and k_shr just above it, which use the same 960-wide dispatch with correctly sized cross-product buffers and pass cleanly. This may also be where the Fault response: Unsuccessful -ENOENT lines in your dmesg come from.
4. ref_mad_hi (lines 46-48) computes a different formula than mad_hi. It evaluates high64(a*b + c), adding c to the full 128-bit product before truncating. OpenCL C defines mad_hi(a, b, c) as mul_hi(a, b) + c, where c is added to the already-truncated high half. We checked all 15 edge values and the GPU matches the spec definition exactly on every one, while the reference matches on none of them.
5. Minor: lines 82, 84, 98, 99, 109 pass NULL for errcode_ret and never check the returned handles, so a genuine driver error there would surface as a null-handle crash rather than a reported error code. Also line 132 creates two buffers from the same host_ptr, which the OpenCL spec leaves undefined.

---

With those addressed, and with no change to the driver, IGC or kernel driver, the sweep passes on our side:

g++ -O2 -g -std=c++17 -Werror=return-type builtin64.cpp -o builtin64 -lOpenCL

./builtin64
build OK
S1 start
S1 buf ok
clz/ctz/popcount done (0 fails so far)
mul_hi/mad_hi done (0 fails so far)
div/rem done (0 fails so far)
shifts done (0 fails total)

ALL 64-BIT BUILTINS CORRECT (0 mismatches)

All ten builtins come out correct on Battlemage.

We are attaching our source code as so you can compare - every change is commented and none of them touch the kernels themselves.

Click to view source code

// 64-bit builtin correctness sweep on B580 vs CPU reference.
// Ops: clz, ctz, popcount, mul_hi, mad_hi, udiv, urem, shl, shr (u64+long).
// Edge inputs; shift amounts kept < 64 (>=64 is UB in OpenCL C).

#define CL_TARGET_OPENCL_VERSION 300
#include <CL/cl.h>
#include <cstdio>
#include <cstring>
#include <cstdint>
#include <vector>
#include <initializer_list>

static const char* SRC = R"CLC(
__kernel void k_clz(__global const ulong* in, __global ulong* out){
    ulong x = in[get_global_id(0)]; if (x) out[get_global_id(0)] = clz(x); else out[get_global_id(0)] = 0xDEAD;
}
__kernel void k_ctz(__global const ulong* in, __global ulong* out){
    ulong x = in[get_global_id(0)]; if (x) out[get_global_id(0)] = ctz(x); else out[get_global_id(0)] = 0xDEAD;
}
__kernel void k_popcount(__global const ulong* in, __global ulong* out){
    out[get_global_id(0)] = popcount(in[get_global_id(0)]);
}
__kernel void k_mul_hi(__global const ulong* a, __global const ulong* b, __global ulong* out){
    out[get_global_id(0)] = mul_hi(a[get_global_id(0)], b[get_global_id(0)]);
}
__kernel void k_mad_hi(__global const ulong* a, __global const ulong* b, __global const ulong* c, __global ulong* out){
    out[get_global_id(0)] = mad_hi(a[get_global_id(0)], b[get_global_id(0)], c[get_global_id(0)]);
}
__kernel void k_div(__global const ulong* a, __global const ulong* b, __global ulong* out){
    if (b[get_global_id(0)]) out[get_global_id(0)] = a[get_global_id(0)] / b[get_global_id(0)]; else out[get_global_id(0)] = 0xDEAD;
}
__kernel void k_rem(__global const ulong* a, __global const ulong* b, __global ulong* out){
    if (b[get_global_id(0)]) out[get_global_id(0)] = a[get_global_id(0)] % b[get_global_id(0)]; else out[get_global_id(0)] = 0xDEAD;
}
__kernel void k_shl(__global const ulong* a, __global const ulong* s, __global ulong* out){
    out[get_global_id(0)] = a[get_global_id(0)] << (s[get_global_id(0)] & 63);
}
__kernel void k_shr(__global const ulong* a, __global const ulong* s, __global ulong* out){
    out[get_global_id(0)] = a[get_global_id(0)] >> (s[get_global_id(0)] & 63);
}
__kernel void k_sshr(__global const long* a, __global const ulong* s, __global long* out){
    out[get_global_id(0)] = a[get_global_id(0)] >> (s[get_global_id(0)] & 63);
}
)CLC";

// Injects "return 1;" - so EVERY function using it must return int explicitly.
#define CK(x) do { cl_int e_=(x); if(e_!=CL_SUCCESS){ printf("OPENCL ERROR %d at %s\n",e_,#x); return 1;} } while(0)

// FIX: never verify results after a failed run.
#define RUN(call) do { if ((call) != 0) { printf("ABORT: %s failed\n", #call); return 1; } } while(0)

// OpenCL C, Integer Functions:
//   mul_hi(x, y) - "Computes x * y and returns the high half of the product of x and y."
//   mad_hi(a, b, c) - "Returns mul_hi(a, b) + c."
static inline uint64_t ref_mul_hi(uint64_t a, uint64_t b){ return (uint64_t)(((__uint128_t)a * b) >> 64); }

// FIX: the original computed high64(a*b + c), i.e. it added c to the FULL
// 128-bit product before taking the high half. That is NOT mad_hi - it is the
// high half of a fused 128-bit multiply-add, and it disagrees with mad_hi for
// any input where c perturbs the low half enough to matter. mad_hi adds c to
// the ALREADY-TRUNCATED high half, with normal unsigned wraparound.
// The GPU matched the spec definition on all 15 edge values; the reference did
// not, which is what produced the 15 "mad_hi FAIL" lines.
static inline uint64_t ref_mad_hi(uint64_t a, uint64_t b, uint64_t c){
    return ref_mul_hi(a, b) + c;          // wraps mod 2^64, as the GPU does
}

int main() {
    setvbuf(stdout, nullptr, _IONBF, 0);
    std::vector<uint64_t> edge = {
        1, 2, 3, 0xFFFFFFFFull, 0x100000000ull, 0x7FFFFFFF00000001ull,
        0x8000000000000000ull, 0x8000000000000001ull, 0xFFFFFFFFFFFFFFFEull,
        0xFFFFFFFFFFFFFFFFull, 0x123456789ABCDEFull, 0xFEDCBA9876543210ull,
        0x5555555555555555ull, 0x00000000FFFFFFFFull, 0x0000000100000000ull
    };
    const size_t M = edge.size();                                   // 15
    std::vector<uint64_t> shifts; for (uint64_t i = 0; i < 64; i++) shifts.push_back(i);
    const size_t S = shifts.size();                                 // 64

    cl_platform_id plats[4]; cl_uint np = 0; CK(clGetPlatformIDs(4, plats, &np));
    cl_platform_id intel = nullptr;
    for (cl_uint i = 0; i < np; i++) {
        char nm[256] = {}; clGetPlatformInfo(plats[i], CL_PLATFORM_NAME, sizeof nm, nm, nullptr);
        if (strstr(nm, "Intel")) { intel = plats[i]; break; }
    }
    if (intel == nullptr) { printf("no Intel OpenCL platform found\n"); return 1; }

    cl_device_id dev; CK(clGetDeviceIDs(intel, CL_DEVICE_TYPE_GPU, 1, &dev, nullptr));
    cl_int err;
    cl_context ctx = clCreateContext(nullptr, 1, &dev, nullptr, nullptr, &err); CK(err);
    // Successor to the deprecated clCreateCommandQueue; nullptr properties == defaults.
    cl_command_queue q = clCreateCommandQueueWithProperties(ctx, dev, nullptr, &err); CK(err);

    const char* s = SRC;
    cl_program prog = clCreateProgramWithSource(ctx, 1, &s, nullptr, &err); CK(err);
    err = clBuildProgram(prog, 1, &dev, nullptr, nullptr, nullptr);
    if (err != CL_SUCCESS) { char log[16384] = {};
        clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, sizeof log, log, nullptr);
        printf("BUILD FAILED %s\n", log); return 1; }
    printf("build OK\n");

    int fails = 0;

    // -------------------------------------------------------------------------
    // FIX: one arity-safe runner replaces run2 and run3.
    //  - explicit "-> int" and an explicit "return 0;"  => no fall-through UB
    //  - inputs bound to args 0..n-1, output buffer bound LAST => any arity
    //  - out.assign(n, 0) rather than resize()          => no stale data
    //  - every clCreate* checked
    // -------------------------------------------------------------------------
    auto runN = [&](const char* kn, std::initializer_list<cl_mem> ins,
                    size_t n, std::vector<uint64_t>& out) -> int {
        cl_int e = CL_SUCCESS;
        cl_kernel k = clCreateKernel(prog, kn, &e);                               CK(e);
        cl_mem   ob = clCreateBuffer(ctx, CL_MEM_READ_WRITE, n * 8, nullptr, &e);  CK(e);

        cl_uint idx = 0;
        for (cl_mem m : ins) CK(clSetKernelArg(k, idx++, sizeof(cl_mem), &m));
        CK(clSetKernelArg(k, idx, sizeof(cl_mem), &ob));          // output = last arg

        CK(clEnqueueNDRangeKernel(q, k, 1, nullptr, &n, nullptr, 0, nullptr, nullptr));

        out.assign(n, 0);
        CK(clEnqueueReadBuffer(q, ob, CL_TRUE, 0, n * 8, out.data(), 0, nullptr, nullptr));

        clReleaseMemObject(ob); clReleaseKernel(k);
        return 0;                                                 // <-- was missing
    };

    auto mkbuf = [&](const std::vector<uint64_t>& v) -> cl_mem {
        cl_int e = CL_SUCCESS;
        cl_mem m = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                  v.size() * 8, (void*)v.data(), &e);
        if (e != CL_SUCCESS || m == nullptr)
            printf("OPENCL ERROR %d at mkbuf(%zu elems)\n", e, v.size());
        return m;
    };

    // ---- clz / ctz / popcount (2-arg kernels: in, out) ----
    {
        printf("S1 start\n");
        cl_mem in = mkbuf(edge); if (in == nullptr) return 1;
        std::vector<uint64_t> out;
        printf("S1 buf ok\n");

        RUN(runN("k_clz", {in}, M, out));
        for (size_t i = 0; i < M; i++) { uint64_t want = edge[i] ? (uint64_t)__builtin_clzll(edge[i]) : 0xDEAD;
            if (out[i] != want) { fails++; printf("clz(0x%016llx) got %llu want %llu FAIL\n",(unsigned long long)edge[i],(unsigned long long)out[i],(unsigned long long)want);} }

        RUN(runN("k_ctz", {in}, M, out));
        for (size_t i = 0; i < M; i++) { uint64_t want = edge[i] ? (uint64_t)__builtin_ctzll(edge[i]) : 0xDEAD;
            if (out[i] != want) { fails++; printf("ctz(0x%016llx) got %llu want %llu FAIL\n",(unsigned long long)edge[i],(unsigned long long)out[i],(unsigned long long)want);} }

        RUN(runN("k_popcount", {in}, M, out));
        for (size_t i = 0; i < M; i++) { uint64_t want = (uint64_t)__builtin_popcountll(edge[i]);
            if (out[i] != want) { fails++; printf("popcount(0x%016llx) got %llu want %llu FAIL\n",(unsigned long long)edge[i],(unsigned long long)out[i],(unsigned long long)want);} }

        clReleaseMemObject(in);
        printf("clz/ctz/popcount done (%d fails so far)\n", fails);
    }

    // ---- mul_hi (3-arg) / mad_hi (4-arg) ----
    // FIX: ONE buffer bound to every input, instead of two buffers built
    // from the same host_ptr. a == b == c == edge is the original intent
    // (see the "mad_hi with c = b (reuse)" comment in the original).
    {
        cl_mem a = mkbuf(edge); if (a == nullptr) return 1;
        std::vector<uint64_t> out;

        RUN(runN("k_mul_hi", {a, a}, M, out));
        for (size_t i = 0; i < M; i++) { uint64_t want = ref_mul_hi(edge[i], edge[i]);
            if (out[i] != want) { fails++; printf("mul_hi(0x%016llx) got %016llx want %016llx FAIL\n",(unsigned long long)edge[i],(unsigned long long)out[i],(unsigned long long)want);} }

        // FIX: four inputs bound (a, b, c) + out - this is the dispatch that
        // previously failed with CL_INVALID_KERNEL_ARGS because arg 3 was unset.
        RUN(runN("k_mad_hi", {a, a, a}, M, out));
        for (size_t i = 0; i < M; i++) { uint64_t want = ref_mad_hi(edge[i], edge[i], edge[i]);
            if (out[i] != want) { fails++; printf("mad_hi(0x%016llx) got %016llx want %016llx FAIL\n",(unsigned long long)edge[i],(unsigned long long)out[i],(unsigned long long)want);} }

        clReleaseMemObject(a);
        printf("mul_hi/mad_hi done (%d fails so far)\n", fails);
    }

    // ---- div / rem (divisors include odd primes, powers of two, near/max) ----
    {
        std::vector<uint64_t> divs = {1,2,3,5,7,61,64,97,255,65537,0xFFFFFFFFull,0x100000000ull,
            0x7FFFFFFFFFFFFFFFull,0x8000000000000000ull,0xFFFFFFFFFFFFFFFFull, 0x987654321ull};
        std::vector<uint64_t> num;
        for (auto d : divs) { num.push_back(d - 1); num.push_back(d); num.push_back(d + 1); num.push_back(0xFFFFFFFFFFFFFFFFull / d * d + d - 1); }
        num.push_back(0x8000000000000000ull); num.push_back(0xFFFFFFFFFFFFFFFFull);
        std::vector<uint64_t> dens; for (size_t i = 0; i < num.size(); i++) dens.push_back(divs[i % divs.size()]);
        const size_t N2 = num.size();

        // num and dens are distinct vectors, so two buffers are legitimate here.
        cl_mem a = mkbuf(num), b = mkbuf(dens);
        if (a == nullptr || b == nullptr) return 1;
        std::vector<uint64_t> out;

        RUN(runN("k_div", {a, b}, N2, out));
        for (size_t i = 0; i < N2; i++) { uint64_t want = num[i] / dens[i];
            if (out[i] != want) { fails++; printf("div(0x%016llx / 0x%016llx) got %016llx want %016llx FAIL\n",(unsigned long long)num[i],(unsigned long long)dens[i],(unsigned long long)out[i],(unsigned long long)want);} }

        RUN(runN("k_rem", {a, b}, N2, out));
        for (size_t i = 0; i < N2; i++) { uint64_t want = num[i] % dens[i];
            if (out[i] != want) { fails++; printf("rem(0x%016llx %% 0x%016llx) got %016llx want %016llx FAIL\n",(unsigned long long)num[i],(unsigned long long)dens[i],(unsigned long long)out[i],(unsigned long long)want);} }

        clReleaseMemObject(a); clReleaseMemObject(b);
        printf("div/rem done (%d fails so far)\n", fails);
    }

    // ---- shifts: u64 shl/shr, long shr (arithmetic) ----
    {
        const size_t N3 = M * S;                                    // 15 * 64 = 960

        // Cross-product inputs: index i -> edge[i % M], shifts[i % S].
        std::vector<uint64_t> A, SH;
        A.reserve(N3); SH.reserve(N3);
        for (size_t i = 0; i < N3; i++) { A.push_back(edge[i % M]); SH.push_back(shifts[i % S]); }

        // Signed inputs for the arithmetic-shift kernel.
        std::vector<int64_t> SA; SA.reserve(M);
        for (auto x : edge) SA.push_back((int64_t)x);
        SA[0] = -1; SA[1] = -2;                                     // exercise negatives

        // FIX: build the FULL N3-element cross-product. The original passed
        // only M = 15 elements while dispatching N3 = 960 work-items, so
        // work-items 15..959 read out of bounds. Building it this way also
        // makes the host reference below (SA[i % M]) match the kernel's a[i].
        std::vector<uint64_t> SAv; SAv.reserve(N3);
        for (size_t i = 0; i < N3; i++) SAv.push_back((uint64_t)SA[i % M]);

        cl_mem bA = mkbuf(A), bSH = mkbuf(SH), bSA = mkbuf(SAv);
        if (bA == nullptr || bSH == nullptr || bSA == nullptr) return 1;

        std::vector<uint64_t> out;

        RUN(runN("k_shl", {bA, bSH}, N3, out));
        for (size_t i = 0; i < N3; i++) { uint64_t want = A[i] << SH[i];
            if (out[i] != want) { fails++; printf("shl(0x%016llx<<%llu) got %016llx want %016llx FAIL\n",(unsigned long long)A[i],(unsigned long long)SH[i],(unsigned long long)out[i],(unsigned long long)want);} }

        RUN(runN("k_shr", {bA, bSH}, N3, out));
        for (size_t i = 0; i < N3; i++) { uint64_t want = A[i] >> SH[i];
            if (out[i] != want) { fails++; printf("shr(0x%016llx>>%llu) got %016llx want %016llx FAIL\n",(unsigned long long)A[i],(unsigned long long)SH[i],(unsigned long long)out[i],(unsigned long long)want);} }

        // Arithmetic shift right on negative longs.
        // Kernel computes out[i] = a[i] >> (s[i] & 63) with a[i] == SA[i % M]
        // and s[i] == SH[i]; SH values are already < 64 so & 63 is identity.
        std::vector<uint64_t> souts;
        RUN(runN("k_sshr", {bSA, bSH}, N3, souts));
        for (size_t i = 0; i < N3; i++) {
            int64_t want = SA[i % M] >> (int)SH[i];
            int64_t got  = (int64_t)souts[i];
            if (got != want) { fails++; printf("sshr(%lld>>%llu) got %lld want %lld FAIL\n",(long long)SA[i % M],(unsigned long long)SH[i],(long long)got,(long long)want);} }

        clReleaseMemObject(bA); clReleaseMemObject(bSH); clReleaseMemObject(bSA);
        printf("shifts done (%d fails total)\n", fails);
    }

    printf("\n%s (%d mismatches)\n", fails ? "MISMATCHES FOUND - investigate before filing" : "ALL 64-BIT BUILTINS CORRECT", fails);

    // Original leaked these; releasing them clears the "still reachable"
    // Valgrind noise so a future Valgrind run is readable.
    clReleaseProgram(prog);
    clReleaseCommandQueue(q);
    clReleaseContext(ctx);

    return fails ? 1 : 0;
}

Could you re-run on your setup and let us know what you see?

Thanks again for the detailed write-up, and for the kind words about the B580.

[SINF-KEN]

I’m gonna check it out :)

Thanks again for taking the time to trace this down. Really appreciate the help!

For a bit of context, I’m building a Goldbach verifier for the B580 using Intel Level Zero, and I’m currently managing ~8e9 numbers/s in the 4e18 number range. :)

I’m definitely pushing the GPU pretty hard and have been running into a few bugs on my side haha. I wasn’t sure if this one was me or the driver 😆 I’m also fairly new to C++ and GPU programming, so I’m still learning a lot along the way.