fix(shader): decode forced degamma on biased RGBA8_UINT samples by eljackjoestar-cmd · Pull Request #5472 · boykopovar/AnyPS5

eljackjoestar-cmd · 2026-10-13T20:39:21Z

What

After #5461, Lost Soul Aside (PPSA26358) reaches a rejection of FORCE_DEGAMMA on guest format 60, RGBA8_UINT. Hardware does not leave these samples as integers: setting the sampler bit makes the same bytes return decoded floating-point sRGB values.

For the measured direct, plain biased pixel-sampling case, materialize the image descriptor as RGBA8_SRGB before selecting the runtime image mode and descriptor heap. This changes both the shader's numeric mode and the sampled image view. The underlying guest descriptor/user-data bytes remain unchanged. Only a one-slice 3D image at mip zero is accepted, and every sampler paired with it must force degamma. The #5461 sampler proof still requires normalized point filtering. Mixed sampler states, other access forms, indirect images and other shapes remain refused.

This depends on #5461's bias-use proof, restricted 3D runtime mode and regression fixture. Its prerequisite commit is included until that PR merges; the new change is the final commit.

Semantics measurement

Direct native RDNA2 measurement on AMD integrated gfx1036 (Ryzen 7 7800X3D), AMD driver 32.0.21036.18, Windows 11. An OpenCL-loaded native kernel executes image_sample_b DIM 2D (f0940f08 00400000) against a 16×16×1 RGBA8_UINT 3D image at mip zero. The script below verifies the compiled instruction/register layout, removes coordinate preprocessing and patches the opcode/DIM and sampler flag before loading the native binary. No emulator shaders run in this measurement.

Controls: wave32; IEEE off; DX10 clamp on; f32/f16 denormals preserved; round to nearest even; FP16 overflow off. Twenty rows combine biases -16/-0.5/0/0.5/16 and four UV pairs, including wrapped coordinates. The RGBA8_UINT texel bytes are 64,128,32,255:

Without FORCE_DEGAMMA, the result words are integers 64,128,32,255.
With it, they are floating-point words 3d520000,3e5d0000,3c6d0000,3f800000, bit-identical to the same RGBA8_UNORM bytes sampled with FORCE_DEGAMMA.
This holds for point weighted-average/min/max reduction and the game's exact sampler words (20175700 03800000 00000000 a0500f2e): 140 RGBA8_UINT rows checked.
R32_UINT is a separating control: it remains integer-valued with the flag. It is not enabled by this change.

Wave64 and PS5 gfx1013 are not measured; the scope and gaps are recorded in shader debt. The Vulkan pixel readback compares decoded channels with a scalar sRGB oracle within one output UNORM code, retaining exact checks for the original non-degamma cases.

Tested

