IGC silently drops kernel at -O2; SPIR-V well-formed, L0 build succeeds, zeKernelCreate fails with INVALID_KERNEL_NAME; -cl-opt-disable recovers (Arc A770, Q_1D=9 double kernel, libCEED GradAtPoints)

## Summary

IGC silently drops a `__global__` kernel from the compiled module at default
optimization. The kernel is declared as `OpEntryPoint Kernel` in the SPIR-V
handed to `zeModuleCreate`, so the SPIR-V is well formed; the IGC backend
just stops exposing it to `zeModuleGetKernelNames` / `zeKernelCreate`.
`clCreateKernel` / `zeKernelCreate` returns `CL_INVALID_KERNEL_NAME` /
`ZE_RESULT_ERROR_INVALID_KERNEL_NAME` for the missing name.

Building the same SPIR-V with `-cl-opt-disable` (passed via
`ze_module_desc_t::pBuildFlags`) recovers the kernel. No finer-grained IGC
env var works; I tried ~40 of them. Details below.

## Environment

- Intel Arc A770 (Xe-HPG, DG2)
- Driver: NEO 26.09.37435.1, Level Zero 1.28
- Compiler frontend: clang HIP frontend via chipStar HIPRTC (the bug reproduces
  from the SPIR-V alone; chipStar is only the producer)
- Source: libCEED `hip-shared-basis-tensor-at-points-templates.h`
  (`GradAtPoints2d`, `Q_1D = 9`)

## Reproducing

Self-contained pure-Level-Zero reproducer (no HIP / chipStar / OpenCL runtime needed). All source inlined below; re-assemble the SPIR-V with `spirv-as` and link the host against `-lze_loader`.

```bash
make && ./repro                   # GradAtPoints MISSING, exit 1
./repro -cl-opt-disable           # GradAtPoints PRESENT, exit 0
```

## Trigger

All of the following are required; removing any one makes the drop disappear:

- Outer `for` loop with trip count ≥ 2 containing an `if/else` calling two
  different `__device__ __forceinline__` helpers that both write to the same
  local `double[QN]` array (`x`).
- A nested reduction reading `x[]` against `extern __shared__ double` indexed
  by `threadIdx.z * QN * QN` (SLM offset that depends on `threadIdx.z` at
  runtime).
- A second call to one of the helpers after the reduction, overwriting `x[]`.
- A final reduction accumulating into `r_V`.
- **`QN == 9`** — exact. Tested 3, 4, 5, 6, 7, 8, 10, 11, 12, 16. None
  reproduces.
- `double` precision.

## IGC shader-dump evidence (`IGC_ShaderDumpEnable=1`)

The dropped `GradAtPoints` kernel is **present in every LLVM stage**:

```
beforeUnification.ll    : 1 spir_kernel declaration  (GradAtPoints)
afterUnification.ll     : 1                         (GradAtPoints)
optimized.ll            : 1                         (GradAtPoints)
codegen.ll              : 1                         (GradAtPoints)
push_analysis.ll        : 1                         (GradAtPoints)
```

…but **no corresponding `entry_*.asm` / `entry_*.visaasm` is produced for it**
at any SIMD width. With `-cl-opt-disable`, `GradAtPoints` compiles as SIMD8
with 9 inlined sub-functions (`_f0..f8`) and a valid VISA/native entry is
emitted.

This points at a late vISA-emit stage failure: `IsValidShader()` returns false
(`m_programSize == 0`) for all SIMD widths, and
`OpenCLKernelCodeGen.cpp::GatherDataForDriver` silently skips the kernel. The
`RetryManager` path exhausts SIMD32 → SIMD16 → SIMD8 without producing a valid
binary, and the kernel is dropped rather than surfacing the error.

The SIMD8 compile of the same SPIR-V with `-cl-opt-disable` succeeds with 9
sub-function splits — evidence that the default IGC pipeline's more aggressive
function inlining for `Q_1D == 9` produces vISA that register allocation can't
satisfy at any SIMD width.

## IGC flags tried that **do NOT** help

Non-exhaustive list (none produces `GradAtPoints`):

`IGC_FunctionControl=0`, `IGC_ForceInlineAll=0`, `IGC_EnableFunctionCloning=0`,
`IGC_ForceFunctionsNoinline=1`, `IGC_EnableSubroutine=1`,
`IGC_EnableFunctionPointer=1`, `IGC_DisableDeadFunctionsEliminationPass=1`,
`IGC_DisableEmptyFunctionEraser=1`, `IGC_DisableInternalization=1`,
`IGC_SIMD{8,16,32}_SpillThreshold=100`, `IGC_ForceAllowSmallSpill=1`,
`IGC_DisableSubroutines=1`, `IGC_FunctionSplittingThreshold={0,100000}`,
`IGC_EnableFunctionSplitting={0,1}`, `IGC_RetryRegAlloc=1`,
`IGC_ExtraRetrySIMD16=1`, `IGC_EnableSIMD8CompileAttemptFirst=1`,
`IGC_AllowStackCallRetry={0,1,2}`, `IGC_DisableLoopUnroll=1`,
`IGC_DisableRegAllocOpt=1`, `IGC_DisableRegAlloc=1`,
`IGC_DisableCompaction=1`, `IGC_forceGlobalRA=1`, `IGC_KeepAllVariables=1`.

Also tried as **build flags** via `ze_module_desc_t::pBuildFlags`:

- `-O0`, `-O1`, `-O2`, `-cl-opt-level=0` — no effect (ignored?)
- `-cl-fast-relaxed-math`, `-cl-no-signed-zeros`, `-cl-uniform-work-group-size` — no effect

Only `-cl-opt-disable` and `-ze-opt-disable` work, and they disable *all* IGC
optimization. That's a ~10× perf hit on the compute-bound kernels this pattern
appears in (libCEED FEM basis evaluation), so it's not a viable shipping
workaround — it only serves as a diagnostic fence.

## Related previously-reported issues

