From 2637832fcae6157f1846ad1ae4d9d8acc88543e2 Mon Sep 17 00:00:00 2001 From: Thomas Weatherly Date: Tue, 1 Sep 2026 14:31:19 -0400 Subject: [PATCH] =?UTF-8?q?tests:=20vm=5Fstress=20=E2=80=94=20TLB=20pressu?= =?UTF-8?q?re=20mixed=20with=20a=20VX=5FMEM=5FPHYS=20identity-mapped=20buf?= =?UTF-8?q?fer?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Every task strides across pages with an odd stride so consecutive touches spread across TLB sets, while also reading a bias through a pinned-slab identity mapping — paged and identity translations contend in the same run, which no existing vm_test mode does. The physical buffer is allocated, released, and allocated again first so the identity map is re-installed over a recycled slab range (the re-map must be idempotent and quiet). Wired into the default regression list and the vm catalog (simx + rtlsim, both XLENs, VORTEX_RANDOMIZE_VA=1). Co-Authored-By: Claude Fable 5 --- ci/testcases/vm.yaml | 12 ++ tests/regression/Makefile | 2 +- tests/regression/vm_stress/Makefile | 16 +++ tests/regression/vm_stress/common.h | 17 +++ tests/regression/vm_stress/kernel.cpp | 21 +++ tests/regression/vm_stress/main.cpp | 185 ++++++++++++++++++++++++++ 6 files changed, 252 insertions(+), 1 deletion(-) create mode 100644 tests/regression/vm_stress/Makefile create mode 100644 tests/regression/vm_stress/common.h create mode 100644 tests/regression/vm_stress/kernel.cpp create mode 100644 tests/regression/vm_stress/main.cpp diff --git a/ci/testcases/vm.yaml b/ci/testcases/vm.yaml index 2dd26aad4b..738eb3851f 100644 --- a/ci/testcases/vm.yaml +++ b/ci/testcases/vm.yaml @@ -137,6 +137,18 @@ tests: OPTS: -t2 -n16 -p4 env: VORTEX_RANDOMIZE_VA: 1 +# TLB-pressure sweep mixed with a VX_MEM_PHYS identity-mapped buffer: every +# task strides across pages (odd stride) while also reading through the +# pinned slab, so paged and identity translations contend in the same run. +- id: vm-stress + via: make-run + drivers: + - simx + - rtlsim + dir: tests/regression/vm_stress + target: run-{driver} + env: + VORTEX_RANDOMIZE_VA: 1 - id: vm-amo via: make-run drivers: diff --git a/tests/regression/Makefile b/tests/regression/Makefile index cccdee5116..42a468a27d 100644 --- a/tests/regression/Makefile +++ b/tests/regression/Makefile @@ -15,7 +15,7 @@ TESTS := \ dxa_copy dxa_copy_mcast dxa_kmajor_check \ async_barrier async_gbarrier packld wgather wsync quad_lod \ occupancy multikernel reopen module_reload single_cta \ - vm_test vm_fault + vm_test vm_fault vm_stress # --- common exclude list --------------------------------------------- EXCLUDE := diff --git a/tests/regression/vm_stress/Makefile b/tests/regression/vm_stress/Makefile new file mode 100644 index 0000000000..b36a07d74d --- /dev/null +++ b/tests/regression/vm_stress/Makefile @@ -0,0 +1,16 @@ +ROOT_DIR := $(realpath ../../..) +include $(ROOT_DIR)/config.mk + +PROJECT := vm_stress + +SRC_DIR := $(VORTEX_HOME)/tests/regression/$(PROJECT) + +SRCS := $(SRC_DIR)/main.cpp + +VX_SRCS := $(SRC_DIR)/kernel.cpp + +OPTS ?= -n256 + +KERNEL_LIB := vortex2 + +include ../common.mk diff --git a/tests/regression/vm_stress/common.h b/tests/regression/vm_stress/common.h new file mode 100644 index 0000000000..a63f80adec --- /dev/null +++ b/tests/regression/vm_stress/common.h @@ -0,0 +1,17 @@ +#ifndef _COMMON_H_ +#define _COMMON_H_ + +#define WORDS_PER_PAGE (4096 / 4) + +typedef struct { + uint32_t num_tasks; + uint32_t pages_per_task; + uint32_t total_pages; + uint32_t stride_pages; + uint32_t phys_words; + uint64_t src_addr; + uint64_t dst_addr; + uint64_t phys_addr; +} kernel_arg_t; + +#endif diff --git a/tests/regression/vm_stress/kernel.cpp b/tests/regression/vm_stress/kernel.cpp new file mode 100644 index 0000000000..c1df67de35 --- /dev/null +++ b/tests/regression/vm_stress/kernel.cpp @@ -0,0 +1,21 @@ +#include +#include "common.h" + +// Each task touches one word in each of its pages, with an odd page stride +// so consecutive touches spread across TLB sets and defeat streaming reuse. +// The physical buffer is identity-mapped (VX_MEM_PHYS) and exercises the +// pinned-slab translation path alongside the paged mappings. +__kernel void kernel_main(kernel_arg_t* __UNIFORM__ arg) { + auto src_ptr = reinterpret_cast(arg->src_addr); + auto dst_ptr = reinterpret_cast(arg->dst_addr); + auto phys_ptr = reinterpret_cast(arg->phys_addr); + + uint32_t task_id = blockIdx.x * blockDim.x + threadIdx.x; + uint32_t bias = phys_ptr[task_id % arg->phys_words]; + + for (uint32_t k = 0; k < arg->pages_per_task; ++k) { + uint32_t page = ((task_id * arg->pages_per_task + k) * arg->stride_pages) % arg->total_pages; + uint32_t word = page * WORDS_PER_PAGE + (task_id % WORDS_PER_PAGE); + dst_ptr[word] = src_ptr[word] + bias; + } +} diff --git a/tests/regression/vm_stress/main.cpp b/tests/regression/vm_stress/main.cpp new file mode 100644 index 0000000000..cbb3a98939 --- /dev/null +++ b/tests/regression/vm_stress/main.cpp @@ -0,0 +1,185 @@ +#include +#include +#include +#include +#include +#include "common.h" + +#define RT_CHECK(_expr) \ + do { \ + int _ret = _expr; \ + if (0 == _ret) \ + break; \ + printf("Error: '%s' returned %d!\n", #_expr, (int)_ret); \ + cleanup(); \ + exit(-1); \ + } while (false) + +const char* kernel_file = "kernel.vxbin"; +uint32_t total_pages = 256; +uint32_t phys_words = 64; + +vx_device_h device = nullptr; +vx_buffer_h src_buffer = nullptr; +vx_buffer_h dst_buffer = nullptr; +vx_buffer_h phys_buffer = nullptr; +vx_queue_h queue = nullptr; +vx_module_h module_ = nullptr; +vx_kernel_h kernel = nullptr; +kernel_arg_t kernel_arg = {}; + +static void show_usage() { + std::cout << "Vortex VM stress test." << std::endl; + std::cout << "Usage: [-k: kernel] [-n pages] [-h: help]" << std::endl; +} + +static void parse_args(int argc, char **argv) { + int c; + while ((c = getopt(argc, argv, "n:k:h")) != -1) { + switch (c) { + case 'n': + total_pages = atoi(optarg); + break; + case 'k': + kernel_file = optarg; + break; + case 'h': { + show_usage(); + exit(0); + } break; + default: + show_usage(); + exit(-1); + } + } +} + +void cleanup() { + if (device) { + if (src_buffer) vx_buffer_release(src_buffer); + if (dst_buffer) vx_buffer_release(dst_buffer); + if (phys_buffer) vx_buffer_release(phys_buffer); + if (kernel) vx_kernel_release(kernel); + if (module_) vx_module_release(module_); + if (queue) vx_queue_release(queue); + vx_device_dump_perf(device, stdout); + vx_device_release(device); + } +} + +int main(int argc, char *argv[]) { + parse_args(argc, argv); + + std::cout << "open device connection" << std::endl; + RT_CHECK(vx_device_open(0, &device)); + + vx_queue_info_t qi = { sizeof(qi), nullptr, VX_QUEUE_PRIORITY_NORMAL, 0 }; + RT_CHECK(vx_queue_create(device, &qi, &queue)); + + uint64_t num_cores, num_warps, num_threads; + RT_CHECK(vx_device_query(device, VX_CAPS_NUM_CORES, &num_cores)); + RT_CHECK(vx_device_query(device, VX_CAPS_NUM_WARPS, &num_warps)); + RT_CHECK(vx_device_query(device, VX_CAPS_NUM_THREADS, &num_threads)); + + uint32_t num_tasks = num_cores * num_warps * num_threads; + uint32_t pages_per_task = (total_pages + num_tasks - 1) / num_tasks; + uint32_t num_words = total_pages * WORDS_PER_PAGE; + uint64_t buf_size = uint64_t(num_words) * sizeof(uint32_t); + + kernel_arg.num_tasks = num_tasks; + kernel_arg.pages_per_task = pages_per_task; + kernel_arg.total_pages = total_pages; + // odd stride so consecutive pages land in different TLB banks + kernel_arg.stride_pages = 17; + kernel_arg.phys_words = phys_words; + + std::cout << "pages: " << total_pages << ", tasks: " << num_tasks + << ", pages/task: " << pages_per_task << std::endl; + + std::cout << "allocate device memory" << std::endl; + RT_CHECK(vx_buffer_create(device, buf_size, VX_MEM_READ, &src_buffer)); + RT_CHECK(vx_buffer_address(src_buffer, &kernel_arg.src_addr)); + RT_CHECK(vx_buffer_create(device, buf_size, VX_MEM_WRITE, &dst_buffer)); + RT_CHECK(vx_buffer_address(dst_buffer, &kernel_arg.dst_addr)); + // physical (identity-mapped) buffer. Allocate, release, and allocate + // again first: the freed slab range is typically handed back for the + // second allocation, so the identity map is re-installed over the same + // PA — the re-map must be idempotent and quiet. + RT_CHECK(vx_buffer_create(device, phys_words * sizeof(uint32_t), + VX_MEM_READ | VX_MEM_PHYS, &phys_buffer)); + RT_CHECK(vx_buffer_release(phys_buffer)); + phys_buffer = nullptr; + RT_CHECK(vx_buffer_create(device, phys_words * sizeof(uint32_t), + VX_MEM_READ | VX_MEM_PHYS, &phys_buffer)); + RT_CHECK(vx_buffer_address(phys_buffer, &kernel_arg.phys_addr)); + + std::cout << "upload buffers" << std::endl; + std::vector h_src(num_words); + std::vector h_dst(num_words, 0); + std::vector h_phys(phys_words); + for (uint32_t i = 0; i < num_words; ++i) { + h_src[i] = i * 2654435761u; + } + for (uint32_t i = 0; i < phys_words; ++i) { + h_phys[i] = 7 + i; + } + RT_CHECK(vx_enqueue_write(queue, src_buffer, 0, h_src.data(), buf_size, 0, nullptr, nullptr)); + RT_CHECK(vx_enqueue_write(queue, dst_buffer, 0, h_dst.data(), buf_size, 0, nullptr, nullptr)); + RT_CHECK(vx_enqueue_write(queue, phys_buffer, 0, h_phys.data(), phys_words * sizeof(uint32_t), 0, nullptr, nullptr)); + + std::cout << "load kernel module" << std::endl; + RT_CHECK(vx_module_load_file(device, kernel_file, &module_)); + RT_CHECK(vx_module_get_kernel(module_, "main", &kernel)); + + std::cout << "launch kernel" << std::endl; + vx_event_h launch_ev = nullptr, read_ev = nullptr; + { + vx_launch_info_t li = {}; + li.struct_size = sizeof(li); + li.kernel = kernel; + li.args_host = &kernel_arg; + li.args_size = sizeof(kernel_arg); + li.ndim = 1; + li.grid_dim[0] = num_tasks / num_threads; + li.block_dim[0] = num_threads; + RT_CHECK(vx_enqueue_launch(queue, &li, 0, nullptr, &launch_ev)); + } + + std::cout << "download destination buffer" << std::endl; + RT_CHECK(vx_enqueue_read(queue, h_dst.data(), dst_buffer, 0, buf_size, 1, &launch_ev, &read_ev)); + + std::cout << "wait for completion" << std::endl; + RT_CHECK(vx_event_wait_value(read_ev, 1, VX_TIMEOUT_INFINITE)); + vx_event_release(read_ev); + vx_event_release(launch_ev); + + std::cout << "verify result" << std::endl; + int errors = 0; + for (uint32_t t = 0; t < num_tasks; ++t) { + uint32_t bias = h_phys[t % phys_words]; + for (uint32_t k = 0; k < pages_per_task; ++k) { + uint32_t page = ((t * pages_per_task + k) * kernel_arg.stride_pages) % total_pages; + uint32_t word = page * WORDS_PER_PAGE + (t % WORDS_PER_PAGE); + uint32_t ref = h_src[word] + bias; + if (h_dst[word] != ref) { + if (errors < 20) { + printf("*** error: task=%u page=%u word=%u expected=0x%x actual=0x%x\n", + t, page, word, ref, h_dst[word]); + } + ++errors; + } + } + } + + std::cout << "cleanup" << std::endl; + cleanup(); + + if (errors != 0) { + std::cout << "Found " << std::dec << errors << " errors!" << std::endl; + std::cout << "FAILED!" << std::endl; + return errors; + } + + std::cout << "PASSED!" << std::endl; + return 0; +}