Rebased onto main 84b01f9cd042e1d2e9e112f8783aa5a9d26399ee on 2026-10-14; PR head 613bb8e90fd7ebe74d35c014016949de719e48c6. Standalone branch verification (#5461 plus this PR), Windows 11 Pro 10.0.26300, WinLibs GCC 15.2, Release, SPIRV-Tools OFF: runtime and selected native test targets built successfully. All 7 focused tests passed without skips in 3.03 s. Conventions and whitespace checks passed. The #5461 prerequisite is rebased onto the same main; it remains the only dependency. The cache format is 32, inherited from #5461. The Vulkan checks above used NVIDIA RTX 5070 / driver 617.14. Three image regressions also passed on AMD integrated gfx1036 / 32.0.21036.18 in 0.53 s, including the existing array-to-volume store regression. No new title/gameplay run, Linux, full-suite or CI result is claimed.

Earlier validation, on the bases stated below:

Full Windows Release build of main a0f5916a2 + fix(shader): sample single-level volumes through biased 2D accesses #5461 + this change alone, WinLibs GCC 15.2.
Six focused tests passed in 1.77 s: shader memory, disk-cache round trip, restricted volume bias/degamma readback, existing FORCE_DEGAMMA sampling, 2D-on-1D implicit LOD and 3D-on-2D addressing.
New readback passes on NVIDIA RTX 5070 (617.14) and AMD integrated gfx1036 (32.0.21036.18), including positive/negative bias, scalar decode comparison and rejection of mixed forced/unforced samplers.
Removing the conversion call makes the new regression fail with the same format-60 FORCE_DEGAMMA rejection; restoring it passes.
Fresh integrated Lost Soul Aside run on the synced main: two X presses pass the title; the previously failing image at 0x2017590000 is now bound as guest format 130 through a sampled sRGB view (vk=43). No renderer rejection occurs before the memory guard stops loading with 1,580 MiB free commit headroom (process private memory 16,943 MiB). Gameplay is still unverified. This integration includes fix(agc): preserve array slice XOR in color target views #4551, fix(shader): map retained SPIR-V from shader cache entries #4806, fix(shader): ignore unused descriptor fields for scalar buffer loads #5049, fix(audio): preserve completed AudioOut2 grains across threads #5198, point-size and depth-storage prerequisites plus the occlusion workaround; those are not new changes in this PR.
Full integration build and ten focused integration tests pass (4.62 s).
Convention and diff checks pass. Pairwise merge simulation with feat(agc): sample non-sRGB textures through FORCE_DEGAMMA samplers as sRGB views #4487 is clean. That PR creates sRGB views of UNORM formats and does not implement the UINT-to-floating-point runtime mode change; no combined runtime validation with it is claimed.

Checklist

Raw native measurements

{
  "device": "gfx1036",
  "driver": "32.0.21036.18",
  "rows": [
    [
      -16.0,
      0.03125,
      0.03125,
      0.0
    ],
    [
      -16.0,
      0.53125,
      0.78125,
      0.0
    ],
    [
      -16.0,
      0.96875,
      0.96875,
      0.0
    ],
    [
      -16.0,
      -0.03125,
      1.03125,
      0.0
    ],
    [
      -0.5,
      0.03125,
      0.03125,
      0.0
    ],
    [
      -0.5,
      0.53125,
      0.78125,
      0.0
    ],
    [
      -0.5,
      0.96875,
      0.96875,
      0.0
    ],
    [
      -0.5,
      -0.03125,
      1.03125,
      0.0
    ],
    [
      0.0,
      0.03125,
      0.03125,
      0.0
    ],
    [
      0.0,
      0.53125,
      0.78125,
      0.0
    ],
    [
      0.0,
      0.96875,
      0.96875,
      0.0
    ],
    [
      0.0,
      -0.03125,
      1.03125,
      0.0
    ],
    [
      0.5,
      0.03125,
      0.03125,
      0.0
    ],
    [
      0.5,
      0.53125,
      0.78125,
      0.0
    ],
    [
      0.5,
      0.96875,
      0.96875,
      0.0
    ],
    [
      0.5,
      -0.03125,
      1.03125,
      0.0
    ],
    [
      16.0,
      0.03125,
      0.03125,
      0.0
    ],
    [
      16.0,
      0.53125,
      0.78125,
      0.0
    ],
    [
      16.0,
      0.96875,
      0.96875,
      0.0
    ],
    [
      16.0,
      -0.03125,
      1.03125,
      0.0
    ]
  ],
  "controls": {
    "wave32": true,
    "ieee": 0,
    "dx10_clamp": 1,
    "denorm32": 3,
    "denorm16": 3,
    "round32": 0,
    "round16": 0,
    "fp16_overflow": 0
  },
  "cases": [
    {
      "format": "R32_UINT",
      "reduction": 0,
      "force_degamma": 0,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ]
      ]
    },
    {
      "format": "R32_UINT",
      "reduction": 0,
      "force_degamma": 1,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ]
      ]
    },
    {
      "format": "R32_UINT",
      "reduction": 1,
      "force_degamma": 0,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ]
      ]
    },
    {
      "format": "R32_UINT",
      "reduction": 1,
      "force_degamma": 1,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ]
      ]
    },
    {
      "format": "R32_UINT",
      "reduction": 2,
      "force_degamma": 0,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ]
      ]
    },
    {
      "format": "R32_UINT",
      "reduction": 2,
      "force_degamma": 1,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ]
      ]
    },
    {
      "format": "RGBA8_UNORM",
      "reduction": 0,
      "force_degamma": 0,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ],
        [
          1048608897,
          1056997505,
          1040220289,
          1065353216
        ]
      ]
    },
    {
      "format": "RGBA8_UNORM",
      "reduction": 0,
      "force_degamma": 1,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ]
      ]
    },
    {
      "format": "R32_UINT",
      "reduction": 1,
      "force_degamma": 1,
      "exact_game_sampler": true,
      "baseline_lz_dim3d": false,
      "output": [
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16780296,
          0,
          0,
          1
        ],
        [
          16781071,
          0,
          0,
          1
        ],
        [
          16777231,
          0,
          0,
          1
        ]
      ]
    },
    {
      "format": "RGBA8_UINT",
      "reduction": 0,
      "force_degamma": 0,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ]
      ]
    },
    {
      "format": "RGBA8_UINT",
      "reduction": 0,
      "force_degamma": 1,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ]
      ]
    },
    {
      "format": "RGBA8_UINT",
      "reduction": 1,
      "force_degamma": 0,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ]
      ]
    },
    {
      "format": "RGBA8_UINT",
      "reduction": 1,
      "force_degamma": 1,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ]
      ]
    },
    {
      "format": "RGBA8_UINT",
      "reduction": 2,
      "force_degamma": 0,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ],
        [
          64,
          128,
          32,
          255
        ]
      ]
    },
    {
      "format": "RGBA8_UINT",
      "reduction": 2,
      "force_degamma": 1,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": false,
      "output": [
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ]
      ]
    },
    {
      "format": "RGBA8_UINT",
      "reduction": 1,
      "force_degamma": 1,
      "exact_game_sampler": true,
      "baseline_lz_dim3d": false,
      "output": [
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ],
        [
          1028784128,
          1046282240,
          1013776384,
          1065353216
        ]
      ]
    },
    {
      "format": "R32_UINT",
      "reduction": 0,
      "force_degamma": 0,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": true,
      "output": [
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16779264,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16777224,
          0,
          0,
          1
        ],
        [
          16779272,
          0,
          0,
          1
        ],
        [
          16781064,
          0,
          0,
          1
        ],
        [
          16781064,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16779264,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16777224,
          0,
          0,
          1
        ],
        [
          16779272,
          0,
          0,
          1
        ],
        [
          16781064,
          0,
          0,
          1
        ],
        [
          16781064,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16779264,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ]
      ]
    },
    {
      "format": "R32_UINT",
      "reduction": 0,
      "force_degamma": 1,
      "exact_game_sampler": false,
      "baseline_lz_dim3d": true,
      "output": [
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16779264,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16777224,
          0,
          0,
          1
        ],
        [
          16779272,
          0,
          0,
          1
        ],
        [
          16781064,
          0,
          0,
          1
        ],
        [
          16781064,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16779264,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16777224,
          0,
          0,
          1
        ],
        [
          16779272,
          0,
          0,
          1
        ],
        [
          16781064,
          0,
          0,
          1
        ],
        [
          16781064,
          0,
          0,
          1
        ],
        [
          16777216,
          0,
          0,
          1
        ],
        [
          16779264,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ],
        [
          16781056,
          0,
          0,
          1
        ]
      ]
    }
  ]
}

