Gen12LP: post-RA spill/fill propagation corrupts OWord spills

Aperta
  • Qualcosa non va?
#471 0 commenti 0 reazioni 0 assegnatari Vedi su GitHub

@maleadt ci sta già lavorando.

Dal 7/11/2026.

  • #473 di @maleadt — aperta
  • #474 di @maleadt — aperta

Valutazione

Difficoltà
4/5
Tempo stimato
3-5 giorni
Idoneità per principianti
58/100
Tipo di issue
Bug
Chiarezza
Specificata chiaramente
Stato di attività
Attiva
Stack tecnologico
cpp
Ambito
compilers

Direzione di ricerca

Start with visa/SpillFillPropagation.cpp, especially SpillFillPropagation and getMaxSpillAreaOffset, then trace the caller that updates spill size. Reproduce with the supplied OpenCL kernel on Gen12LP and compare pass dumps with -postRAspillCleanup 0. Done means invalid-offset spill/fill intrinsics no longer merge distinct slots or reduce spill size, while valid-offset propagation still works; also account for OWord push/pop from BoundedRA.

Scritto dal modello di indicizzazione a partire dal testo della issue.

Descrizione

On platforms without LSC (Gen12LP and older), post-RA spill/fill propagation treats distinct OWord spill slots as the same slot, and the kernel silently returns wrong results.

// Each work-item loads 8 uint16 vectors and keeps them live until the stores
// (`out` may alias `in`), so at SIMD32 some of them get spilled. The private
// array `p` is dynamically indexed, so it lives in scratch: 640 * 4 B * 32 lanes
// = 80 KB per thread, which makes vISA spill with OWord block messages.
#ifndef P
#define P 640
#endif

#define R8(X) X(0) X(1) X(2) X(3) X(4) X(5) X(6) X(7)
#define LOAD(a) uint16 v##a = vload16(a * S + g, in);
#define STORE(a) vstore16(v##a + p[(a * s) % P], (7 - a) * S + g, out);

__attribute__((intel_reqd_sub_group_size(32)))
__kernel void k(__global uint *out, __global const uint *in, uint s) {
    uint g = get_global_id(0), S = get_global_size(0);
    R8(LOAD)
    uint p[P];
    for (uint j = 0; j < P; j++)
        p[(j * s) % P] = j;
    R8(STORE)
}

With 256 work-items in groups of 32 and s = 3 (host program below), on a Tiger Lake Iris Xe with IGC 2.41.5:

device: Intel(R) Iris(R) Xe Graphics, driver 26.35.039758
mismatch: vector 0, item 0, lane 1: got 0x00000000, want 0x9e3779b1
...
FAIL: 29695 of 32768 values wrong

With -DP=613 (78464 B of private memory, just under the threshold below) it passes. The kernel as given passes on an Arc A770 (DG2).

Root cause

When private memory takes at least 60% of the 128 KB that HWord scratch messages can address, GlobalRA::coloringRegAlloc spills with OWord block messages instead. Their offset goes in the message header, so their spill/fill intrinsics carry InvalidOffset.

SpillFillPropagation, added in v2.33.0 and enabled by default through vISA_GRFPostRASpillCodeCleanup, identifies spill slots by the intrinsic offset. These OWord intrinsics all carry InvalidOffset, so the pass treats distinct slots as one. In this reproducer, -dumpPassesAll shows 116 spills and 232 fills after Spill_GRF, but 1 spill and 4 fills after spillFillPropagation, with 228 more movs. getMaxSpillAreaOffset ignores invalid offsets and returns 0, so the caller also reduces the spill size to 0.

With IGC_VISAOptions="-postRAspillCleanup 0", the reproducer passes.

The same bug is on master (0ac43c0cf). DG2 uses LSC instead of OWord spills. The HWord path on Gen12LP has a separate problem, #469.

A fix could skip the propagation, and keep the spill size, whenever a spill/fill intrinsic has no valid offset. That also covers the OWord push/pop that BoundedRA generates.

We hit this in oneAPI.jl (Julia) with a kernel that has a large private array.

host.c

Build with cc host.c -lOpenCL -o host, then run ./host kernel.cl.

#define CL_TARGET_OPENCL_VERSION 300
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>

