Arc A770 / UHD 730: GPU hang (or wrong values) on a double add of a sub-group shuffle under fp64 emulation

Start by compiling and running the provided repro.c with fp64 emulation enabled on an Arc A770 or UHD 730; compare variant 0 with the working as_long variant 1. The issue does not name repository files or tests, so trace how the OpenCL compiler handles double sub-group shuffle and addition under fp64 emulation. Done means the original kernel completes and returns v[i] + v[i ^ 1] for all 32 lanes on both GPUs.

Pre-submission Checklist

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

Driver 26.22.38646.4 with IGC 2.36.3 was tested; 26.35.39758.10 with IGC 2.41.5 was not, because the machine is a shared CI host whose driver I cannot change. The closest existing reports are https://github.com/intel/intel-graphics-compiler/issues/397 and https://github.com/intel/compute-runtime/issues/907; neither covers this case.

GPU Hardware

Intel Arc A770 (DG2), also reproduces on Intel UHD Graphics 730 (Alder Lake-S GT1)

DRI Devices Information

0 crw-rw----+ 1 root video  226,   0 spal.  18 14:34 /dev/dri/card0
0 crw-rw----+ 1 root video  226,   1 spal.  18 14:34 /dev/dri/card1
0 crw-rw----+ 1 root render 226, 128 spal.  18 14:34 /dev/dri/renderD128
0 crw-rw----+ 1 root render 226, 129 spal.  18 14:34 /dev/dri/renderD129

/dev/dri/by-path:
total 0
0 lrwxrwxrwx 1 root root  8 spal.  18 14:34 pci-0000:00:02.0-card -> ../card0
0 lrwxrwxrwx 1 root root 13 spal.  18 14:34 pci-0000:00:02.0-render -> ../renderD128
0 lrwxrwxrwx 1 root root  8 spal.  18 14:34 pci-0000:03:00.0-card -> ../card1
0 lrwxrwxrwx 1 root root 13 spal.  18 14:34 pci-0000:03:00.0-render -> ../renderD129

GPU Detailed Information (lspci output)

03:00.0 VGA compatible controller: Intel Corporation DG2 [Arc A770] (rev 08) (prog-if 00 [VGA controller])
	Subsystem: ASRock Incorporation DG2 [Arc A770]
	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
	Interrupt: pin ? routed to IRQ 136
	IOMMU group: 12
	Region 0: Memory at 71000000 (64-bit, non-prefetchable) [size=16M]
	Region 2: Memory at 6000000000 (64-bit, prefetchable) [size=16G]
	Expansion ROM at 72000000 [disabled] [size=2M]
	Capabilities: <access denied>
	Kernel driver in use: i915
	Kernel modules: i915, xe
00:02.0 VGA compatible controller: Intel Corporation Alder Lake-S GT1 [UHD Graphics 730] (rev 0c) (prog-if 00 [VGA controller])

Driver Version

26.22.38646.4 (IGC 2.36.3)

Installed GPU Driver Packages

ii intel-igc-core-2 2.36.3 amd64
ii intel-igc-opencl-2 2.36.3 amd64
ii intel-ocloc 26.22.38646.4-0 amd64
ii intel-opencl-icd 26.22.38646.4-0 amd64
ii level-zero 1.14.0 amd64
ii libigdgmm12:amd64 22.10.0 amd64
ii libze-intel-gpu1 26.22.38646.4-0 amd64
ii ocl-icd-libopencl1:amd64 2.3.2-1build1 amd64

Driver Installation Details

Installed as .deb packages, versions as listed above. Kernel driver i915.

Linux Distribution

Ubuntu 24.04 LTS (24.04.3)

Kernel Version & Boot Parameters

7.0.0-38-generic
BOOT_IMAGE=/boot/vmlinuz-7.0.0-38-generic root=UUID=68ad0da8-ccd8-416b-81d9-68cc95c35399 ro quiet splash i915.enable_hangcheck=0 i915.request_timeout_ms=0 vt.handoff=7
xe                   4366336  0
i915                 5025792  8

Actual Behavior

With fp64 emulation enabled (OverrideDefaultFP64Settings=1 IGC_EnableDPEmulation=1), a kernel with a required sub-group size that adds a double to a sub-group shuffle of that same double never completes on the Arc A770 (host blocks in clEnqueueReadBuffer; with default i915 hangcheck the context is reset and the read returns CL_OUT_OF_RESOURCES). On the UHD 730 the same kernel completes but returns wrong values (lane 0 = 0, expected 1).