Reproduction script (Windows OpenCL, gfx1036)

import ctypes as c
import json
import struct
from pathlib import Path

cl = c.WinDLL('OpenCL.dll')
P = c.c_void_p
U = c.c_uint
I = c.c_int
Z = c.c_size_t
Q = c.c_ulonglong

def bind(name, ret, args):
    f = getattr(cl, name)
    f.restype = ret
    f.argtypes = args
    return f

platforms = bind('clGetPlatformIDs', I, [U, P, P])
devices = bind('clGetDeviceIDs', I, [P, Q, U, P, P])
device_info = bind('clGetDeviceInfo', I, [P, U, Z, P, P])
create_context = bind('clCreateContext', P, [P, U, P, P, P, P])
create_source = bind('clCreateProgramWithSource', P, [P, U, P, P, P])
build = bind('clBuildProgram', I, [P, U, P, c.c_char_p, P, P])
program_info = bind('clGetProgramInfo', I, [P, U, Z, P, P])
build_info = bind('clGetProgramBuildInfo', I, [P, P, U, Z, P, P])
release_program = bind('clReleaseProgram', I, [P])
release_context = bind('clReleaseContext', I, [P])

def check(result, label):
    if result != 0:
        raise RuntimeError(f'{label}: OpenCL error {result}')

count = U()
check(platforms(0, None, c.byref(count)), 'platform count')
platform_ids = (P * count.value)()
check(platforms(count, platform_ids, None), 'platforms')
device = None
for platform in platform_ids:
    n = U()
    check(devices(platform, 0xffffffff, 0, None, c.byref(n)), 'device count')
    ids = (P * n.value)()
    check(devices(platform, 0xffffffff, n, ids, None), 'devices')
    for candidate in ids:
        name = c.create_string_buffer(1024)
        check(device_info(candidate, 0x102b, len(name), name, None), 'device name')
        if name.value == b'gfx1036':
            device = P(candidate)
if not device:
    raise RuntimeError('gfx1036 is not available')
