Gen12LP: OWord spills of indirectly accessed variables use the wrong slot

Aperta
  • Qualcosa non va?
#472 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
53/100
Tipo di issue
Bug
Chiarezza
Specificata chiaramente
Stato di attività
Attiva
Stack tecnologico
cpp
Ambito
compilers

Direzione di ricerca

Start with insertAddrTakenSpillAndFillCode and its LS variant, then inspect sendInSpilledRegVarPortions and sendOutSpilledRegVarPortions in visa/SpillManagerGMRF.cpp, especially how they derive OWord offsets. Compare those helpers with the HWord and LSC paths described in the issue. Reproduce with the supplied OpenCL kernel on Gen12LP; done means dynamic reads use the same scratch slot as direct spills and the kernel passes.

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

Descrizione

On platforms without LSC (Gen12LP and older), OWord fills for indirect register accesses (r[a0]) can read a different scratch slot from the one written by direct spills of the same variable, producing 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. Each vector is
// also read at a dynamic lane index, which vISA lowers to an indirect access.
// The private array `p` takes 80 KB per thread, so spills use OWord 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); \
                 out[16 * 8 * S + a * S + g] = v##a[(a + s + g) & 15];

__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)
}

On a Tiger Lake Iris Xe with IGC 2.41.5 and IGC_VISAOptions="-postRAspillCleanup 0" (to avoid #471), with 256 work-items in groups of 32 and s = 3 (host program below):

mismatch: vector 1, item 23, dynamic lane: got 0xc83884a9, want 0xb3bc390b
...
FAIL: 1664 of 34816 values wrong

All 1664 mismatches are dynamically indexed reads. With -DP=613, which selects HWord spills, the kernel passes, and it passes on an Arc A770 (DG2). Disabling spill space compression (-postRAspillCleanup 0 -nospillcompression) leaves the same 1664 mismatches.

We also hit this in oneAPI.jl (Julia) as a GPU hang: a loop selects one lane's address and clears the lanes with matching addresses. An indirect fill reads the wrong slot, so no lanes match and the loop never terminates. Correcting those fill offsets in the binary makes the kernel pass.

Root cause

insertAddrTakenSpillAndFillCode and its LS variant fill and spill indirectly accessed variables through temporary ADDR_SP_FL ranges. For OWord messages, sendInSpilledRegVarPortions and sendOutSpilledRegVarPortions derive the header offset from the temporary's displacement.

These temporaries are G4_RegVarTmp, not transient spill/fill ranges, so getDisp can allocate a separate slot instead of resolving the original variable's slot. With compression enabled, it calls calculateSpillDisp for the temporary; without compression, it reuses the base displacement only if that is already assigned.

The HWord and LSC paths instead use the base variable's displacement plus the row offset. A fix is to do the same in both OWord helpers: getDisp(range->getRegVar()->getBaseRegVar()) + srcRegOff * GRF.

The same code is on master (0ac43c0cf).

host.c 
cc host.c -lOpenCL -o host
IGC_VISAOptions='-postRAspillCleanup 0' ./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, nout = n + (size_t)V * G;
    cl_uint *in = malloc(n * sizeof *in), *out = calloc(nout, 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, nout * 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, nout * 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);
            }
    // out[V * G * W + v * G + g] should be lane (v + s + g) % W of vector v
    for (size_t v = 0; v < V; v++)
        for (size_t g = 0; g < G; g++) {
            cl_uint got = out[n + v * G + g];
            cl_uint want = in[(v * G + g) * W + (v + s + g) % W];
            if (got != want && bad++ < 10)
                printf("mismatch: vector %zu, item %zu, dynamic lane: got 0x%08x, want 0x%08x\n",
                       v, g, got, want);
        }
    printf("%s: %zu of %zu values wrong\n", bad ? "FAIL" : "PASS", bad, nout);
    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.