__attribute__((intel_reqd_sub_group_size(32)))
__kernel void k(__global const double *v, __global double *y) {
  size_t i = get_global_id(0); double s = v[i];
  s += sub_group_shuffle(s, get_sub_group_local_id() ^ 1);
  y[i] = s;
}

Shuffling the bits as long instead (s += as_double(sub_group_shuffle(as_long(s), get_sub_group_local_id() ^ 1));) gives correct results on both GPUs.

Expected Behavior

The kernel completes and y[i] == v[i] + v[i ^ 1] for all 32 lanes, as with the as_long variant.

Reproduction Rate

Always reproduces - 100%

Steps to Reproduce

1. gcc repro.c -o repro -lOpenCL (source below)
2. OverrideDefaultFP64Settings=1 IGC_EnableDPEmulation=1 NEO_CACHE_PERSISTENT=0 ./repro 0 Arc hangs; ./repro 1 Arc prints PASS
3. OverrideDefaultFP64Settings=1 IGC_EnableDPEmulation=1 NEO_CACHE_PERSISTENT=0 ./repro 0 UHD prints a wrong value; ./repro 1 UHD prints PASS

Output:

=== dev=Arc variant=0 start 2026-10-26 22:34:35
Intel(R) Arc(TM) A770 Graphics, driver 26.22.38646.4, variant 0
rc=124 end 2026-10-26 22:35:35
=== dev=Arc variant=1 start 2026-10-26 22:35:35
Intel(R) Arc(TM) A770 Graphics, driver 26.22.38646.4, variant 1
PASS
rc=0 end 2026-10-26 22:35:35
=== dev=UHD variant=0 start 2026-10-26 22:35:35
Intel(R) UHD Graphics 730, driver 26.22.38646.4, variant 0
FAIL: lane 0 = 0, expected 1
rc=1 end 2026-10-26 22:35:36
=== dev=UHD variant=1 start 2026-10-26 22:35:36
Intel(R) UHD Graphics 730, driver 26.22.38646.4, variant 1
PASS
rc=0 end 2026-10-26 22:35:36

Is this a regression?

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

Unknown; only this driver version was tested.

Last Known Working Driver Version

Unknown

First Known Failing Driver Version

26.22.38646.4 (earliest tested)

API Call Logs

Not captured. The only calls are the ones in the reproducer; the host blocks in clEnqueueReadBuffer of the output buffer.

strace Logs

Not captured.

System Logs / dmesg Output

From an earlier run of the same kernel (on the A770 through Level Zero, same driver):

i915 0000:03:00.0: [drm] GPU HANG: ecode 12:10:85defdfb, in rocsparse-test [4137288]
i915 0000:03:00.0: [drm] rocsparse-test[4137288] context reset due to GPU hang

Backtrace (if crash or hang occurred)

Not applicable beyond the host waiting in clEnqueueReadBuffer (blocking read).

Source Code / Reproducer

gcc repro.c -o repro -lOpenCL