err = I()
context = create_context(None, 1, c.byref(device), None, None, c.byref(err))
check(err.value, 'context')
source = b'''
__kernel void probe(read_only image3d_t image, __global const float4* input, __global uint4* output) {
    const uint id = get_global_id(0);
    const float4 coord = input[id];
    output[id] = read_imageui(image, CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_REPEAT | CLK_FILTER_NEAREST, coord);
}
'''
source_ptr = c.c_char_p(source)
source_len = Z(len(source))
program = create_source(context, 1, c.byref(source_ptr), c.byref(source_len), c.byref(err))
check(err.value, 'program')
result = build(program, 1, c.byref(device), b'-cl-std=CL2.0', None, None)
log_size = Z()
check(build_info(program, device, 0x1183, 0, None, c.byref(log_size)), 'build log size')
log = c.create_string_buffer(log_size.value)
check(build_info(program, device, 0x1183, len(log), log, None), 'build log')
print(log.value.decode())
check(result, 'build')
binary_size = Z()
check(program_info(program, 0x1165, c.sizeof(binary_size), c.byref(binary_size), None), 'binary size')
binary = c.create_string_buffer(binary_size.value)
binary_ptr = c.cast(binary, P)
check(program_info(program, 0x1166, c.sizeof(binary_ptr), c.byref(binary_ptr), None), 'binary')
path = Path(__file__).with_name('amd-opencl-texture-original.bin')
path.write_bytes(binary.raw)
print('binary', binary_size.value, binary.raw[:16].hex())
check(release_program(program), 'release program')
check(release_context(context), 'release context')

create_binary = bind('clCreateProgramWithBinary', P, [P, U, P, P, P, P, P])
create_queue = bind('clCreateCommandQueue', P, [P, P, Q, P])
create_buffer = bind('clCreateBuffer', P, [P, Q, Z, P, P])
create_image = bind('clCreateImage', P, [P, Q, P, P, P, P])
create_kernel = bind('clCreateKernel', P, [P, c.c_char_p, P])
set_arg = bind('clSetKernelArg', I, [P, U, Z, P])
enqueue = bind('clEnqueueNDRangeKernel', I, [P, P, U, P, P, P, U, P, P])
read_buffer = bind('clEnqueueReadBuffer', I, [P, P, U, Z, Z, P, U, P, P])
finish = bind('clFinish', I, [P])
release_mem = bind('clReleaseMemObject', I, [P])
release_kernel = bind('clReleaseKernel', I, [P])
release_queue = bind('clReleaseCommandQueue', I, [P])

class ImageFormat(c.Structure):
    _fields_ = [('order', U), ('type', U)]

class ImageDescriptor(c.Structure):
    _fields_ = [('type', U), ('width', Z), ('height', Z), ('depth', Z), ('arraySize', Z), ('rowPitch', Z), ('slicePitch', Z), ('mips', U), ('samples', U), ('buffer', P)]

native = bytearray(path.read_bytes())
sections_offset = struct.unpack_from('<Q', native, 40)[0]
section_size, section_count, names_index = struct.unpack_from('<HHH', native, 58)
headers = [struct.unpack_from('<IIQQQQIIQQ', native, sections_offset + i * section_size) for i in range(section_count)]
names_header = headers[names_index]
names = native[names_header[4]:names_header[4] + names_header[5]]
sections = {names[s[0]:].split(b'\0', 1)[0].decode(): s for s in headers}
text_start, text_size = sections['.text'][4:6]
matches = [o for o in range(text_start, text_start + text_size, 4) if struct.unpack_from('<I', native, o)[0] == 0xf09c0f10]
if len(matches) != 1:
    raise RuntimeError(f'expected one 3D image_sample_lz, found {matches}')
sample_offset = matches[0]
if struct.unpack_from('<I', native, sample_offset + 4)[0] != 0x00400000:
    raise RuntimeError('unexpected image/sampler/address register encoding')
expected_init = {0x9ac: 0xbe8a0b84, 0x9b0: 0xbe8903ff, 0x9b4: 0x00fff000, 0x9c4: 0xbe8803ff, 0x9c8: 0x10003000}
if sample_offset != 0x9f0:
    raise RuntimeError('compiler layout changed')
for offset, value in expected_init.items():
    if struct.unpack_from('<I', native, offset)[0] != value:
        raise RuntimeError(f'compiler sampler initialization changed at {offset:x}')
for offset in range(0x9a0, sample_offset, 4):
    if offset not in expected_init:
        struct.pack_into('<I', native, offset, 0xbf800000)