- [intel/intel-graphics-compiler#403](https://github.com/intel/intel-graphics-compiler/issues/403)
  — same failure class (kernels silently omitted from final native object
  after `RetryManager` exhausts SIMD widths). Different trigger (rocThrust
  lookback template filter at link step). Open.
- [intel/compute-runtime#683](https://github.com/intel/compute-runtime/issues/683)
  — **same surface symptom against the same codebase** (CHIP-SPV / libCEED
  `CeedKernel*` missing from module, `zeKernelCreate` fails by name).
  Reporter diagnosed *"register spills + user local memory allocation usage
  exceeds the hardware SLM size"*, matching this fingerprint. Intel triager
  requested a zebin, reporter didn't respond, bot auto-closed. **This is that
  bug with the zebin-and-then-some.**

Not the same as #403 (different trigger) nor as compute-runtime#683 (wrong
repo — failure is in IGC codegen, not NEO). Filing fresh here with the
zebin-equivalent evidence Intel asked for in 2023.

## The kernel (HIP source, 52 LOC)

```cpp
// kernel.hip — compiled via chipStar HIPRTC at -O2 into the SPIR-V below
#include <hip/hip_runtime.h>

#define QN 9
#define NP 4

__device__ __forceinline__ void polyA(const double x, double *cx) {
  cx[0] = 1.0; cx[1] = 2 * x;
  for (int i = 2; i < QN; i++) cx[i] = 2 * x * cx[i - 1] - cx[i - 2];
}
__device__ __forceinline__ void polyB(const double x, double *cx) {
  cx[0] = x; cx[1] = 1.0;
  for (int i = 2; i < QN; i++) cx[i] = x * cx[i - 1] + cx[i - 2];
}

__device__ __forceinline__ void GradAtPoints2d(double *slice, int p,
                                               const double *r_X,
                                               double *r_V) {
  double buf[QN], x[QN];
  for (int dim = 0; dim < 3; dim++) {
    if (p == 0) polyA(r_X[0], x);
    else polyB(r_X[0], x);
    for (int i = 0; i < QN; i++) {
      buf[i] = 0.0;
      for (int j = 0; j < QN; j++) buf[i] += x[j] * slice[j + i * QN];
    }
    polyB(r_X[1], x);
    for (int i = 0; i < QN; i++) r_V[0] += x[i] * buf[i];
  }
}

extern "C" __global__
void GradAtPoints(const double *__restrict__ d_X, double *__restrict__ d_V) {
  extern __shared__ double slice_raw[];
  double *slice = slice_raw + threadIdx.z * QN * QN;
  int p = threadIdx.x;
  double r_X[2] = { d_X[p], d_X[p + NP] };
  double r_V[2] = { 0, 0 };
  GradAtPoints2d(slice, p, r_X, r_V);
  d_V[p] = r_V[0];
  d_V[p + NP] = r_V[1];
}
```

## SPIR-V (718 LOC, result of `clang -O2` + `llvm-spirv` + `spirv-opt --eliminate-dead-*` + line-bisection)

<details>
<summary>Click to expand SPIR-V assembly</summary>

```
               OpEntryPoint Kernel %2 "GradAtPoints" %gl_LocalInvocationID
               OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
               OpDecorate %6 LinkageAttributes "GradAtPoints" Export
               OpDecorate %8 FuncParamAttr NoAlias
               OpDecorate %9 FuncParamAttr NoAlias
               OpDecorate %10 FuncParamAttr NoAlias
      %ulong = OpTypeInt 64 0
      %uchar = OpTypeInt 8 0
       %uint = OpTypeInt 32 0
    %ulong_0 = OpConstant %ulong 0
    %uchar_1 = OpConstant %uchar 1
    %uint_81 = OpConstant %uint 81
     %uint_4 = OpConstant %uint 4
     %uint_0 = OpConstant %uint 0
    %ulong_8 = OpConstant %ulong 8
   %ulong_16 = OpConstant %ulong 16
   %ulong_24 = OpConstant %ulong 24
   %ulong_32 = OpConstant %ulong 32
   %ulong_40 = OpConstant %ulong 40
   %ulong_48 = OpConstant %ulong 48
   %ulong_56 = OpConstant %ulong 56
   %ulong_64 = OpConstant %ulong 64
   %ulong_72 = OpConstant %ulong 72
   %ulong_80 = OpConstant %ulong 80
   %ulong_88 = OpConstant %ulong 88
   %ulong_96 = OpConstant %ulong 96
  %ulong_104 = OpConstant %ulong 104
  %ulong_112 = OpConstant %ulong 112
  %ulong_120 = OpConstant %ulong 120
  %ulong_128 = OpConstant %ulong 128
  %ulong_136 = OpConstant %ulong 136
  %ulong_144 = OpConstant %ulong 144
  %ulong_152 = OpConstant %ulong 152
  %ulong_160 = OpConstant %ulong 160
  %ulong_168 = OpConstant %ulong 168
  %ulong_176 = OpConstant %ulong 176
  %ulong_184 = OpConstant %ulong 184
  %ulong_192 = OpConstant %ulong 192
  %ulong_200 = OpConstant %ulong 200
  %ulong_208 = OpConstant %ulong 208
  %ulong_216 = OpConstant %ulong 216
  %ulong_224 = OpConstant %ulong 224
  %ulong_232 = OpConstant %ulong 232
  %ulong_240 = OpConstant %ulong 240
  %ulong_248 = OpConstant %ulong 248
  %ulong_256 = OpConstant %ulong 256
  %ulong_264 = OpConstant %ulong 264
  %ulong_272 = OpConstant %ulong 272
  %ulong_280 = OpConstant %ulong 280
  %ulong_288 = OpConstant %ulong 288
  %ulong_296 = OpConstant %ulong 296
  %ulong_304 = OpConstant %ulong 304
  %ulong_312 = OpConstant %ulong 312
  %ulong_320 = OpConstant %ulong 320
  %ulong_328 = OpConstant %ulong 328
  %ulong_336 = OpConstant %ulong 336
  %ulong_344 = OpConstant %ulong 344
  %ulong_352 = OpConstant %ulong 352
  %ulong_360 = OpConstant %ulong 360
  %ulong_368 = OpConstant %ulong 368
  %ulong_376 = OpConstant %ulong 376
  %ulong_384 = OpConstant %ulong 384
  %ulong_392 = OpConstant %ulong 392
  %ulong_400 = OpConstant %ulong 400
  %ulong_408 = OpConstant %ulong 408
  %ulong_416 = OpConstant %ulong 416
  %ulong_424 = OpConstant %ulong 424
  %ulong_432 = OpConstant %ulong 432
  %ulong_440 = OpConstant %ulong 440
  %ulong_448 = OpConstant %ulong 448
  %ulong_456 = OpConstant %ulong 456
  %ulong_464 = OpConstant %ulong 464
  %ulong_472 = OpConstant %ulong 472
  %ulong_480 = OpConstant %ulong 480
  %ulong_488 = OpConstant %ulong 488
  %ulong_496 = OpConstant %ulong 496
  %ulong_504 = OpConstant %ulong 504
  %ulong_512 = OpConstant %ulong 512
  %ulong_520 = OpConstant %ulong 520
  %ulong_528 = OpConstant %ulong 528
  %ulong_536 = OpConstant %ulong 536
  %ulong_544 = OpConstant %ulong 544
  %ulong_552 = OpConstant %ulong 552
  %ulong_560 = OpConstant %ulong 560
  %ulong_568 = OpConstant %ulong 568
  %ulong_576 = OpConstant %ulong 576
  %ulong_584 = OpConstant %ulong 584
  %ulong_592 = OpConstant %ulong 592
  %ulong_600 = OpConstant %ulong 600
  %ulong_608 = OpConstant %ulong 608
  %ulong_616 = OpConstant %ulong 616
  %ulong_624 = OpConstant %ulong 624
  %ulong_632 = OpConstant %ulong 632
  %ulong_640 = OpConstant %ulong 640
     %uint_1 = OpConstant %uint 1
     %uint_2 = OpConstant %uint 2
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
    %v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
       %void = OpTypeVoid
     %double = OpTypeFloat 64
%_ptr_Workgroup_double = OpTypePointer Workgroup %double
        %108 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar %_ptr_Workgroup_double
%_ptr_Generic_double = OpTypePointer Generic %double
       %bool = OpTypeBool
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
        %112 = OpTypeFunction %void %_ptr_CrossWorkgroup_ulong
        %113 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar
        %114 = OpTypeFunction %void
          %4 = OpVariable %_ptr_CrossWorkgroup_ulong CrossWorkgroup %ulong_0
          %5 = OpVariable %_ptr_CrossWorkgroup_uchar CrossWorkgroup %uchar_1
%gl_LocalInvocationID = OpVariable %_ptr_Input_v3ulong Input
       %true = OpConstantTrue %bool
   %double_2 = OpConstant %double 2
   %double_1 = OpConstant %double 1
  %double_n1 = OpConstant %double -1
   %double_0 = OpConstant %double 0
          %6 = OpFunction %void None %108
          %7 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
          %8 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
        %120 = OpFunctionParameter %_ptr_Workgroup_double
        %121 = OpLabel
        %122 = OpConvertPtrToU %ulong %7
        %123 = OpConvertUToPtr %_ptr_Generic_double %122
        %124 = OpLoad %v3ulong %gl_LocalInvocationID Aligned 32
        %125 = OpCompositeExtract %ulong %124 2
        %126 = OpSelect %ulong %true %125 %ulong_0
        %127 = OpUConvert %uint %126
        %128 = OpIMul %uint %127 %uint_81
        %129 = OpUConvert %ulong %128
        %130 = OpInBoundsPtrAccessChain %_ptr_Workgroup_double %120 %129
        %131 = OpLoad %v3ulong %gl_LocalInvocationID Aligned 32
        %132 = OpCompositeExtract %ulong %131 0
        %133 = OpSelect %ulong %true %132 %ulong_0
        %134 = OpUConvert %uint %133
        %135 = OpSConvert %ulong %134
        %136 = OpInBoundsPtrAccessChain %_ptr_Generic_double %123 %135
        %137 = OpLoad %double %136 Aligned 8
        %138 = OpIAdd %uint %134 %uint_4
        %139 = OpSConvert %ulong %138
        %140 = OpInBoundsPtrAccessChain %_ptr_Generic_double %123 %139
        %141 = OpLoad %double %140 Aligned 8
        %142 = OpIEqual %bool %134 %uint_0
        %143 = OpFMul %double %137 %double_2
        %144 = OpFAdd %double %137 %137
        %145 = OpFMul %double %137 %144
        %146 = OpFAdd %double %145 %double_1
        %147 = OpFMul %double %137 %146
        %148 = OpFAdd %double %147 %144
        %149 = OpFMul %double %137 %148
        %150 = OpFAdd %double %149 %146
        %151 = OpFMul %double %137 %150
        %152 = OpFAdd %double %151 %148
        %153 = OpFMul %double %137 %152
        %154 = OpFAdd %double %153 %150
        %155 = OpFMul %double %137 %154
        %156 = OpFAdd %double %155 %152
        %157 = OpFMul %double %143 %143
        %158 = OpFAdd %double %157 %double_n1
        %159 = OpFMul %double %143 %158
        %160 = OpFSub %double %159 %143
        %161 = OpFMul %double %143 %160
        %162 = OpFSub %double %161 %158
        %163 = OpFMul %double %143 %162
        %164 = OpFSub %double %163 %160
        %165 = OpFMul %double %143 %164
        %166 = OpFSub %double %165 %162
        %167 = OpFMul %double %143 %166
        %168 = OpFSub %double %167 %164
        %169 = OpFMul %double %143 %168
        %170 = OpFSub %double %169 %166
        %171 = OpLoad %double %130 Aligned 8
        %172 = OpBitcast %_ptr_Workgroup_uchar %130
        %173 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %172 %ulong_8
        %174 = OpBitcast %_ptr_Workgroup_double %173
        %175 = OpLoad %double %174 Aligned 8
        %176 = OpBitcast %_ptr_Workgroup_uchar %130
        %177 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %176 %ulong_16
        %178 = OpBitcast %_ptr_Workgroup_double %177
        %179 = OpLoad %double %178 Aligned 8
        %180 = OpBitcast %_ptr_Workgroup_uchar %130
        %181 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %180 %ulong_24
        %182 = OpBitcast %_ptr_Workgroup_double %181
        %183 = OpLoad %double %182 Aligned 8
        %184 = OpBitcast %_ptr_Workgroup_uchar %130
        %185 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %184 %ulong_32
        %186 = OpBitcast %_ptr_Workgroup_double %185
        %187 = OpLoad %double %186 Aligned 8
        %188 = OpBitcast %_ptr_Workgroup_uchar %130
        %189 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %188 %ulong_40
        %190 = OpBitcast %_ptr_Workgroup_double %189
        %191 = OpLoad %double %190 Aligned 8
        %192 = OpBitcast %_ptr_Workgroup_uchar %130
        %193 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %192 %ulong_48
        %194 = OpBitcast %_ptr_Workgroup_double %193
        %195 = OpLoad %double %194 Aligned 8
        %196 = OpBitcast %_ptr_Workgroup_uchar %130
        %197 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %196 %ulong_56
        %198 = OpBitcast %_ptr_Workgroup_double %197
        %199 = OpLoad %double %198 Aligned 8
        %200 = OpBitcast %_ptr_Workgroup_uchar %130
        %201 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %200 %ulong_64
        %202 = OpBitcast %_ptr_Workgroup_double %201
        %203 = OpLoad %double %202 Aligned 8
        %204 = OpBitcast %_ptr_Workgroup_uchar %130
        %205 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %204 %ulong_72
        %206 = OpBitcast %_ptr_Workgroup_double %205
        %207 = OpLoad %double %206 Aligned 8
        %208 = OpBitcast %_ptr_Workgroup_uchar %130
        %209 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %208 %ulong_80
        %210 = OpBitcast %_ptr_Workgroup_double %209
        %211 = OpLoad %double %210 Aligned 8
        %212 = OpBitcast %_ptr_Workgroup_uchar %130
        %213 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %212 %ulong_88
        %214 = OpBitcast %_ptr_Workgroup_double %213
        %215 = OpLoad %double %214 Aligned 8
        %216 = OpBitcast %_ptr_Workgroup_uchar %130
        %217 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %216 %ulong_96
        %218 = OpBitcast %_ptr_Workgroup_double %217
        %219 = OpLoad %double %218 Aligned 8
        %220 = OpBitcast %_ptr_Workgroup_uchar %130
        %221 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %220 %ulong_104
        %222 = OpBitcast %_ptr_Workgroup_double %221
        %223 = OpLoad %double %222 Aligned 8
        %224 = OpBitcast %_ptr_Workgroup_uchar %130
        %225 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %224 %ulong_112
        %226 = OpBitcast %_ptr_Workgroup_double %225
        %227 = OpLoad %double %226 Aligned 8
        %228 = OpBitcast %_ptr_Workgroup_uchar %130
        %229 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %228 %ulong_120
        %230 = OpBitcast %_ptr_Workgroup_double %229
        %231 = OpLoad %double %230 Aligned 8
        %232 = OpBitcast %_ptr_Workgroup_uchar %130
        %233 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %232 %ulong_128
        %234 = OpBitcast %_ptr_Workgroup_double %233
        %235 = OpLoad %double %234 Aligned 8
        %236 = OpBitcast %_ptr_Workgroup_uchar %130
        %237 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %236 %ulong_136
        %238 = OpBitcast %_ptr_Workgroup_double %237
        %239 = OpLoad %double %238 Aligned 8
        %240 = OpBitcast %_ptr_Workgroup_uchar %130
        %241 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %240 %ulong_144
        %242 = OpBitcast %_ptr_Workgroup_double %241
        %243 = OpLoad %double %242 Aligned 8
        %244 = OpBitcast %_ptr_Workgroup_uchar %130
        %245 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %244 %ulong_152
        %246 = OpBitcast %_ptr_Workgroup_double %245
        %247 = OpLoad %double %246 Aligned 8
        %248 = OpBitcast %_ptr_Workgroup_uchar %130
        %249 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %248 %ulong_160
        %250 = OpBitcast %_ptr_Workgroup_double %249
        %251 = OpLoad %double %250 Aligned 8
        %252 = OpBitcast %_ptr_Workgroup_uchar %130
        %253 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %252 %ulong_168
        %254 = OpBitcast %_ptr_Workgroup_double %253
        %255 = OpLoad %double %254 Aligned 8
        %256 = OpBitcast %_ptr_Workgroup_uchar %130
        %257 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %256 %ulong_176
        %258 = OpBitcast %_ptr_Workgroup_double %257
        %259 = OpLoad %double %258 Aligned 8
        %260 = OpBitcast %_ptr_Workgroup_uchar %130
        %261 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %260 %ulong_184
        %262 = OpBitcast %_ptr_Workgroup_double %261
        %263 = OpLoad %double %262 Aligned 8
        %264 = OpBitcast %_ptr_Workgroup_uchar %130
        %265 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %264 %ulong_192
        %266 = OpBitcast %_ptr_Workgroup_double %265
        %267 = OpLoad %double %266 Aligned 8
        %268 = OpBitcast %_ptr_Workgroup_uchar %130
        %269 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %268 %ulong_200
        %270 = OpBitcast %_ptr_Workgroup_double %269
        %271 = OpLoad %double %270 Aligned 8
        %272 = OpBitcast %_ptr_Workgroup_uchar %130
        %273 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %272 %ulong_208
        %274 = OpBitcast %_ptr_Workgroup_double %273
        %275 = OpLoad %double %274 Aligned 8
        %276 = OpBitcast %_ptr_Workgroup_uchar %130
        %277 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %276 %ulong_216
        %278 = OpBitcast %_ptr_Workgroup_double %277
        %279 = OpLoad %double %278 Aligned 8
        %280 = OpBitcast %_ptr_Workgroup_uchar %130
        %281 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %280 %ulong_224
        %282 = OpBitcast %_ptr_Workgroup_double %281
        %283 = OpLoad %double %282 Aligned 8
        %284 = OpBitcast %_ptr_Workgroup_uchar %130
        %285 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %284 %ulong_232
        %286 = OpBitcast %_ptr_Workgroup_double %285
        %287 = OpLoad %double %286 Aligned 8
        %288 = OpBitcast %_ptr_Workgroup_uchar %130
        %289 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %288 %ulong_240
        %290 = OpBitcast %_ptr_Workgroup_double %289
        %291 = OpLoad %double %290 Aligned 8
        %292 = OpBitcast %_ptr_Workgroup_uchar %130
        %293 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %292 %ulong_248
        %294 = OpBitcast %_ptr_Workgroup_double %293
        %295 = OpLoad %double %294 Aligned 8
        %296 = OpBitcast %_ptr_Workgroup_uchar %130
        %297 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %296 %ulong_256
        %298 = OpBitcast %_ptr_Workgroup_double %297
        %299 = OpLoad %double %298 Aligned 8
        %300 = OpBitcast %_ptr_Workgroup_uchar %130
        %301 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %300 %ulong_264
        %302 = OpBitcast %_ptr_Workgroup_double %301
        %303 = OpLoad %double %302 Aligned 8
        %304 = OpBitcast %_ptr_Workgroup_uchar %130
        %305 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %304 %ulong_272
        %306 = OpBitcast %_ptr_Workgroup_double %305
        %307 = OpLoad %double %306 Aligned 8
        %308 = OpBitcast %_ptr_Workgroup_uchar %130
        %309 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %308 %ulong_280
        %310 = OpBitcast %_ptr_Workgroup_double %309
        %311 = OpLoad %double %310 Aligned 8
        %312 = OpBitcast %_ptr_Workgroup_uchar %130
        %313 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %312 %ulong_288
        %314 = OpBitcast %_ptr_Workgroup_double %313
        %315 = OpLoad %double %314 Aligned 8
        %316 = OpBitcast %_ptr_Workgroup_uchar %130
        %317 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %316 %ulong_296
        %318 = OpBitcast %_ptr_Workgroup_double %317
        %319 = OpLoad %double %318 Aligned 8
        %320 = OpBitcast %_ptr_Workgroup_uchar %130
        %321 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %320 %ulong_304
        %322 = OpBitcast %_ptr_Workgroup_double %321
        %323 = OpLoad %double %322 Aligned 8
        %324 = OpBitcast %_ptr_Workgroup_uchar %130
        %325 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %324 %ulong_312
        %326 = OpBitcast %_ptr_Workgroup_double %325
        %327 = OpLoad %double %326 Aligned 8
        %328 = OpBitcast %_ptr_Workgroup_uchar %130
        %329 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %328 %ulong_320
        %330 = OpBitcast %_ptr_Workgroup_double %329
        %331 = OpLoad %double %330 Aligned 8
        %332 = OpBitcast %_ptr_Workgroup_uchar %130
        %333 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %332 %ulong_328
        %334 = OpBitcast %_ptr_Workgroup_double %333
        %335 = OpLoad %double %334 Aligned 8
        %336 = OpBitcast %_ptr_Workgroup_uchar %130
        %337 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %336 %ulong_336
        %338 = OpBitcast %_ptr_Workgroup_double %337
        %339 = OpLoad %double %338 Aligned 8
        %340 = OpBitcast %_ptr_Workgroup_uchar %130
        %341 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %340 %ulong_344
        %342 = OpBitcast %_ptr_Workgroup_double %341
        %343 = OpLoad %double %342 Aligned 8
        %344 = OpBitcast %_ptr_Workgroup_uchar %130
        %345 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %344 %ulong_352
        %346 = OpBitcast %_ptr_Workgroup_double %345
        %347 = OpLoad %double %346 Aligned 8
        %348 = OpBitcast %_ptr_Workgroup_uchar %130
        %349 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %348 %ulong_360
        %350 = OpBitcast %_ptr_Workgroup_double %349
        %351 = OpLoad %double %350 Aligned 8
        %352 = OpBitcast %_ptr_Workgroup_uchar %130
        %353 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %352 %ulong_368
        %354 = OpBitcast %_ptr_Workgroup_double %353
        %355 = OpLoad %double %354 Aligned 8
        %356 = OpBitcast %_ptr_Workgroup_uchar %130
        %357 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %356 %ulong_376
        %358 = OpBitcast %_ptr_Workgroup_double %357
        %359 = OpLoad %double %358 Aligned 8
        %360 = OpBitcast %_ptr_Workgroup_uchar %130
        %361 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %360 %ulong_384
        %362 = OpBitcast %_ptr_Workgroup_double %361
        %363 = OpLoad %double %362 Aligned 8
        %364 = OpBitcast %_ptr_Workgroup_uchar %130
        %365 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %364 %ulong_392
        %366 = OpBitcast %_ptr_Workgroup_double %365
        %367 = OpLoad %double %366 Aligned 8
        %368 = OpBitcast %_ptr_Workgroup_uchar %130
        %369 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %368 %ulong_400
        %370 = OpBitcast %_ptr_Workgroup_double %369
        %371 = OpLoad %double %370 Aligned 8
        %372 = OpBitcast %_ptr_Workgroup_uchar %130
        %373 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %372 %ulong_408
        %374 = OpBitcast %_ptr_Workgroup_double %373
        %375 = OpLoad %double %374 Aligned 8
        %376 = OpBitcast %_ptr_Workgroup_uchar %130
        %377 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %376 %ulong_416
        %378 = OpBitcast %_ptr_Workgroup_double %377
        %379 = OpLoad %double %378 Aligned 8
        %380 = OpBitcast %_ptr_Workgroup_uchar %130
        %381 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %380 %ulong_424
        %382 = OpBitcast %_ptr_Workgroup_double %381
        %383 = OpLoad %double %382 Aligned 8
        %384 = OpBitcast %_ptr_Workgroup_uchar %130
        %385 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %384 %ulong_432
        %386 = OpBitcast %_ptr_Workgroup_double %385
        %387 = OpLoad %double %386 Aligned 8
        %388 = OpBitcast %_ptr_Workgroup_uchar %130
        %389 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %388 %ulong_440
        %390 = OpBitcast %_ptr_Workgroup_double %389
        %391 = OpLoad %double %390 Aligned 8
        %392 = OpBitcast %_ptr_Workgroup_uchar %130
        %393 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %392 %ulong_448
        %394 = OpBitcast %_ptr_Workgroup_double %393
        %395 = OpLoad %double %394 Aligned 8
        %396 = OpBitcast %_ptr_Workgroup_uchar %130
        %397 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %396 %ulong_456
        %398 = OpBitcast %_ptr_Workgroup_double %397
        %399 = OpLoad %double %398 Aligned 8
        %400 = OpBitcast %_ptr_Workgroup_uchar %130
        %401 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %400 %ulong_464
        %402 = OpBitcast %_ptr_Workgroup_double %401
        %403 = OpLoad %double %402 Aligned 8
        %404 = OpBitcast %_ptr_Workgroup_uchar %130
        %405 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %404 %ulong_472
        %406 = OpBitcast %_ptr_Workgroup_double %405
        %407 = OpLoad %double %406 Aligned 8
        %408 = OpBitcast %_ptr_Workgroup_uchar %130
        %409 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %408 %ulong_480
        %410 = OpBitcast %_ptr_Workgroup_double %409
        %411 = OpLoad %double %410 Aligned 8
        %412 = OpBitcast %_ptr_Workgroup_uchar %130
        %413 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %412 %ulong_488
        %414 = OpBitcast %_ptr_Workgroup_double %413
        %415 = OpLoad %double %414 Aligned 8
        %416 = OpBitcast %_ptr_Workgroup_uchar %130
        %417 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %416 %ulong_496
        %418 = OpBitcast %_ptr_Workgroup_double %417
        %419 = OpLoad %double %418 Aligned 8
        %420 = OpBitcast %_ptr_Workgroup_uchar %130
        %421 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %420 %ulong_504
        %422 = OpBitcast %_ptr_Workgroup_double %421
        %423 = OpLoad %double %422 Aligned 8
        %424 = OpBitcast %_ptr_Workgroup_uchar %130
        %425 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %424 %ulong_512
        %426 = OpBitcast %_ptr_Workgroup_double %425
        %427 = OpLoad %double %426 Aligned 8
        %428 = OpBitcast %_ptr_Workgroup_uchar %130
        %429 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %428 %ulong_520
        %430 = OpBitcast %_ptr_Workgroup_double %429
        %431 = OpLoad %double %430 Aligned 8
        %432 = OpBitcast %_ptr_Workgroup_uchar %130
        %433 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %432 %ulong_528
        %434 = OpBitcast %_ptr_Workgroup_double %433
        %435 = OpLoad %double %434 Aligned 8
        %436 = OpBitcast %_ptr_Workgroup_uchar %130
        %437 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %436 %ulong_536
        %438 = OpBitcast %_ptr_Workgroup_double %437
        %439 = OpLoad %double %438 Aligned 8
        %440 = OpBitcast %_ptr_Workgroup_uchar %130
        %441 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %440 %ulong_544
        %442 = OpBitcast %_ptr_Workgroup_double %441
        %443 = OpLoad %double %442 Aligned 8
        %444 = OpBitcast %_ptr_Workgroup_uchar %130
        %445 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %444 %ulong_552
        %446 = OpBitcast %_ptr_Workgroup_double %445
        %447 = OpLoad %double %446 Aligned 8
        %448 = OpBitcast %_ptr_Workgroup_uchar %130
        %449 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %448 %ulong_560
        %450 = OpBitcast %_ptr_Workgroup_double %449
        %451 = OpLoad %double %450 Aligned 8
        %452 = OpBitcast %_ptr_Workgroup_uchar %130
        %453 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %452 %ulong_568
        %454 = OpBitcast %_ptr_Workgroup_double %453
        %455 = OpLoad %double %454 Aligned 8
        %456 = OpBitcast %_ptr_Workgroup_uchar %130
        %457 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %456 %ulong_576
        %458 = OpBitcast %_ptr_Workgroup_double %457
        %459 = OpLoad %double %458 Aligned 8
        %460 = OpBitcast %_ptr_Workgroup_uchar %130
        %461 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %460 %ulong_584
        %462 = OpBitcast %_ptr_Workgroup_double %461
        %463 = OpLoad %double %462 Aligned 8
        %464 = OpBitcast %_ptr_Workgroup_uchar %130
        %465 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %464 %ulong_592
        %466 = OpBitcast %_ptr_Workgroup_double %465
        %467 = OpLoad %double %466 Aligned 8
        %468 = OpBitcast %_ptr_Workgroup_uchar %130
        %469 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %468 %ulong_600
        %470 = OpBitcast %_ptr_Workgroup_double %469
        %471 = OpLoad %double %470 Aligned 8
        %472 = OpBitcast %_ptr_Workgroup_uchar %130
        %473 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %472 %ulong_608
        %474 = OpBitcast %_ptr_Workgroup_double %473
        %475 = OpLoad %double %474 Aligned 8
        %476 = OpBitcast %_ptr_Workgroup_uchar %130
        %477 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %476 %ulong_616
        %478 = OpBitcast %_ptr_Workgroup_double %477
        %479 = OpLoad %double %478 Aligned 8
        %480 = OpBitcast %_ptr_Workgroup_uchar %130
        %481 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %480 %ulong_624
        %482 = OpBitcast %_ptr_Workgroup_double %481
        %483 = OpLoad %double %482 Aligned 8
        %484 = OpBitcast %_ptr_Workgroup_uchar %130
        %485 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %484 %ulong_632
        %486 = OpBitcast %_ptr_Workgroup_double %485
        %487 = OpLoad %double %486 Aligned 8
        %488 = OpBitcast %_ptr_Workgroup_uchar %130
        %489 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %488 %ulong_640
        %490 = OpBitcast %_ptr_Workgroup_double %489
        %491 = OpLoad %double %490 Aligned 8
        %492 = OpFAdd %double %141 %141
        %493 = OpFMul %double %141 %492
        %494 = OpFAdd %double %493 %double_1
        %495 = OpFMul %double %141 %494
        %496 = OpFAdd %double %495 %492
        %497 = OpFMul %double %141 %496
        %498 = OpFAdd %double %497 %494
        %499 = OpFMul %double %141 %498
        %500 = OpFAdd %double %499 %496
        %501 = OpFMul %double %141 %500
        %502 = OpFAdd %double %501 %498
        %503 = OpFMul %double %141 %502
        %504 = OpFAdd %double %503 %500
        %505 = OpSelect %double %142 %170 %156
        %506 = OpSelect %double %142 %168 %154
        %507 = OpSelect %double %142 %166 %152
        %508 = OpSelect %double %142 %164 %150
        %509 = OpSelect %double %142 %162 %148
        %510 = OpSelect %double %142 %160 %146
        %511 = OpSelect %double %142 %158 %144
        %512 = OpSelect %double %142 %143 %double_1
        %513 = OpSelect %double %142 %double_1 %137
        %514 = OpFMul %double %513 %171
        %515 = OpFAdd %double %514 %double_0
        %516 = OpFMul %double %512 %175
        %517 = OpFAdd %double %515 %516
        %518 = OpFMul %double %511 %179
        %519 = OpFAdd %double %517 %518
        %520 = OpFMul %double %510 %183
        %521 = OpFAdd %double %519 %520
        %522 = OpFMul %double %509 %187
        %523 = OpFAdd %double %521 %522
        %524 = OpFMul %double %508 %191
        %525 = OpFAdd %double %523 %524
        %526 = OpFMul %double %507 %195
        %527 = OpFAdd %double %525 %526
        %528 = OpFMul %double %506 %199
        %529 = OpFAdd %double %527 %528
        %530 = OpFMul %double %505 %203
        %531 = OpFAdd %double %529 %530
        %532 = OpFMul %double %513 %207
        %533 = OpFAdd %double %532 %double_0
        %534 = OpFMul %double %512 %211
        %535 = OpFAdd %double %533 %534
        %536 = OpFMul %double %511 %215
        %537 = OpFAdd %double %535 %536
        %538 = OpFMul %double %510 %219
        %539 = OpFAdd %double %537 %538
        %540 = OpFMul %double %509 %223
        %541 = OpFAdd %double %539 %540
        %542 = OpFMul %double %508 %227
        %543 = OpFAdd %double %541 %542
        %544 = OpFMul %double %507 %231
        %545 = OpFAdd %double %543 %544
        %546 = OpFMul %double %506 %235
        %547 = OpFAdd %double %545 %546
        %548 = OpFMul %double %505 %239
        %549 = OpFAdd %double %547 %548
        %550 = OpFMul %double %513 %243
        %551 = OpFAdd %double %550 %double_0
        %552 = OpFMul %double %512 %247
        %553 = OpFAdd %double %551 %552
        %554 = OpFMul %double %511 %251
        %555 = OpFAdd %double %553 %554
        %556 = OpFMul %double %510 %255
        %557 = OpFAdd %double %555 %556
        %558 = OpFMul %double %509 %259
        %559 = OpFAdd %double %557 %558
        %560 = OpFMul %double %508 %263
        %561 = OpFAdd %double %559 %560
        %562 = OpFMul %double %507 %267
        %563 = OpFAdd %double %561 %562
        %564 = OpFMul %double %506 %271
        %565 = OpFAdd %double %563 %564
        %566 = OpFMul %double %505 %275
        %567 = OpFAdd %double %565 %566
        %568 = OpFMul %double %513 %279
        %569 = OpFAdd %double %568 %double_0
        %570 = OpFMul %double %512 %283
        %571 = OpFAdd %double %569 %570
        %572 = OpFMul %double %511 %287
        %573 = OpFAdd %double %571 %572
        %574 = OpFMul %double %510 %291
        %575 = OpFAdd %double %573 %574
        %576 = OpFMul %double %509 %295
        %577 = OpFAdd %double %575 %576
        %578 = OpFMul %double %508 %299
        %579 = OpFAdd %double %577 %578
        %580 = OpFMul %double %507 %303
        %581 = OpFAdd %double %579 %580
        %582 = OpFMul %double %506 %307
        %583 = OpFAdd %double %581 %582
        %584 = OpFMul %double %505 %311
        %585 = OpFAdd %double %583 %584
        %586 = OpFMul %double %513 %315
        %587 = OpFAdd %double %586 %double_0
        %588 = OpFMul %double %512 %319
        %589 = OpFAdd %double %587 %588
        %590 = OpFMul %double %511 %323
        %591 = OpFAdd %double %589 %590
        %592 = OpFMul %double %510 %327
        %593 = OpFAdd %double %591 %592
        %594 = OpFMul %double %509 %331
        %595 = OpFAdd %double %593 %594
        %596 = OpFMul %double %508 %335
        %597 = OpFAdd %double %595 %596
        %598 = OpFMul %double %507 %339
        %599 = OpFAdd %double %597 %598
        %600 = OpFMul %double %506 %343
        %601 = OpFAdd %double %599 %600
        %602 = OpFMul %double %505 %347
        %603 = OpFAdd %double %601 %602
        %604 = OpFMul %double %513 %351
        %605 = OpFAdd %double %604 %double_0
        %606 = OpFMul %double %512 %355
        %607 = OpFAdd %double %605 %606
        %608 = OpFMul %double %511 %359
        %609 = OpFAdd %double %607 %608
        %610 = OpFMul %double %510 %363
        %611 = OpFAdd %double %609 %610
        %612 = OpFMul %double %509 %367
        %613 = OpFAdd %double %611 %612
        %614 = OpFMul %double %508 %371
        %615 = OpFAdd %double %613 %614
        %616 = OpFMul %double %507 %375
        %617 = OpFAdd %double %615 %616
        %618 = OpFMul %double %506 %379
        %619 = OpFAdd %double %617 %618
        %620 = OpFMul %double %505 %383
        %621 = OpFAdd %double %619 %620
        %622 = OpFMul %double %513 %387
        %623 = OpFAdd %double %622 %double_0
        %624 = OpFMul %double %512 %391
        %625 = OpFAdd %double %623 %624
        %626 = OpFMul %double %511 %395
        %627 = OpFAdd %double %625 %626
        %628 = OpFMul %double %510 %399
        %629 = OpFAdd %double %627 %628
        %630 = OpFMul %double %509 %403
        %631 = OpFAdd %double %629 %630
        %632 = OpFMul %double %508 %407
        %633 = OpFAdd %double %631 %632
        %634 = OpFMul %double %507 %411
        %635 = OpFAdd %double %633 %634
        %636 = OpFMul %double %506 %415
        %637 = OpFAdd %double %635 %636
        %638 = OpFMul %double %505 %419
        %639 = OpFAdd %double %637 %638
        %640 = OpFMul %double %513 %423
        %641 = OpFAdd %double %640 %double_0
        %642 = OpFMul %double %512 %427
        %643 = OpFAdd %double %641 %642
        %644 = OpFMul %double %511 %431
        %645 = OpFAdd %double %643 %644
        %646 = OpFMul %double %510 %435
        %647 = OpFAdd %double %645 %646
        %648 = OpFMul %double %509 %439
        %649 = OpFAdd %double %647 %648
        %650 = OpFMul %double %508 %443
        %651 = OpFAdd %double %649 %650
        %652 = OpFMul %double %507 %447
        %653 = OpFAdd %double %651 %652
        %654 = OpFMul %double %506 %451
        %655 = OpFAdd %double %653 %654
        %656 = OpFMul %double %505 %455
        %657 = OpFAdd %double %655 %656
        %658 = OpFMul %double %513 %459
        %659 = OpFAdd %double %658 %double_0
        %660 = OpFMul %double %512 %463
        %661 = OpFAdd %double %659 %660
        %662 = OpFMul %double %511 %467
        %663 = OpFAdd %double %661 %662
        %664 = OpFMul %double %510 %471
        %665 = OpFAdd %double %663 %664
        %666 = OpFMul %double %509 %475
        %667 = OpFAdd %double %665 %666
        %668 = OpFMul %double %508 %479
        %669 = OpFAdd %double %667 %668
        %670 = OpFMul %double %507 %483
        %671 = OpFAdd %double %669 %670
        %672 = OpFMul %double %506 %487
        %673 = OpFAdd %double %671 %672
        %674 = OpFMul %double %505 %491
        %675 = OpFAdd %double %673 %674
        %676 = OpFMul %double %141 %531
        %677 = OpFMul %double %492 %567
        %678 = OpFMul %double %494 %585
        %679 = OpFMul %double %496 %603
        %680 = OpFMul %double %498 %621
        %681 = OpFMul %double %500 %639
        %682 = OpFMul %double %502 %657
        %683 = OpFMul %double %504 %675
               OpBranch %684
        %684 = OpLabel
        %685 = OpPhi %uint %uint_0 %121 %686 %684
        %687 = OpPhi %double %double_0 %121 %688 %684
        %689 = OpFAdd %double %687 %676
        %690 = OpFAdd %double %689 %549
        %691 = OpFAdd %double %690 %677
        %692 = OpFAdd %double %691 %678
        %693 = OpFAdd %double %692 %679
        %694 = OpFAdd %double %693 %680
        %695 = OpFAdd %double %694 %681
        %696 = OpFAdd %double %695 %682
        %688 = OpFAdd %double %696 %683
        %686 = OpIAdd %uint %685 %uint_1
        %697 = OpULessThan %bool %685 %uint_2
               OpBranchConditional %697 %684 %698
        %698 = OpLabel
        %699 = OpConvertPtrToU %ulong %8
        %700 = OpConvertUToPtr %_ptr_Generic_double %699
        %701 = OpInBoundsPtrAccessChain %_ptr_Generic_double %700 %135
               OpStore %701 %688 Aligned 8
        %702 = OpInBoundsPtrAccessChain %_ptr_Generic_double %700 %139
               OpStore %702 %double_0 Aligned 8
               OpReturn
               OpFunctionEnd
          %2 = OpFunction %void None %108
          %9 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %10 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
        %703 = OpFunctionParameter %_ptr_Workgroup_double
        %704 = OpLabel
        %705 = OpFunctionCall %void %6 %9 %10 %703
               OpReturn
               OpFunctionEnd
```

</details>

Re-assemble with `spirv-as trigger.spvasm -o trigger.spv`.

## Host (pure Level Zero, no HIP)

See `repro.cpp` in the reproducer repo (104 LOC, `lze_loader` the only
runtime dep). Core loop:

```cpp
ze_module_desc_t md = {ZE_STRUCTURE_TYPE_MODULE_DESC, nullptr,
                       ZE_MODULE_FORMAT_IL_SPIRV,
                       spv_size, (const uint8_t*)spv.data(),
                       build_flags /* "" or "-cl-opt-disable" */, nullptr};
zeModuleCreate(ctx, dev, &md, &module, &log);
zeModuleGetKernelNames(module, &kc, names);  // <-- GradAtPoints not in names
```

## Ask

1. Confirm this is the same failure class as #403 (RetryManager exhausting
   all SIMD widths and silently dropping the kernel rather than emitting a
   compile error).
