summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--kernel/gsp_sema_service.c487
-rw-r--r--kernel/main.c89
-rw-r--r--tools/cuda-ab/README.md32
-rwxr-xr-xtools/cuda-ab/ab.sh88
-rw-r--r--tools/cuda-ab/run/host.txt3
-rwxr-xr-xtools/cuda-ab/run/vector_addbin0 -> 820856 bytes
-rw-r--r--tools/cuda-ab/vector_add.cu93
7 files changed, 789 insertions, 3 deletions
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<gpa>,0x<len> > .../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 <gpu>/')
+ g_norm=$(echo "$g" | sed 's/^DEVICE .*/DEVICE <gpu>/')
+ 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
--- /dev/null
+++ b/tools/cuda-ab/run/vector_add
Binary files 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 <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ * RESULT n=<count> bad=<mismatches> checksum=<sum>
+ * 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 <cstdio>
+#include <cstdlib>
+#include <cuda_runtime.h>
+
+#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;
+}
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-12 21:40:08 +0000