struct.pack_into('<I', native, sample_offset, 0xf0940f08)
descriptor_offset = sections['.rodata'][4]
rsrc1 = struct.unpack_from('<I', native, descriptor_offset + 48)[0]
mask = (0xff << 12) | (1 << 21) | (1 << 23) | (1 << 26)
rsrc1 = (rsrc1 & ~mask) | (0xf0 << 12) | (1 << 21)
struct.pack_into('<I', native, descriptor_offset + 48, rsrc1)
properties = struct.unpack_from('<H', native, descriptor_offset + 56)[0]
print('native controls', hex(rsrc1), hex(properties), 'wave32', bool(properties & (1 << 10)))

size = 16
texels = (U * (size * size))(*(0x1000000 + y * 256 + x for y in range(size) for x in range(size)))
rows = [(bias, u, v, 0.0) for bias in (-16.0, -0.5, 0.0, 0.5, 16.0) for u, v in ((0.03125, 0.03125), (0.53125, 0.78125), (0.96875, 0.96875), (-0.03125, 1.03125))]
inputs = (c.c_float * (len(rows) * 4))(*(v for row in rows for v in row))
report = {'device': 'gfx1036', 'driver': '32.0.21036.18', 'rows': rows, 'controls': {'wave32': bool(properties & (1 << 10)), 'ieee': 0, 'dx10_clamp': 1, 'denorm32': 3, 'denorm16': 3, 'round32': 0, 'round16': 0, 'fp16_overflow': 0}, 'cases': []}
cases = [(r, d, "R32_UINT", False) for r in (0, 1, 2) for d in (0, 1)] + [(0, d, "RGBA8_UNORM", False) for d in (0, 1)] + [(1, 1, "R32_UINT", True)]
cases += [(r, d, "RGBA8_UINT", False) for r in (0, 1, 2) for d in (0, 1)]
cases.append((1, 1, "RGBA8_UINT", True))
cases.append((0, 0, "R32_UINT", False))
cases.append((0, 1, "R32_UINT", False))
for case_index, (reduction, degamma, format_name, exact_sampler) in enumerate(cases):
    data = bytearray(native)
    baseline = case_index >= len(cases) - 2
    if baseline:
        struct.pack_into("<I", data, sample_offset, 0xf09c0f10)
    struct.pack_into('<I', data, 0x9c8, 0x10003000 | (reduction << 29) | (degamma << 20))
    if exact_sampler:
        struct.pack_into('<I', data, 0x9c8, 0x20175700)
        struct.pack_into('<I', data, 0x9b4, 0x03800000)
        struct.pack_into('<I', data, 0x9ac, 0xbe8a0b80)
        struct.pack_into('<II', data, 0x9e4, 0xbe8b03ff, 0xa0500f2e)
    case_path = path.with_name(f'amd-opencl-bias-dim2d-{format_name}-r{reduction}-d{degamma}-exact{int(exact_sampler)}-baseline{int(baseline)}.bin')
    case_path.write_bytes(data)
    context = create_context(None, 1, c.byref(device), None, None, c.byref(err))
    check(err.value, 'context')
    data_buffer = c.create_string_buffer(bytes(data))
    data_pointer = c.cast(data_buffer, P)
    data_size = Z(len(data))
    binary_status = I()
    program = create_binary(context, 1, c.byref(device), c.byref(data_size), c.byref(data_pointer), c.byref(binary_status), c.byref(err))
    check(err.value, 'load native binary')
    check(binary_status.value, 'native binary status')
    check(build(program, 1, c.byref(device), None, None, None), 'build native binary')
    queue = create_queue(context, device, 0, c.byref(err))
    check(err.value, 'queue')
    image_format = ImageFormat(0x10b0, 0x10dc) if format_name == "R32_UINT" else ImageFormat(0x10b5, 0x10da if format_name == "RGBA8_UINT" else 0x10d2)
    case_texels = texels if format_name == "R32_UINT" else (U * (size * size))(*(0xff208040 for _ in range(size * size)))
    image_descriptor = ImageDescriptor(0x10f2, size, size, 1, 0, 0, 0, 0, 0, None)
    image = P(create_image(context, 4 | 32, c.byref(image_format), c.byref(image_descriptor), case_texels, c.byref(err)))
    check(err.value, 'image')
    input_buffer = P(create_buffer(context, 4 | 32, c.sizeof(inputs), inputs, c.byref(err)))
    check(err.value, 'input')
    output = (U * (len(rows) * 4))()
    output_buffer = P(create_buffer(context, 1, c.sizeof(output), None, c.byref(err)))
    check(err.value, 'output')
    kernel = create_kernel(program, b'probe', c.byref(err))
    check(err.value, 'kernel')
    for index, value in enumerate((image, input_buffer, output_buffer)):
        check(set_arg(kernel, index, c.sizeof(value), c.byref(value)), 'argument')
    global_size = Z(len(rows))
    check(enqueue(queue, kernel, 1, None, c.byref(global_size), None, 0, None, None), 'enqueue')
    check(read_buffer(queue, output_buffer, 1, 0, c.sizeof(output), output, 0, None, None), 'read output')
    check(finish(queue), 'finish')
    results = [list(output[i:i + 4]) for i in range(0, len(output), 4)]
    print('measurement', format_name, reduction, degamma, 'exact', exact_sampler, 'baseline', baseline, results)
    report['cases'].append({'format': format_name, 'reduction': reduction, 'force_degamma': degamma, 'exact_game_sampler': exact_sampler, 'baseline_lz_dim3d': baseline, 'output': results})
    for memory in (image, input_buffer, output_buffer):
        check(release_mem(memory), 'release memory')
    check(release_kernel(kernel), 'release kernel')
    check(release_queue(queue), 'release queue')
    check(release_program(program), 'release program')
    check(release_context(context), 'release context')
