[GSD-13102] [xe/BMG] Hung compute kernel is never TDR-killed while the process lives (job_timeout_ms ignored); observed escalation: exit-time engine reset fails on aged state, wedging the GPU machine-wide until reboot

[pvelesko]

GPU Hardware

Intel Arc B570 (Battlemage, PCI 8086:e20c, 0000:09:00.0), kernel driver xe. Same box also has UHD 770 on i915 (used as behavioral contrast below).

OS / Kernel / Driver

• Ubuntu 24.04, kernel 6.17.0-14-generic
• intel-opencl-icd 26.22.38646.4, intel-igc-core-2 / intel-igc-opencl-2 2.36.3
• All engines report job_timeout_ms=5000, timeslice_duration_us=1000 (defaults) under /sys/class/drm/cardN/device/tile0/gt0/engines/*/

Summary

A trivially hung compute kernel (while (*flag == 0); where flag is never set — ~35-line plain-OpenCL reproducer below) is never killed while the submitting process is alive:

1. clFinish blocks indefinitely — measured 180 s+ with zero dmesg activity, despite job_timeout_ms=5000. No error, no reset, no log line. The application has no way to detect or recover from the hang (found via HeCBench tqs-hip, where an out-of-bounds work-queue read produced exactly this silent permanent hang; reported on the app side as OpenCL: out-of-bounds device read silently wedges the queue (no error, no diagnostic) — Arc B570 CHIP-SPV/chipStar#1364).
2. Only when the process exits does the driver act: xe 0000:09:00.0: [drm] GT0: Engine reset: engine_class=ccs, logical_mask: 0x1, guc_id=2 — on a freshly booted system this recovers the GPU for other processes.
3. Observed escalation (once, ~6 days uptime with prior device faults in the session): after the same class of hang, the exit-time reset did not recover the engine. Every subsequent context from every process hung on any kernel (even a 64-thread x[gid]=gid), dmesg stayed completely silent, and neither xe unbind/rebind nor PCI function reset (echo 1 > .../reset) recovered the device — only a reboot did. I am reporting this as observed-once with those conditions rather than a claimed invariant; it matches the 'permanently wedged, restart does not recover' flavor of [GSD-13010] Sporadic permanent GPU wedge on dual Arc Pro B70 (BMG G31): ccs/bcs engine reset + "Fault response: Unsuccessful -ENOENT/-EINVAL" under sustained Level-Zero inference load #948.

For contrast: the same hang class submitted to the UHD 770 (i915, same box, same day) is detected and context-reset within seconds, with proper GPU HANG ... context reset dmesg lines.

Expected Behavior

The GuC/driver job timeout should kill a hung compute job within job_timeout_ms while the process is alive and return CL_OUT_OF_RESOURCES-class errors to the client — as i915 does — and an engine reset must reliably return the engine to a usable state.

Minimal reproducer (deterministic, ~30 s)

/* gcc hang.c -o hang -lOpenCL && timeout 200 ./hang ; echo rc=$? */
/* rc=42 => no TDR within 180 s (bug). A healthy driver returns an error in ~seconds. */
#define CL_TARGET_OPENCL_VERSION 300
#include <CL/cl.h>
#include <signal.h>
#include <stdio.h>
#include <string.h>
#include <unistd.h>
static void wd(int s){(void)s;(void)!write(1,"WEDGED: no TDR within 180 s\n",27);_exit(42);}
int main(void) {
  setbuf(stdout, NULL);
  cl_platform_id plats[8]; cl_uint np; clGetPlatformIDs(8, plats, &np);
  cl_device_id dev = NULL;
  for (cl_uint i = 0; i < np && !dev; i++) {
    cl_device_id d[4]; cl_uint nd;
    if (clGetDeviceIDs(plats[i], CL_DEVICE_TYPE_GPU, 4, d, &nd) != CL_SUCCESS) continue;
    for (cl_uint j = 0; j < nd; j++) { char n[256]={0}; clGetDeviceInfo(d[j], CL_DEVICE_NAME, 256, n, NULL);
      if (strstr(n, "Arc")) { dev = d[j]; break; } } }
  if (!dev) { printf("no Arc\n"); return 1; }
  cl_int e; cl_context c = clCreateContext(NULL,1,&dev,NULL,NULL,&e);
  cl_command_queue q = clCreateCommandQueueWithProperties(c,dev,NULL,&e);
  const char *src = "__kernel void spin(__global volatile int *flag) { while (*flag == 0); }\n";
  cl_program p = clCreateProgramWithSource(c,1,&src,NULL,&e);
  if (clBuildProgram(p,1,&dev,"",NULL,NULL)!=CL_SUCCESS){printf("build fail\n");return 1;}
  cl_kernel k = clCreateKernel(p,"spin",&e);
  int zero = 0;
  cl_mem flag = clCreateBuffer(c, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, 4, &zero, &e);
  clSetKernelArg(k,0,sizeof(cl_mem),&flag);
  size_t g=64,l=64;
  clEnqueueNDRangeKernel(q,k,1,NULL,&g,&l,0,NULL,NULL);
  signal(SIGALRM, wd); alarm(180);
  e = clFinish(q); alarm(0);
  printf("clFinish returned %d (driver handled the hang)\n", (int)e);
  return e == CL_SUCCESS ? 1 : 0;
}

Steps:

1. gcc hang.c -o hang -lOpenCL && timeout 200 ./hang → prints WEDGED: no TDR within 180 s, rc=42, dmesg silent (defect 1).
2. After the process exits, dmesg logs GT0: Engine reset: engine_class=ccs; on fresh boot a subsequent trivial kernel from a new process works again.
3. The escalated machine-wide wedge (defect 3 above) was captured after repeated device faults on ~6 days uptime; I can attempt to re-age into that state if useful.

Reproduction Rate

Defect 1 (no in-flight TDR): 100%, every run. Defect 3 (failed exit-time recovery): observed once under the stated conditions.

[kgibala]

Hi @pvelesko,

Thank you for your contribution.

Could you please use the GitHub issue template when creating tickets, as recommended and established here:

• [GSD-12610] clEnqueueSVMMemcpy with invalid destination pointer and committed source buffer wedges process unkillably (i915 kernel regression) #914 (comment)
• [GSD-12610] clEnqueueSVMMemcpy with invalid destination pointer and committed source buffer wedges process unkillably (i915 kernel regression) #914 (comment)

We designed the templates to ensure all essential information is captured, so following them really helps us out. You're welcome to edit this issue accordingly.

Thank you for your understanding.