2. A finer-grained workaround than `-cl-opt-disable`. Ideally an IGC env var
   that disables only the function-inlining pass that produces the
   unregister-allocatable vISA at `Q_1D = 9`.
3. A real fix in IGC.

Happy to bisect further, rebuild IGC from source for pass-level logging, or
test patches.


Provide feedback

Saved searches

Use saved searches to filter your results more quickly

IGC silently drops kernel at -O2; SPIR-V well-formed, L0 build succeeds, zeKernelCreate fails with INVALID_KERNEL_NAME; -cl-opt-disable recovers (Arc A770, Q_1D=9 double kernel, libCEED GradAtPoints) #404

Summary

Environment

Reproducing

Trigger

IGC shader-dump evidence (`IGC_ShaderDumpEnable=1`)

IGC flags tried that do NOT help

Related previously-reported issues

The kernel (HIP source, 52 LOC)

SPIR-V (718 LOC, result of `clang -O2` + `llvm-spirv` + `spirv-opt --eliminate-dead-*` + line-bisection)

Host (pure Level Zero, no HIP)

Ask

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

IGC silently drops kernel at -O2; SPIR-V well-formed, L0 build succeeds, zeKernelCreate fails with INVALID_KERNEL_NAME; -cl-opt-disable recovers (Arc A770, Q_1D=9 double kernel, libCEED GradAtPoints) #404

Description

Summary

Environment

Reproducing

Trigger

IGC shader-dump evidence (IGC_ShaderDumpEnable=1)

IGC flags tried that do NOT help

Related previously-reported issues

The kernel (HIP source, 52 LOC)

SPIR-V (718 LOC, result of clang -O2 + llvm-spirv + spirv-opt --eliminate-dead-* + line-bisection)

Host (pure Level Zero, no HIP)

Ask

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions

IGC shader-dump evidence (`IGC_ShaderDumpEnable=1`)

SPIR-V (718 LOC, result of `clang -O2` + `llvm-spirv` + `spirv-opt --eliminate-dead-*` + line-bisection)