| Commit message (Collapse) | Author | Age | |
|---|---|---|---|
| * | cuda-ab: blk512 w256 kernel index fix | Arthur Hanson Rasmusson | 44 hours |
| | | | | | | | | | | | | | | | 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. | ||
| * | cuda-ab: blk512 - minimal 2-D-grid CTA band-bug reproducer | Arthur Hanson Rasmusson | 44 hours |
| 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 | |||
 |
index : open-iov.git | |
| Open-IOV vGPU-IO: host-side GSP emulation for unmodified NVIDIA guest drivers | root |
| summaryrefslogtreecommitdiffstats |