#define CHECK(x) do { cl_int e_ = (x); if (e_ != CL_SUCCESS) { \
    fprintf(stderr, "%s failed: %d\n", #x, e_); exit(1); } } while (0)

int main(int argc, char **argv) {
    const char *path = argc > 1 ? argv[1] : "kernel.cl";
    const char *opts = argc > 2 ? argv[2] : "";
    FILE *f = fopen(path, "rb");
    if (!f) { perror(path); return 1; }
    fseek(f, 0, SEEK_END); long len = ftell(f); fseek(f, 0, SEEK_SET);
    char *src = malloc(len + 1); fread(src, 1, len, f); src[len] = 0; fclose(f);

    cl_platform_id plat; cl_device_id dev; cl_int err;
    CHECK(clGetPlatformIDs(1, &plat, NULL));
    CHECK(clGetDeviceIDs(plat, CL_DEVICE_TYPE_GPU, 1, &dev, NULL));
    char name[256]; clGetDeviceInfo(dev, CL_DEVICE_NAME, sizeof name, name, NULL);
    char drv[256]; clGetDeviceInfo(dev, CL_DRIVER_VERSION, sizeof drv, drv, NULL);
    printf("device: %s, driver %s\n", name, drv);

    cl_context ctx = clCreateContext(NULL, 1, &dev, NULL, NULL, &err); CHECK(err);
    cl_command_queue q = clCreateCommandQueueWithProperties(ctx, dev, NULL, &err); CHECK(err);
    cl_program prog = clCreateProgramWithSource(ctx, 1, (const char **)&src, NULL, &err); CHECK(err);
    if (clBuildProgram(prog, 1, &dev, opts, NULL, NULL) != CL_SUCCESS) {
        char log[16384]; clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, sizeof log, log, NULL);
        fprintf(stderr, "build failed:\n%s\n", log); return 1;
    }
    cl_kernel k = clCreateKernel(prog, "k", &err); CHECK(err);

    enum { G = 256, V = 8, W = 16 };
    size_t n = (size_t)V * G * W;
    cl_uint *in = malloc(n * sizeof *in), *out = calloc(n, sizeof *out);
    for (size_t i = 0; i < n; i++) in[i] = (cl_uint)(i * 2654435761u);
    cl_mem bin = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, n * 4, in, &err); CHECK(err);
    cl_mem bout = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, n * 4, out, &err); CHECK(err);
    cl_uint s = 3;
    CHECK(clSetKernelArg(k, 0, sizeof bout, &bout));
    CHECK(clSetKernelArg(k, 1, sizeof bin, &bin));
    CHECK(clSetKernelArg(k, 2, sizeof s, &s));
    size_t gsz = G, lsz = 32;
    CHECK(clEnqueueNDRangeKernel(q, k, 1, NULL, &gsz, &lsz, 0, NULL, NULL));
    CHECK(clEnqueueReadBuffer(q, bout, CL_TRUE, 0, n * 4, out, 0, NULL, NULL));

    // out[7 - v] should be in[v] + p[v * s % P] = in[v] + v
    size_t bad = 0;
    for (size_t v = 0; v < V; v++)
        for (size_t g = 0; g < G; g++)
            for (size_t c = 0; c < W; c++) {
                cl_uint got = out[((V - 1 - v) * G + g) * W + c];
                cl_uint want = in[(v * G + g) * W + c] + (cl_uint)v;
                if (got != want && bad++ < 5)
                    printf("mismatch: vector %zu, item %zu, lane %zu: got 0x%08x, want 0x%08x\n",
                           v, g, c, got, want);
            }
    printf("%s: %zu of %zu values wrong\n", bad ? "FAIL" : "PASS", bad, n);
    return bad != 0;
}
Lingua principale
C++
Stelle
725
Fork
200
Nessuna PR unita negli ultimi 30g

Preparare l'ambiente

Come iniziare

  1. Leggi tutta la issue e poi la guida ai contributi del progetto.
  2. Commenta sulla issue per dire che te ne occupi tu — evita che due persone facciano lo stesso lavoro.
  3. Fai un fork del repository e lavora su un branch.
  4. Apri una pull request che faccia riferimento al numero della issue.

Altre issue di intel/intel-graphics-compiler

Tutte le issue di intel/intel-graphics-compiler

Issue simili

Altre issue su C++

Ricevi le nuove issue nella tua casella

Un breve riepilogo di issue GitHub adatte ai principianti.