path.with_name('amd-opencl-texture-measurement.json').write_text(json.dumps(report, indent=2), encoding='utf-8')

* Metal semantics: run upstream image execution tests through the Metal driver The existing semantics shim binds only guest buffers. A second shim (tests/metal_semantics/driver_shim) runs the upstream execution tests that bind images, samplers and color targets through the native Metal driver: dispatches bind their descriptors with MetalShaderResources and draws go through MetalDraw::DrawSynchronously, as in a title. The tests' host arrays are the guest memory; the device borrows every range the shader's buffer and image descriptors and the color targets name. ImagePlaneAddressOnLine and ImageVolumeAddressOnPlane (already on main) now run on Metal and pass. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DgW4n3qYgi5jAAohEmXRQj * fix(shader): capture image descriptors read from constant buffers at uniform offsets Removes the tracker check that refused a T# whose dwords come from a ReadConstBuffer, so a T# loaded with s_buffer_load from a table record that user data selects is captured like any other runtime descriptor. Metal: tests/metal_semantics/ImageTableRecord.cpp runs the upstream program through the Metal driver (MetalShaderResources image binding) and checks that the record user data selects is the texture sampled. It fails on main with the removed "is not a valid runtime value" error. Ported from upstream PR boykopovar#5188 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DgW4n3qYgi5jAAohEmXRQj * fix(shader): accept TYPE=0 images during materialization A T# with TYPE=0 is null on AMD hardware even when its base and other words are nonzero; the materializer now accepts it as a null image instead of refusing it as invalid (reserved TYPE=1 is still refused). Upstream's frontend test is ported to tests/ShaderMemory.cpp. Metal: the Metal image binding already treats TYPE=0 as null (IsNullTextureResource), so such an image binds no texture. The Metal ImageTableRecord test now also samples a TYPE=0 record with a nonzero base and reads zero; on main it fails with "runtime image descriptor is invalid". Debt note added to docs/dev/TechnicalDebt.md. Ported from upstream PR boykopovar#5473 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DgW4n3qYgi5jAAohEmXRQj * fix(shader): read 1D T#s through 3D image_sample and image_load accesses A DIM 3D image_sample_lz, plain image_sample or mip-less image_load whose T# is 1D gets a runtime image mode that binds the 1D texture: the sample reads the texel u selects, and the load reads the x texel when z is 0 and zero otherwise, as measured on RDNA1. Other 3D accesses of such an image still throw. ImageResource gains flatVolumeLineCompatible, so the shader disk cache FormatVersion goes from 35 to 36. This is the only bump in this series: the later image-mode fields reuse it. Metal: the 1D mode binds through the existing Image1D shape (a 2D texture of height 1) with no new view. Upstream's ImageVolumeAddressOnLine execution test runs unchanged through the Metal driver (anyps5_metal_driver_semantics_imagevolumeaddressonline). It fails without the materializer hunk ("incompatible with the static runtime image interface"). Debt note added to docs/dev/TechnicalDebt.md. Ported from upstream PR boykopovar#5375 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DgW4n3qYgi5jAAohEmXRQj * fix(shader): sample single-level volumes through biased 2D accesses A pixel-stage, direct, plain 32-bit image_sample_b with DIM 2D whose T# is a one-slice, single-level 3D image gets the 3D runtime image mode: the sample supplies a zero third coordinate and keeps the bias, as measured on RDNA2. The binding requires a normalized point sampler and the descriptor one slice and mip zero; other cases throw. ImageResource gains singleMipPlaneBiasCompatible; it is encoded by the shader disk cache under the FormatVersion 36 bump of this series (upstream's own bump to 31 is not taken). Metal: the 3D mode binds the volume as the Image3D shape it already has, with no new view. Upstream's ImagePlaneBiasOnSingleLevelVolume execution test (a pixel draw) runs unchanged through MetalDraw (anyps5_metal_driver_semantics_imageplanebiasonsinglelevelvolume). It fails without the tracker hunk ("incompatible with the static runtime image interface"). Debt note added to docs/dev/TechnicalDebt.md. Ported from upstream PR boykopovar#5461 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DgW4n3qYgi5jAAohEmXRQj * fix(shader): decode forced degamma on biased RGBA8_UINT samples A direct RGBA8_UINT image sampled only by plain biased pixel accesses (the one-slice volume case of boykopovar#5461) through samplers that all set FORCE_DEGAMMA is materialized as RGBA8_SRGB, which returns the same decoded floats as the hardware. Mixed sampler states, indirect images, other accesses and formats still throw. Metal: MetalSampler refused every FORCE_DEGAMMA S#, so this (and the existing sRGB case of ImageSampleForceDegamma) could not bind. Metal now does what the Vulkan driver does: MetalShaderResources::Bindings creates the sampler with forceDegammaPaired and refuses a sampled image of a non-sRGB format paired with a forcing sampler (the same message as RequireDegammaFormat). The rewritten descriptor binds through the existing sRGB pixel-format view. Upstream's extended ImagePlaneBiasOnSingleLevelVolume and the existing ImageSampleForceDegamma tests now pass on Metal; both failed with "forces sRGB decoding which is not implemented". Debt note added to docs/dev/TechnicalDebt.md. Ported from upstream PR boykopovar#5472 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DgW4n3qYgi5jAAohEmXRQj * fix(shader): read cube T#s through DIM 2D image_sample and image_load as face 0 A DIM 2D image_sample_lz, plain image_sample, image_sample_l or mip-less image_load whose T# is a cube gets the cube runtime image mode and reads the T#'s BASE_ARRAY face with face coordinate 0, as measured on RDNA1. Other DIM 2D accesses through a cube T# still throw. ImageResource gains cubeFaceCompatible; encoded under the series' FormatVersion 36 bump. Metal: a cube T# is already a 2D-array texture of faces and the cube mode binds the Image2DArray shape, so no new view is needed. Upstream's ImagePlaneAddressOnCube execution test runs unchanged through the Metal driver (anyps5_metal_driver_semantics_imageplaneaddressoncube) and fails without the materializer hunk. Debt note added to docs/dev/TechnicalDebt.md. Ported from upstream PR boykopovar#5509 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DgW4n3qYgi5jAAohEmXRQj * fix(shader): compare through 2D image_sample_c on 1D T#s The flat-line (2D access through a 1D T#) compatibility now ignores the compare flag, so image_sample_c and image_sample_c_lz with DIM 2D on a 1D T# take the 1D mode like their non-compare forms. Metal: the 1D mode binds through the existing Image1D shape (a 2D texture of height 1), with no new view. Upstream's ImagePlaneCompareOnLine execution test runs unchanged through the Metal driver (anyps5_metal_driver_semantics_imageplanecompareonline) and fails without the tracker hunk. The debt note in docs/dev/TechnicalDebt.md is updated as upstream's. Ported from upstream PR boykopovar#5456 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DgW4n3qYgi5jAAohEmXRQj * fix(shader): apply the V# DST_SEL to formatted buffer stores buffer_store_format_* (D16 included) wrote memory component i from data VGPR i whatever the V#'s DST_SEL. A formatted store now writes each format component from the slot its selector names, only for the DST_SELs measured on RDNA2 or covered by conformance tests: the identity, Z,Y,X,W on 8_8_8_8_UNORM and 16_16_16_16_UINT, and 0,0,0,X on 8_UNORM and 8_UINT. Every other DST_SEL on a formatted store is refused: a host V# throws in DescriptorBindingBuilder::Prepare, and a GPU-selected V# records an invalid-descriptor BDA fault and skips the store. BufferResource gains formattedWritten; encoded under the series' FormatVersion 36 bump. Metal: the store is lowered in the shared SPIR-V (SpirvRuntimeBuffer), so no Metal binding change. Upstream's BufferStoreFormatDstSel execution test runs unchanged through the Metal driver (anyps5_metal_driver_semantics_bufferstoreformatdstsel). Without the SpirvRuntimeBuffer hunk it fails (Z,Y,X,W on 8_8_8_8_UNORM writes 0x4033bfff, expected 0x40ffbf33). Debt note added to docs/dev/TechnicalDebt.md. Ported from upstream PR boykopovar#5172 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DgW4n3qYgi5jAAohEmXRQj --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

