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

オープン
  • 何か問題がありますか？
#472 コメント 0 件 リアクション 0 件 担当者 0 名 GitHub で見る

@maleadt がすでに取り組んでいます。

2026年11月7日 から。

  • #473 @maleadt による — オープン
  • #474 @maleadt による — オープン

評価

難易度
4/5
見積もり時間
3〜5日
初心者へのやさしさ
53/100
issue の種類
バグ
明瞭さ
明確に書かれている
活発さ
活発
技術スタック
cpp
領域
compilers

調査の方向性

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.

索引モデルが issue の本文から書いたものです。

説明

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;
}
主要言語
C++
スター
725
フォーク
200
30日以内にマージされた PR はありません

環境構築

はじめの一歩

  1. issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
  2. 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
  3. リポジトリをフォークし、ブランチを切って変更します。
  4. issue 番号を参照したプルリクエストを送ります。

intel/intel-graphics-compiler のほかの issue

intel/intel-graphics-compiler の issue をすべて見る

似ている issue

C++ の issue をもっと見る

新しい issue をメールで受け取る

初心者向けの GitHub issue を短くまとめたダイジェスト。