Open-IOV vGPU-IO: host-side GSP emulation for unmodified NVIDIA guest drivers https://git.vgpu.io/open-iov.git/atom?h=bigpath 2026-10-19T13:27:19+00:00 2026-10-19T13:27:19+00:00 Arthur Hanson Rasmusson arthur@vgpu.io 2026-10-19T13:27:19+00:00 urn:sha1:8f953e86e73a621b1d28fd39d4517ba2de538f65 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. 2026-10-19T12:55:48+00:00 Arthur Hanson Rasmusson arthur@vgpu.io 2026-10-19T12:55:48+00:00 urn:sha1:44132d7887222a300dcc155e0376f445dca41655 w512<<<8,512>>> writes 16384 ints out[i]=i perfectly (bad=0 across repeated runs); w256<<<(16,4),(16,16)>>> with the same write does bad=12288 deterministically: rows 0..15 correct, rows >= 16 wrong. This is the transformer/row-band failure without any probe machinery, and the host scan reports the engine wrote 31/31 tracked pages for the failing launch, i.e. the by>=1 stores exist but do not reach the pages the client reads. Build: nvcc -O2 --cudart=static -arch=sm_75 -o blk512 blk512.cu