adriadam10 · 2026-10-14T21:33:38Z

Measured on RDNA2 (gfx1035, wave32, IEEE=0): image_sample_b DIM 2D over all 256 codes per channel, on 2D and one-slice 3D T#s, biases -16/0/16, clamp/wrap/min/max and the title's S#: RGBA8_UINT with FORCE_DEGAMMA returns the RGBA8_SRGB result bit for bit (0 differing rows), integers without the bit. The new readback passes on RADV and CI's lavapipe and fails without the change.

#5461 is merged, so I rebased the final commit alone onto main: the debt line now sits next to #5461's, materializeTables lost its tables argument, and indirectRoot != ImageResource::NoIndirectImage became table != NoTable (main renamed it, as in the merged #5461). Merging once the full build is green.

github-actions Bot added shader core/shader docs agc libSceAgc* needs rebase Conflicts with main, rebase needed labels Oct 13, 2026

eljackjoestar-cmd force-pushed the codex/rgba8-uint-force-degamma branch from c25423a to e2101fc Compare October 14, 2026 01:29

github-actions Bot mentioned this pull request Oct 14, 2026

fix(agc): accept 2D-array image loads and stores on 3D T#s #5555

Merged

11 tasks done

github-actions Bot removed the needs rebase Conflicts with main, rebase needed label Oct 14, 2026

