diff options
| author | Arthur Hanson Rasmusson <arthur@vgpu.io> | 2026-10-19 09:27:19 -0400 |
|---|---|---|
| committer | Arthur Hanson Rasmusson <arthur@vgpu.io> | 2026-10-19 09:27:19 -0400 |
| commit | 8f953e86e73a621b1d28fd39d4517ba2de538f65 (patch) | |
| tree | 12991f9b67a028ce4e2bc16161c2601b940656a2 /legacy-driver-guest/kernel-tests | |
| parent | 44132d7887222a300dcc155e0376f445dca41655 (diff) | |
| download | open-iov-8f953e86e73a621b1d28fd39d4517ba2de538f65.tar.gz open-iov-8f953e86e73a621b1d28fd39d4517ba2de538f65.zip | |
cuda-ab: blk512 w256 kernel index fix
The committed w256 kernel mixed element and block strides
(bx*1024 + tx*4) and could not cover the buffer. Replace with the
well-formed index i = (blockIdx.y*16 + threadIdx.y)*256 +
blockIdx.x*16 + threadIdx.x, writing out[i..i+3] = i..i+3.
Verified on the emulated vGPU: w512 bad=0; w256 bad=0 with all 16
blocks zero-bad=256/full-bad=0; span sweep firstbad=-1, incl. the
page-15 tail (16376..16383). Two consecutive boots green. This
isolates the earlier w256 bad=12288 to the test kernel itself and
retires blk512 as an index-math control; A-read (arow) failures
remain tracked via ctaid_probe PROBE2.
Diffstat (limited to 'legacy-driver-guest/kernel-tests')
0 files changed, 0 insertions, 0 deletions