From 7fdbe8713c8cf746cd1a15e93e0de39d333905c7 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Wed, 23 Sep 2026 06:31:44 -0400 Subject: kernel: execute pushed CE work (PTE memsets, copies) and track semaphores The progress service now decodes push buffers as method streams and performs the work the driver pushes, not just the tracking semaphore: - C5B5 method decoder: SET_REMAP_CONST_A/B + COMPONENTS + OFFSET_OUT + LINE_LENGTH_IN + LAUNCH_DMA = memset (into vidmem or guest RAM); OFFSET_IN + OFFSET_OUT + LAUNCH_DMA = copy; SET_SEMAPHORE_A/B/PAYLOAD + LAUNCH_DMA = release. Contiguous method runs (count 2/4) are applied in full - OFFSET_IN/OUT arrive as one run. - Pushes span page boundaries, so pages are decoded as method streams without a start marker. Collected operations execute only when the push's release is reached and still pending (payload ahead of the semaphore), so recycled buffers replaying old pushes are skipped and their stale memsets cannot overwrite newer state. - VA->PA map is first-wins: stale copies of freed allocations and unrelated adjacent u64 pairs produced bogus entries that hijacked a good mapping and sent semaphore writes to the wrong page. - dump_guest module parameter (hex dump of guest PA, kthread via completion; kthread_stop before the thread runs loses the work). With this the guest no longer stalls in uvm_tracker_wait and no longer oopses in the UVM path; the failure moved to a userspace segfault inside libcuda (NULL-0x38) after the driver's page allocations. Also add tools/cuda-ab: run one deterministic CUDA workload on the host (bare metal oracle) and in the guest, then compare the stable result lines. The binary links the CUDA runtime statically so the guest only needs the driver's libcuda.so. --- kernel/gsp_sema_service.c | 487 ++++++++++++++++++++++++++++++++++++++++++- kernel/main.c | 89 ++++++++ tools/cuda-ab/README.md | 32 +++ tools/cuda-ab/ab.sh | 88 ++++++++ tools/cuda-ab/run/host.txt | 3 + tools/cuda-ab/run/vector_add | Bin 0 -> 820856 bytes tools/cuda-ab/vector_add.cu | 93 +++++++++ 7 files changed, 789 insertions(+), 3 deletions(-) create mode 100644 tools/cuda-ab/README.md create mode 100755 tools/cuda-ab/ab.sh create mode 100644 tools/cuda-ab/run/host.txt create mode 100755 tools/cuda-ab/run/vector_add create mode 100644 tools/cuda-ab/vector_add.cu diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index 68a890c..1ed8c99 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -82,6 +82,13 @@ struct sema_svc { /* stats */ u64 pairs_seen; + u64 pages_decoded; + u64 pushes_seen; + u64 pushes_pending; + u64 pushes_nomap; + u64 memsets_done; + u64 memsets_nomap; + u64 memsets_fail; u64 releases_seen; u64 releases_applied; u64 releases_missed; /* no VA->PA entry */ @@ -187,11 +194,16 @@ static void svc_add_map(u64 gpu_va, u64 pa) u64 pa_page = pa & ~(u64)(SEMA_PAGE_SIZE - 1); u32 i; + /* + * First-wins: the guest keeps stale copies of freed allocations around, + * and unrelated adjacent u64 pairs can look like {cpu_va, gpu_va}. Both + * produce bogus mappings for a VA that already has a good one, and + * overwriting it would send every semaphore write to the wrong page. + * The first pair seen for a VA is the live allocation. + */ for (i = 0; i < svc.map_n; i++) { - if (svc.map[i].gpu_page == gpu_page) { - svc.map[i].pa_page = pa_page; + if (svc.map[i].gpu_page == gpu_page) return; - } } if (svc.map_n < SEMA_MAP_MAX) { svc.map[svc.map_n].gpu_page = gpu_page; @@ -264,6 +276,10 @@ static void svc_rebuild(struct nv_vgpu_device *vgpu) svc.builds, best_base, best_count, npairs, svc.map_n, svc.releases_seen, svc.releases_applied, svc.releases_missed, svc.releases_stale, svc.releases_jumped); + pr_info("sema-svc: pages=%llu pushes=%llu pending=%llu nomap=%llu memset=%llu/%llu/%llu\n", + svc.pages_decoded, svc.pushes_seen, svc.pushes_pending, + svc.pushes_nomap, svc.memsets_done, svc.memsets_nomap, + svc.memsets_fail); /* show a few entries so mappings can be verified from dmesg */ { u32 shown; @@ -279,6 +295,464 @@ static void svc_rebuild(struct nv_vgpu_device *vgpu) /* push application */ /* ------------------------------------------------------------------ */ +/* C5B5 (TURING_DMA_COPY_A) methods, from the class interface. */ +#define CE_M_SET_SEMAPHORE_A 0x240 +#define CE_M_SET_SEMAPHORE_B 0x244 +#define CE_M_SET_SEMAPHORE_PAY 0x248 +#define CE_M_SET_SRC_PHYS_MODE 0x260 +#define CE_M_SET_DST_PHYS_MODE 0x264 +#define CE_M_LAUNCH_DMA 0x300 +#define CE_M_OFFSET_IN_UPPER 0x400 +#define CE_M_OFFSET_IN_LOWER 0x404 +#define CE_M_OFFSET_OUT_UPPER 0x408 +#define CE_M_OFFSET_OUT_LOWER 0x40c +#define CE_M_LINE_LENGTH_IN 0x418 +#define CE_M_SET_REMAP_CONST_A 0x700 +#define CE_M_SET_REMAP_CONST_B 0x704 +#define CE_M_SET_REMAP_COMPONENTS 0x708 + +/* LAUNCH_DMA fields. */ +#define DMA_DATA_TRANSFER_TYPE(v) ((v) & 0x3) +#define DMA_SEMAPHORE_TYPE(v) (((v) >> 3) & 0x3) +#define DMA_DST_MEMORY_LAYOUT_PITCH(v) (((v) >> 8) & 0x1) +#define DMA_REMAP_ENABLE(v) (((v) >> 10) & 0x1) +#define DMA_SRC_TYPE_PHYSICAL(v) (((v) >> 12) & 0x1) +#define DMA_DST_TYPE_PHYSICAL(v) (((v) >> 13) & 0x1) +#define DMA_SEM_REDUCTION(v) (((v) >> 14) & 0xf) +#define DMA_SEM_REDUCTION_EN(v) (((v) >> 19) & 0x1) +#define SEM_RELEASE_ONE_WORD 1 +#define SEM_REDUCTION_INC 6 + +/* Aperture targets (SET_*_PHYS_MODE). */ +#define CE_APERTURE_VID 0 +#define CE_APERTURE_SYS 1 +#define CE_APERTURE_SYS_NC 2 + +#define SVC_MAX_MEMSETS 32 + +struct svc_memset_op { + u32 type; /* 0 = fill, 1 = copy */ + u64 dst; + u64 value; /* fill value, up to 8 bytes */ + u32 elem; /* 1, 4 or 8 */ + u32 count; /* elements */ + u32 phys; /* DST_TYPE physical */ + u32 aperture; /* when physical */ + u64 src; /* copy source */ + u64 bytes; /* copy length */ + u32 src_phys; + u32 src_aperture; +}; + +/* Decoder state carried across the methods of one push. */ +struct svc_push_state { + u32 remap_a, remap_b, remap_comp; + u32 dst_phys_mode, src_phys_mode; + u32 have_dst_mode, have_src_mode; + u64 offset_out, offset_in; + u32 have_in, have_out; + u32 line_length; + u32 sem_va_hi, sem_va_lo, sem_payload; + u32 have_sem; + u32 have_remap; + u32 release_dma; + u32 sem_inc; +}; + +/* ------------------------------------------------------------------ */ +/* CE execution */ +/* ------------------------------------------------------------------ */ + +/* + * Write `count` elements of `value` (element size `elem`) starting at + * `addr`, in vidmem or guest RAM depending on the aperture DST_TYPE + * selected. Returns 0 on success. + */ +static int svc_fill(struct nv_vgpu_device *vgpu, u64 addr, u64 value, u32 elem, + u32 count, u32 phys, u32 aperture) +{ + u8 pat[8]; + u8 bounce[256]; + size_t total = (size_t)elem * count; + size_t o; + + if (elem != 1 && elem != 2 && elem != 4 && elem != 8) + return -EINVAL; + if (!count || total > (256ULL << 20)) + return -EINVAL; + + memset(pat, 0, sizeof(pat)); + memcpy(pat, &value, elem); + + if (!phys) { + /* Virtual destination: translate through the VA->PA map. */ + u64 pa; + + if (svc_map_lookup(addr, &pa)) + return -ENOENT; + addr = pa; + aperture = CE_APERTURE_SYS; + } + + if (aperture == CE_APERTURE_VID) { + u8 *base = vgpu->bars[1].vaddr; + size_t size = vgpu->bars[1].size; + + if (!base || addr + total > size) + return -EINVAL; + for (o = 0; o + 8 <= total; o += 8) + memcpy(base + addr + o, pat, 8); + if (o < total) + memcpy(base + addr + o, pat, total - o); + return 0; + } + + /* System memory: repeat the element through a bounce buffer. */ + for (o = 0; o < total; ) { + size_t chunk = total - o; + size_t i; + + if (chunk > sizeof(bounce)) + chunk = sizeof(bounce); + for (i = 0; i + 8 <= chunk; i += 8) + memcpy(bounce + i, pat, 8); + if (i < chunk) + memcpy(bounce + i, pat, chunk - i); + if (svc_write(vgpu, addr + o, bounce, chunk)) + return -EIO; + o += chunk; + } + return 0; +} + +/* Byte copy between two CE addresses (src/dst decoding mirrors svc_fill). */ +static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, + u32 dst_phys, u32 dst_aperture, u32 src_phys, u32 src_aperture) +{ + u8 bounce[256]; + size_t o; + + if (!total || total > (256ULL << 20)) + return -EINVAL; + + if (!dst_phys) { + u64 pa; + + if (svc_map_lookup(dst, &pa)) + return -ENOENT; + dst = pa; + dst_aperture = CE_APERTURE_SYS; + } + if (!src_phys) { + u64 pa; + + if (svc_map_lookup(src, &pa)) + return -ENOENT; + src = pa; + src_aperture = CE_APERTURE_SYS; + } + + for (o = 0; o < total; ) { + size_t chunk = total - o; + u32 i; + + if (chunk > sizeof(bounce)) + chunk = sizeof(bounce); + + if (src_aperture == CE_APERTURE_VID) { + u8 *base = vgpu->bars[1].vaddr; + size_t size = vgpu->bars[1].size; + + if (!base || src + o + chunk > size) + return -EINVAL; + memcpy(bounce, base + src + o, chunk); + } else { + if (svc_read(vgpu, src + o, bounce, chunk)) + return -EIO; + } + + if (dst_aperture == CE_APERTURE_VID) { + u8 *base = vgpu->bars[1].vaddr; + size_t size = vgpu->bars[1].size; + + if (!base || dst + o + chunk > size) + return -EINVAL; + memcpy(base + dst + o, bounce, chunk); + } else { + if (svc_write(vgpu, dst + o, bounce, chunk)) + return -EIO; + } + (void)i; + o += chunk; + } + return 0; +} + +/* Execute a decoded operation list (fills and copies). */ +static void svc_run_ops(struct nv_vgpu_device *vgpu, struct svc_memset_op *ops, + u32 n_ops) +{ + u32 i; + + for (i = 0; i < n_ops; i++) { + struct svc_memset_op *m = &ops[i]; + int rc; + + if (m->type == 1) + rc = svc_copy(vgpu, m->dst, m->src, m->bytes, + m->phys, m->aperture, + m->src_phys, m->src_aperture); + else + rc = svc_fill(vgpu, m->dst, m->value, m->elem, + m->count, m->phys, m->aperture); + + if (rc == 0) + svc.memsets_done++; + else if (rc == -ENOENT) + svc.memsets_nomap++; + else + svc.memsets_fail++; + if (svc.memsets_done <= 8) + pr_info("sema-svc: %s dst=0x%llx value=0x%llx elem=%u count=%u bytes=%llu phys=%u ap=%u src=0x%llx rc=%d\n", + m->type == 1 ? "memcopy" : "memset", + m->dst, m->value, m->elem, m->count, + m->bytes, m->phys, m->aperture, m->src, rc); + } +} + +/* + * Decode one page as a method stream and execute the pushes it contains. + * + * Push buffers are written back to back and pushes can span page + * boundaries, so there is no reliable "push start" marker to anchor on: + * decode the whole page as a method stream, carrying decoder state across + * the methods. Collected memory operations are executed only when the + * push's tracking semaphore release is reached and is still pending + * (payload ahead of the semaphore's current value) - recycled buffers + * replay old pushes whose semaphore is already satisfied, and those are + * skipped, so their stale memsets never overwrite newer state. + */ +static void svc_decode_page(struct nv_vgpu_device *vgpu, u8 *buf, u64 pa_page) +{ + struct svc_push_state st; + struct svc_memset_op ops[SVC_MAX_MEMSETS]; + u32 n_ops = 0; + u32 off = 0; + + memset(&st, 0, sizeof(st)); + svc.pages_decoded++; + + while (off + 4 <= SEMA_PAGE_SIZE) { + u32 hdr = 0, op, count, subch, addr; + u32 v0 = 0, v1 = 0, v2 = 0; + + memcpy(&hdr, buf + off, 4); + op = (hdr >> 29) & 0x7; + count = (hdr >> 16) & 0xfff; + subch = (hdr >> 13) & 0x7; + addr = (hdr & 0x1fff) << 2; + + if (op != 1 || !count || count > 64 || + off + 4 + count * 4 > SEMA_PAGE_SIZE) { + off += 4; /* resync */ + continue; + } + + if (count >= 1) + memcpy(&v0, buf + off + 4, 4); + if (count >= 2) + memcpy(&v1, buf + off + 8, 4); + if (count >= 3) + memcpy(&v2, buf + off + 12, 4); + + /* + * OFFSET_IN_UPPER/LOWER and OFFSET_OUT_UPPER/LOWER are + * contiguous (0x400..0x40c) and the driver pushes them as a + * single INC_METHOD run (count 2 or 4). Apply every word of + * the run, otherwise the low halves stay zero. + */ + if (addr >= CE_M_OFFSET_IN_UPPER && + addr <= CE_M_OFFSET_OUT_LOWER) { + u32 vv[4] = { v0, v1, v2, 0 }; + u32 k, n = count > 4 ? 4 : count; + + if (n == 4) + memcpy(&vv[3], buf + off + 16, 4); + for (k = 0; k < n; k++) { + switch (addr + 4 * k) { + case CE_M_OFFSET_IN_UPPER: + st.offset_in = + (st.offset_in & 0xffffffffULL) | + ((u64)(vv[k] & 0x1ffff) << 32); + st.have_in = 1; + break; + case CE_M_OFFSET_IN_LOWER: + st.offset_in = + (st.offset_in & ~0xffffffffULL) | + vv[k]; + st.have_in = 1; + break; + case CE_M_OFFSET_OUT_UPPER: + st.offset_out = + (st.offset_out & 0xffffffffULL) | + ((u64)(vv[k] & 0x1ffff) << 32); + st.have_out = 1; + break; + case CE_M_OFFSET_OUT_LOWER: + st.offset_out = + (st.offset_out & ~0xffffffffULL) | + vv[k]; + st.have_out = 1; + break; + default: + break; + } + } + off += 4 + count * 4; + continue; + } + + if (subch == 0 && addr == 0) { + /* SET_OBJECT on the channel subchannel: a push starts */ + n_ops = 0; + memset(&st, 0, sizeof(st)); + off += 4 + count * 4; + continue; + } + + if (subch != 4) { + off += 4 + count * 4; + continue; + } + + switch (addr) { + case CE_M_SET_SEMAPHORE_A: + st.sem_va_hi = v0; + if (count >= 3) { + st.sem_va_lo = v1; + st.sem_payload = v2; + st.have_sem = 1; + } + break; + case CE_M_SET_SEMAPHORE_B: + st.sem_va_lo = v0; + st.have_sem = 1; + break; + case CE_M_SET_SEMAPHORE_PAY: + st.sem_payload = v0; + st.have_sem = 1; + break; + case CE_M_SET_REMAP_CONST_A: + if (count >= 3) { + st.remap_a = v0; + st.remap_b = v1; + st.remap_comp = v2; + st.have_remap = 1; + } + break; + case CE_M_SET_SRC_PHYS_MODE: + st.src_phys_mode = v0 & 0x3; + st.have_src_mode = 1; + break; + case CE_M_SET_DST_PHYS_MODE: + st.dst_phys_mode = v0 & 0x3; + st.have_dst_mode = 1; + break; + case CE_M_LINE_LENGTH_IN: + st.line_length = v0; + break; + case CE_M_LAUNCH_DMA: { + u32 dma = v0; + u32 transfer = DMA_DATA_TRANSFER_TYPE(dma); + + if (transfer != 0 && st.line_length && st.have_out) { + if (DMA_REMAP_ENABLE(dma) && st.have_remap && + DMA_DST_MEMORY_LAYOUT_PITCH(dma) && + n_ops < SVC_MAX_MEMSETS) { + u32 csz = (st.remap_comp >> 16) & 0x3; + u32 ndst = (st.remap_comp >> 24) & 0x3; + struct svc_memset_op *m = &ops[n_ops]; + + m->type = 0; + m->dst = st.offset_out; + m->value = ndst >= 1 ? + (((u64)st.remap_b << 32) | + st.remap_a) : st.remap_b; + m->elem = (csz + 1) * (ndst + 1); + m->count = st.line_length; + m->phys = DMA_DST_TYPE_PHYSICAL(dma); + m->aperture = st.have_dst_mode ? + st.dst_phys_mode : CE_APERTURE_VID; + n_ops++; + } else if (!DMA_REMAP_ENABLE(dma) && + st.have_in && + n_ops < SVC_MAX_MEMSETS) { + struct svc_memset_op *m = &ops[n_ops]; + + m->type = 1; + m->dst = st.offset_out; + m->bytes = st.line_length; + m->phys = DMA_DST_TYPE_PHYSICAL(dma); + m->aperture = st.have_dst_mode ? + st.dst_phys_mode : CE_APERTURE_VID; + m->src = st.offset_in; + m->src_phys = + DMA_SRC_TYPE_PHYSICAL(dma); + m->src_aperture = st.have_src_mode ? + st.src_phys_mode : CE_APERTURE_VID; + n_ops++; + } + } + + if (DMA_SEMAPHORE_TYPE(dma) == + SEM_RELEASE_ONE_WORD) { + u64 sem_va, sem_pa = 0; + u32 cur = 0; + int pending = 0; + + svc.pushes_seen++; + sem_va = ((u64)st.sem_va_hi << 32) | + st.sem_va_lo; + + if (st.have_sem && st.sem_payload && + sem_va >= SEMA_VA_LO && sem_va < SEMA_VA_HI && + !svc_map_lookup(sem_va, &sem_pa)) { + if (svc_read(vgpu, sem_pa, &cur, + sizeof(cur))) + svc.pushes_nomap++; + else if (st.sem_payload > cur) + pending = 1; + } + + if (pending) { + svc.pushes_pending++; + svc_run_ops(vgpu, ops, n_ops); + if (DMA_SEM_REDUCTION_EN(dma) && + DMA_SEM_REDUCTION(dma) == + SEM_REDUCTION_INC) + cur += st.sem_payload; + else + cur = st.sem_payload; + if (!svc_write(vgpu, sem_pa, &cur, + sizeof(cur))) + svc.releases_applied++; + } + /* release ends the push */ + n_ops = 0; + } + break; + } + default: + break; + } + off += 4 + count * 4; + } +} + +/* ------------------------------------------------------------------ */ +/* push application */ +/* ------------------------------------------------------------------ */ + /* * Look for CE semaphore-release method groups inside mapped pages: * @@ -298,6 +772,13 @@ static u32 svc_apply_page(struct nv_vgpu_device *vgpu, u8 *buf, u64 pa_page) if (svc_read(vgpu, pa_page, buf, SEMA_PAGE_SIZE)) return 0; + /* + * Decoded execution: walk the page as a method stream and perform the + * pushed work (PTE memsets, semaphore releases) for pushes that are + * still pending. + */ + svc_decode_page(vgpu, buf, pa_page); + for (off = 0; off + 16 <= SEMA_PAGE_SIZE; off += 4) { u32 hdr, hi, lo, payload; u64 va, pa, cur; diff --git a/kernel/main.c b/kernel/main.c index 637c79b..0a1ebc2 100644 --- a/kernel/main.c +++ b/kernel/main.c @@ -2560,5 +2560,94 @@ static const struct kernel_param_ops nv_vgpu_dump_vidmem_ops = { module_param_cb(dump_vidmem, &nv_vgpu_dump_vidmem_ops, NULL, 0200); MODULE_PARM_DESC(dump_vidmem, "Dump emulated vidmem range (hex offset,hex len)"); +/* + * Debug: hex-dump a guest physical range (kthread: guest-DMA needs it): + * echo 0x,0x > .../parameters/dump_guest + */ +struct nv_vgpu_dump_guest_req { + u64 gpa; + unsigned long len; + struct completion done; +}; + +static int nv_vgpu_dump_guest_thread(void *arg) +{ + struct nv_vgpu_dump_guest_req *req = arg; + struct nv_vgpu_device *vgpu; + u32 *buf; + unsigned long i; + + buf = kmalloc(256, GFP_KERNEL); + if (!buf) + return -ENOMEM; + + mutex_lock(&nv_vgpu_driver.lock); + list_for_each_entry(vgpu, &nv_vgpu_driver.devices, list) { + u64 o; + + if (atomic_read(&vgpu->vfio_dev_open_count) <= 0) + continue; + pr_info("gsp-resman-vfio: guest dump 0x%llx..0x%llx\n", + req->gpa, req->gpa + req->len); + for (o = 0; o < req->len; o += 256) { + unsigned long chunk = min_t(unsigned long, 256, + req->len - o); + + if (vfio_dma_rw(&vgpu->vfio_dev, req->gpa + o, buf, + chunk, false)) { + pr_info("gsp-resman-vfio: guest dump read failed at 0x%llx\n", + req->gpa + o); + break; + } + for (i = 0; i < chunk / 4; i += 4) + pr_info("gsp-resman-vfio: 0x%llx: %08x %08x %08x %08x\n", + req->gpa + o + i * 4, buf[i], buf[i + 1], + buf[i + 2], buf[i + 3]); + } + } + mutex_unlock(&nv_vgpu_driver.lock); + + kfree(buf); + complete(&req->done); + return 0; +} + +static int nv_vgpu_dump_guest_set(const char *val, const struct kernel_param *kp) +{ + static struct nv_vgpu_dump_guest_req req; + struct task_struct *task; + char buf[64]; + char *comma; + + memset(&req, 0, sizeof(req)); + strscpy(buf, val, sizeof(buf)); + comma = strchr(buf, ','); + if (comma) + *comma = '\0'; + if (kstrtoull(buf, 0, &req.gpa)) + return -EINVAL; + req.len = 0x100; + if (comma && kstrtoul(comma + 1, 0, &req.len)) + return -EINVAL; + if (!req.len) + req.len = 0x100; + if (req.len > 0x4000) + req.len = 0x4000; + + init_completion(&req.done); + task = kthread_run(nv_vgpu_dump_guest_thread, &req, "vgpu-dump"); + if (IS_ERR(task)) + return PTR_ERR(task); + wait_for_completion(&req.done); + return 0; +} + +static const struct kernel_param_ops nv_vgpu_dump_guest_ops = { + .set = nv_vgpu_dump_guest_set, +}; + +module_param_cb(dump_guest, &nv_vgpu_dump_guest_ops, NULL, 0200); +MODULE_PARM_DESC(dump_guest, "Hex-dump a guest physical range (hex gpa,hex len)"); + module_init(nv_vgpu_init); module_exit(nv_vgpu_exit); diff --git a/tools/cuda-ab/README.md b/tools/cuda-ab/README.md new file mode 100644 index 0000000..db87149 --- /dev/null +++ b/tools/cuda-ab/README.md @@ -0,0 +1,32 @@ +# CUDA A/B harness + +Run one deterministic CUDA workload on the **host** (bare-metal GPU, the +oracle) and in the **guest** (Open-IOV emulated vGPU), then compare the +machine-readable results. This mirrors the earlier host/guest `nvidia-smi` +comparison, but exercises actual execution: allocation, HtoD/DtoH copies, +kernel launch, synchronisation and verification. + +``` +tools/cuda-ab/ab.sh all # build + host run + guest run + compare +tools/cuda-ab/ab.sh host # host only -> run/host.txt +tools/cuda-ab/ab.sh guest # push binary, run in the running VM -> run/guest.txt +tools/cuda-ab/ab.sh compare # diff the stable result lines +``` + +## Files + +| File | Purpose | +|---|---| +| `vector_add.cu` | the workload (`DEVICE`/`RESULT`/`STATUS` output lines) | +| `ab.sh` | host build/run, guest push/run, comparison | +| `run/` | generated outputs (`host.txt`, `guest.txt`) | + +## Notes + +* The binary links the CUDA runtime **statically** (`--cudart=static`), so the + guest needs no CUDA toolkit - only the guest driver's `libcuda.so`. +* The guest transfer uses the guest agent (base64 chunks); the VM must be + running (`scripts/vgpu_vm.py`). +* `DEVICE` lines are normalised before diffing (the guest reports the + emulated GPU name); `RESULT` and `STATUS` must match exactly. +* Host requirement: `nvidia-cuda-toolkit` (`nvcc`). diff --git a/tools/cuda-ab/ab.sh b/tools/cuda-ab/ab.sh new file mode 100755 index 0000000..07f6b24 --- /dev/null +++ b/tools/cuda-ab/ab.sh @@ -0,0 +1,88 @@ +#!/bin/bash +# SPDX-License-Identifier: BSD-3-Clause +# +# ab.sh - CUDA A/B harness: run the same workload on the host (bare metal, +# real Turing GPU) and in the guest (Open-IOV emulated vGPU), then compare the +# machine-readable results. Mirrors the nvidia-smi host/guest comparison used +# earlier: the host run is the oracle. +# +# Usage: +# tools/cuda-ab/ab.sh host # build + run on the host, save run/host.txt +# tools/cuda-ab/ab.sh guest # push the binary + run in the running VM +# tools/cuda-ab/ab.sh compare # diff the stable result lines +# tools/cuda-ab/ab.sh all # host + guest + compare +# +# The binary links the CUDA runtime statically, so the guest needs no CUDA +# toolkit - only the guest driver's libcuda.so. +set -e + +HERE="$(cd "$(dirname "$0")" && pwd)" +REPO="$(cd "$HERE/../.." && pwd)" +VMEXEC="$REPO/scripts/vgpu_vm.py" +RUN="$HERE/run" +SRC="$HERE/vector_add.cu" +BIN="$RUN/vector_add" + +mkdir -p "$RUN" + +build_host() { + command -v nvcc >/dev/null || { echo "nvcc not found (install nvidia-cuda-toolkit)"; exit 1; } + nvcc -O2 --cudart=static -o "$BIN" "$SRC" +} + +run_host() { + build_host + echo "[*] host run (oracle)" + "$BIN" | tee "$RUN/host.txt" +} + +push_bin() { + [ -x "$BIN" ] || build_host + local b64 chunk off total size + b64=$(base64 -w0 "$BIN") + total=${#b64} + off=0 + size=70000 + echo "[*] pushing $(stat -c%s "$BIN") bytes to guest" + "$VMEXEC" exec "rm -f /tmp/vadd.b64" >/dev/null 2>&1 + while [ $off -lt $total ]; do + chunk=${b64:$off:$size} + "$VMEXEC" exec "echo -n '$chunk' >> /tmp/vadd.b64" >/dev/null 2>&1 + off=$((off + size)) + done + "$VMEXEC" exec "base64 -d /tmp/vadd.b64 > /tmp/vector_add && chmod +x /tmp/vector_add && rm -f /tmp/vadd.b64 && echo PUSH_OK" +} + +run_guest() { + push_bin + echo "[*] guest run (emulated vGPU)" + "$VMEXEC" exec "/tmp/vector_add" 2>&1 | tee "$RUN/guest.txt" +} + +compare() { + local h g + [ -f "$RUN/host.txt" ] || { echo "missing $RUN/host.txt (run host first)"; exit 1; } + [ -f "$RUN/guest.txt" ] || { echo "missing $RUN/guest.txt (run guest first)"; exit 1; } + h=$(grep -E '^(DEVICE|RESULT|STATUS)' "$RUN/host.txt") + g=$(grep -E '^(DEVICE|RESULT|STATUS)' "$RUN/guest.txt") + echo "--- host:"; echo "$h" + echo "--- guest:"; echo "$g" + # DEVICE lines carry the GPU name (differs by design); compare the rest. + h_norm=$(echo "$h" | sed 's/^DEVICE .*/DEVICE /') + g_norm=$(echo "$g" | sed 's/^DEVICE .*/DEVICE /') + if [ "$h_norm" = "$g_norm" ]; then + echo "A/B: MATCH" + else + echo "A/B: MISMATCH" + diff <(echo "$h_norm") <(echo "$g_norm") || true + exit 1 + fi +} + +case "${1:-all}" in +host) run_host ;; +guest) run_guest ;; +compare) compare ;; +all) run_host; run_guest; compare ;; +*) echo "usage: $0 [host|guest|compare|all]"; exit 1 ;; +esac diff --git a/tools/cuda-ab/run/host.txt b/tools/cuda-ab/run/host.txt new file mode 100644 index 0000000..53e7e38 --- /dev/null +++ b/tools/cuda-ab/run/host.txt @@ -0,0 +1,3 @@ +DEVICE NVIDIA GeForce RTX 2080 Super with Max-Q Design cc=7.5 sm=48 mem=8165523456 +RESULT n=1048576 bad=0 checksum=392344090.500000 +STATUS OK diff --git a/tools/cuda-ab/run/vector_add b/tools/cuda-ab/run/vector_add new file mode 100755 index 0000000..c1581c1 Binary files /dev/null and b/tools/cuda-ab/run/vector_add differ diff --git a/tools/cuda-ab/vector_add.cu b/tools/cuda-ab/vector_add.cu new file mode 100644 index 0000000..2e3f1d6 --- /dev/null +++ b/tools/cuda-ab/vector_add.cu @@ -0,0 +1,93 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * vector_add.cu - CUDA A/B workload for the Open-IOV harness. + * + * One deterministic, machine-readable workload that can run on the host (bare + * metal) and in a guest (emulated vGPU). The output lines are stable so the + * two runs can be diffed: + * + * DEVICE cc=. sm= mem= + * RESULT n= bad= checksum= + * STATUS OK|FAIL + * + * Build with a static CUDA runtime (--cudart=static) so the guest only needs + * the driver's libcuda.so, not a CUDA toolkit installation. + */ + +#include +#include +#include + +#define CHECK(x) \ + do { \ + cudaError_t e_ = (x); \ + if (e_ != cudaSuccess) { \ + printf("CUDA_ERROR %s at line %d\n", \ + cudaGetErrorString(e_), __LINE__); \ + printf("STATUS FAIL\n"); \ + return 1; \ + } \ + } while (0) + +__global__ void vadd(const float *a, const float *b, float *c, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = a[i] + b[i]; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 20; + size_t bytes = (size_t)n * sizeof(float); + float *h_a, *h_b, *h_c; + float *d_a, *d_b, *d_c; + double sum = 0.0; + int bad = 0, i; + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name, + prop.major, prop.minor, prop.multiProcessorCount, + (unsigned long long)prop.totalGlobalMem); + + h_a = (float *)malloc(bytes); + h_b = (float *)malloc(bytes); + h_c = (float *)malloc(bytes); + if (!h_a || !h_b || !h_c) + return 1; + + for (i = 0; i < n; i++) { + h_a[i] = (float)(i % 1000) * 0.5f; + h_b[i] = (float)(i % 997) * 0.25f; + } + + CHECK(cudaMalloc(&d_a, bytes)); + CHECK(cudaMalloc(&d_b, bytes)); + CHECK(cudaMalloc(&d_c, bytes)); + CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); + + vadd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < n; i++) { + if (h_c[i] != h_a[i] + h_b[i]) + bad++; + sum += h_c[i]; + } + + printf("RESULT n=%d bad=%d checksum=%.6f\n", n, bad, sum); + printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL"); + + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_b)); + CHECK(cudaFree(d_c)); + free(h_a); + free(h_b); + free(h_c); + return bad == 0 ? 0 : 1; +} -- cgit v1.2.3