// GPU hang: double add of a sub-group shuffle under IGC fp64 emulation.
// Build: gcc repro.c -o repro -lOpenCL
// Run:   OverrideDefaultFP64Settings=1 IGC_EnableDPEmulation=1 ./repro [0|1] [device name substring]
//        0 = s + sub_group_shuffle(s, lane)  (hangs)
//        1 = same, shuffling as_long(s)      (works)
#define CL_TARGET_OPENCL_VERSION 300
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#define CHECK(x)                                                               \
  do {                                                                         \
    cl_int e_ = (x);                                                           \
    if (e_ != CL_SUCCESS) {                                                    \
      printf("FAIL: %s = %d (line %d)\n", #x, e_, __LINE__);                   \
      return 1;                                                                \
    }                                                                          \
  } while (0)

static const char *Src[] = {
    "__attribute__((intel_reqd_sub_group_size(32)))\n"
    "__kernel void k(__global const double *v, __global double *y) {\n"
    "  size_t i = get_global_id(0); double s = v[i];\n"
    "  s += sub_group_shuffle(s, get_sub_group_local_id() ^ 1);\n"
    "  y[i] = s;\n"
    "}\n",
    "__attribute__((intel_reqd_sub_group_size(32)))\n"
    "__kernel void k(__global const double *v, __global double *y) {\n"
    "  size_t i = get_global_id(0); double s = v[i];\n"
    "  s += as_double(sub_group_shuffle(as_long(s), get_sub_group_local_id() ^ 1));\n"
    "  y[i] = s;\n"
    "}\n"};

int main(int argc, char **argv) {
  setvbuf(stdout, NULL, _IONBF, 0);
  int V = argc > 1 ? atoi(argv[1]) : 0;
  cl_platform_id P[16];
  cl_uint NP;
  cl_device_id D = 0;
  CHECK(clGetPlatformIDs(16, P, &NP));
  char Name[256], Drv[64];
  for (cl_uint i = 0; i < NP && !D; i++) {
    cl_device_id Ds[16];
    cl_uint ND = 0;
    if (clGetDeviceIDs(P[i], CL_DEVICE_TYPE_GPU, 16, Ds, &ND) != CL_SUCCESS)
      continue;
    for (cl_uint j = 0; j < ND && !D; j++) {
      clGetDeviceInfo(Ds[j], CL_DEVICE_NAME, sizeof Name, Name, NULL);
      if (argc < 3 || strstr(Name, argv[2]))
        D = Ds[j];
    }
  }
  if (!D) {
    printf("FAIL: no GPU\n");
    return 1;
  }
  CHECK(clGetDeviceInfo(D, CL_DEVICE_NAME, sizeof Name, Name, NULL));
  CHECK(clGetDeviceInfo(D, CL_DRIVER_VERSION, sizeof Drv, Drv, NULL));
  printf("%s, driver %s, variant %d\n", Name, Drv, V);
  cl_int E;
  cl_context C = clCreateContext(NULL, 1, &D, NULL, NULL, &E);
  CHECK(E);
  cl_command_queue Q = clCreateCommandQueueWithProperties(C, D, NULL, &E);
  CHECK(E);
  cl_program Pr = clCreateProgramWithSource(C, 1, &Src[V], NULL, &E);
  CHECK(E);
  if (clBuildProgram(Pr, 1, &D, "-cl-std=CL3.0", NULL, NULL) != CL_SUCCESS) {
    char Log[8192];
    clGetProgramBuildInfo(Pr, D, CL_PROGRAM_BUILD_LOG, sizeof Log, Log, NULL);
    printf("FAIL: build\n%s\n", Log);
    return 1;
  }
  cl_kernel K = clCreateKernel(Pr, "k", &E);
  CHECK(E);
  double H[32], R[32];
  for (int i = 0; i < 32; i++)
    H[i] = i;
  cl_mem In = clCreateBuffer(C, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                             sizeof H, H, &E);
  CHECK(E);
  cl_mem Out = clCreateBuffer(C, CL_MEM_WRITE_ONLY, sizeof R, NULL, &E);
  CHECK(E);
  CHECK(clSetKernelArg(K, 0, sizeof In, &In));
  CHECK(clSetKernelArg(K, 1, sizeof Out, &Out));
  size_t G = 32, L = 32;
  CHECK(clEnqueueNDRangeKernel(Q, K, 1, NULL, &G, &L, 0, NULL, NULL));
  CHECK(clEnqueueReadBuffer(Q, Out, CL_TRUE, 0, sizeof R, R, 0, NULL, NULL));
  for (int i = 0; i < 32; i++)
    if (R[i] != H[i] + H[i ^ 1]) {
      printf("FAIL: lane %d = %g, expected %g\n", i, R[i], H[i] + H[i ^ 1]);
      return 1;
    }
  printf("PASS\n");
  return 0;
}

Command Line / Application Details

./repro <variant 0|1> <device name substring>, environment OverrideDefaultFP64Settings=1 IGC_EnableDPEmulation=1. Found through rocSPARSE's f64 bsrmv test running on chipStar (https://github.com/CHIP-SPV/chipStar), whose wavefront reduction does sum += __shfl_xor(sum, i) on doubles; tracked at https://github.com/CHIP-SPV/chipStar/issues/2106.

oneAPI Version (if applicable)

Not used.

Screenshots / Video

None.

Additional Notes

Variants of the kernel body, all with 32 work items in one work group and v[i] = i (Arc A770):

reqd 32  y = shuffle(v, lid^1)                        OK
reqd 32  loop m=4,2,1: s += shuffle(s, lid^m)         OK
reqd 32  s += shuffle(s, lid^1)                       hang
reqd 32  s += sub_group_shuffle_xor(s, 1)             hang
reqd 16  s += shuffle(s, lid^1)                       hang
reqd 32  y = shuffle(v, lid^1) + 1.0                  OK
reqd 32  s += as_double(shuffle(as_long(s), lid^1))   OK
no reqd  s += shuffle(s, lid^1)                       OK

IGC_EmulationFunctionControl=1, 2 and 3 (with NEO_CACHE_PERSISTENT=0) do not change the result, unlike the related https://github.com/intel/intel-graphics-compiler/issues/397. Without fp64 emulation the kernel does not build on these GPUs.

This issue was generated using AI