Celegans12 mentioned this pull request Oct 14, 2026

feat(agc): implement image and sampler tables indexed at runtime offsets #757

Merged

11 tasks done

github-actions Bot mentioned this pull request Oct 14, 2026

feat(agc): implement image tables loaded through 64-bit pointers #758

Merged

11 tasks done

eljackjoestar-cmd mentioned this pull request Oct 14, 2026

perf(shader): keep only evaluable branch conditions in resource plans #5181

Merged

10 tasks done

RINNECODER mentioned this pull request Oct 14, 2026

Upstream PR batch 1 U8a: Metal image-mode views RINNECODER/AnyPS5#485

Merged

github-actions Bot added the needs rebase Conflicts with main, rebase needed label Oct 14, 2026

eljackjoestar-cmd force-pushed the codex/rgba8-uint-force-degamma branch from e2101fc to b7855ec Compare October 14, 2026 09:53

github-actions Bot added needs rebase Conflicts with main, rebase needed and removed needs rebase Conflicts with main, rebase needed labels Oct 14, 2026

eljackjoestar-cmd force-pushed the codex/rgba8-uint-force-degamma branch from b7855ec to 613bb8e Compare October 14, 2026 15:26

github-actions Bot added needs rebase Conflicts with main, rebase needed and removed needs rebase Conflicts with main, rebase needed labels Oct 14, 2026

adriadam10 self-assigned this Oct 14, 2026

fix(shader): decode forced degamma on biased RGBA8_UINT samples

76eb60a

adriadam10 force-pushed the codex/rgba8-uint-force-degamma branch from 613bb8e to 76eb60a Compare October 14, 2026 21:33

github-actions Bot removed the needs rebase Conflicts with main, rebase needed label Oct 14, 2026

adriadam10 merged commit 596a0d7 into boykopovar:main Oct 14, 2026
8 checks passed

RINNECODER mentioned this pull request Oct 15, 2026

Upstream sync 11D: shader recompiler (13 commits) RINNECODER/AnyPS5#511

Merged

Navigation Menu

Repository navigation

fix(shader): decode forced degamma on biased RGBA8_UINT samples - #5472

fix(shader): decode forced degamma on biased RGBA8_UINT samples#5472
adriadam10 merged 1 commit into
boykopovar:mainfrom
eljackjoestar-cmd:codex/rgba8-uint-force-degamma

eljackjoestar-cmd commented Oct 13, 2026 •

edited

Loading

Uh oh!

adriadam10 commented Oct 14, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

eljackjoestar-cmd commented Oct 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What

Semantics measurement

Tested

Checklist

Uh oh!

adriadam10 commented Oct 14, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

eljackjoestar-cmd commented Oct 13, 2026 •

edited

Loading