diff --git a/CMakeLists.txt b/CMakeLists.txt index c315e83..b5f2db2 100755 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -32,7 +32,7 @@ else () target_compile_options(qdma PRIVATE -O3 -fPIC -mavx512f -no-pie) endif () -add_subdirectory(example) +add_subdirectory(tb) install( TARGETS qdma diff --git a/README.md b/README.md index aad30c2..121ff28 100644 --- a/README.md +++ b/README.md @@ -275,6 +275,137 @@ around 8M bridge write, when ops exceeds this, the latency increase a lot. --- +## CPU && GPU +### Throughput +- pool_size = 1GB +- total_cmds = 256 * 1024 +#### H2C + +| CPU | Bytes | Speed | | GPU | Bytes | Speed | +| --- | :---: | :---: | --- | --- | :---: | :---: | +| | 64 | 2.10 | | | 64 | 1.58 | +| | 128 | 3.94 | | | 128 | 2.96 | +| | 256 | 7.78 | | | 256 | 5.64 | +| | 512 | 12.09 | | | 512 | 10.68 | +| | 1024 | 12.14 | | | 1024 | 10.61 | + +#### C2H + +| CPU | Bytes | Speed | | GPU | Bytes | Speed | +| --- | :---: | :---: | --- | --- | :---: | :---: | +| | 64 | 4.97 | | | 64 | 4.97 | +| | 128 | 9.93 | | | 128 | 9.93 | +| | 256 | 11.91 | | | 256 | 11.92 | +| | 512 | 12.19 | | | 512 | 13.00 | +| | 1024 | 12.18 | | | 1024 | 13.00 | +--- +### random +- pool_size = 1GB +- total_cmds = 256 * 1024 +#### H2C + +| CPU | Bytes | Speed | OPS | | GPU | Bytes | Speed | OPS | +| :-: | :---: | :---: | :----: | --- | --- | :---: | :---: | :----: | +| | 64 | 2.01 | 33.732 | | | 64 | 1.52 | 25.472 | +| | 128 | 3.98 | 33.417 | | | 128 | 2.92 | 24.535 | +| | 256 | 7.30 | 30.629 | | | 256 | 5.82 | 24.400 | +| | 512 | 12.00 | 25.176 | | | 512 | 10.48 | 21.971 | +| | 1024 | 11.95 | 12.530 | | | 1024 | 10.63 | 11.151 | + +#### C2H + +| CPU | Bytes | Speed | OPS | | GPU | Bytes | Speed | OPS | +| :-: | :---: | :---: | :----: | --- | --- | :---: | :---: | :----: | +| | 64 | 4.97 | 83.333 | | | 64 | 4.97 | 83.332 | +| | 128 | 9.90 | 83.035 | | | 128 | 9.93 | 83.332 | +| | 256 | 11.68 | 48.991 | | | 256 | 11.92 | 50.000 | +| | 512 | 11.76 | 24.660 | | | 512 | 13.00 | 27.269 | +| | 1024 | 11.75 | 12.3 | | | 1024 | 13.00 | 13.633 | +--- +### latency +- pool_size = 1GB +- total_cmds = 256 * 1024 +- Wait cycles: minimum cycles between two cmds +#### H2C + +| CPU | Bytes | Wait cycles | OPS limit (Mops) | Throughput (Mops) | Throughput (GB/s) | Latency (us) | +| :-----: | :-------: | :-------------: | :------------------: | :-------------------: | :-------------------: | :--------------: | +| | 64 | 50 | 5 | 4.6 | 0.3 | 1.17 | +| | | 25 | 10 | 8.8 | 0.6 | 1.12 | +| | | 12 | 20.8 | 17.0 | 1.1 | 1.16 | +| | | 6 | 41.7 | 29.8 | 1.9 | 1.17 | +| | | 0 | - | 31.6 | 2.0 | 3.15 | +| | | | | | | | +| | 4 * 1024 | 100 | 2.5 | 2.3 | 9.3 | 1.60 | +| | | 90 | 2.8 | 2.6 | 10.4 | 1.62 | +| | | 85 | 2.9 | 2.7 | 11.0 | 1.67 | +| | | 80 | 3.1 | 2.9 | 11.6 | 1.91 | +| | | 75 | 3.3 | 3.1 | 12.3 | 13.99 | +| | | 70 | 3.6 | 3.1 | 12.3 | 14.16 | +| | | 50 | 5.0 | 3.1 | 12.3 | 14.33 | +| | | | | | | | +| **GPU** | **Bytes** | **Wait cycles** | **OPS limit (Mops)** | **Throughput (Mops)** | **Throughput (GB/s)** | **Latency (us)** | +| | 64 | 50 | 5 | 4.6 | 0.3 | 1.23 | +| | | 25 | 10 | 8.8 | 0.6 | 1.22 | +| | | 12 | 20.8 | 17.0 | 1.1 | 1.23 | +| | | 6 | 41.7 | 29.8 | 1.9 | 1.21 | +| | | 0 | - | 30.4 | 1.9 | 3.30 | +| | | | | | | | +| | 4 * 1024 | 100 | 2.5 | 2.3 | 9.3 | 2.94 | +| | | 90 | 2.8 | 2.6 | 10.2 | 3.28 | +| | | 85 | 2.9 | 2.6 | 10.6 | 16.38 | +| | | 80 | 3.1 | 2.6 | 10.6 | 16.51 | +| | | 75 | 3.3 | 2.6 | 10.6 | 16.53 | +| | | 70 | 3.6 | 2.6 | 10.6 | 16.56 | +| | | 50 | 5.0 | 2.6 | 10.6 | 16.65 | + + +#### C2H +- Latency CMD: duration between cmd issues and axibridge return +- Latency DATA: duration between last data issues and axibridge return +- \* : this latency can be thousands us sometimes, because write latency use bridge channel to reply, single thread can issue around 8M bridge write, when ops exceeds this, the latency increase a lot. + +| CPU | Bytes | Wait cycles | OPS limit (Mops) | Throughput (Mops) | Throughput (GB/s) | Latency CMD (us) | Latency DATA (us) | +| :-----: | :-------: | :-------------: | :------------------: | :-------------------: | :-------------------: | :------------------: | :-------------------: | +| | 64 | 50 | 5 | 4.6 | 0.29 | 1.6 | 1.6 | +| | | 25 | 10 | 5.8 | 0.34 | * | * | +| | | | | | | | | +| | 4 * 1024 | 100 | 2.5 | 2.3 | 9.35 | 1.5 | 1.2 | +| | | 90 | 2.8 | 2.6 | 10.37 | 1.5 | 1.2 | +| | | 85 | 2.9 | 2.7 | 10.96 | 1.6 | 1.3 | +| | | 80 | 3.1 | 2.9 | 11.63 | 1.7 | 1.4 | +| | | 75 | 3.3 | 3.0 | 12.17 | 11.3 | 6.9 | +| | | 70 | 3.6 | 3.0 | 12.17 | 11.3 | 6.9 | +| | | 50 | 5.0 | 3.0 | 12.17 | 11.4 | 7.0 | +| | | | | | | | | +| **GPU** | **Bytes** | **Wait cycles** | **OPS limit (Mops)** | **Throughput (Mops)** | **Throughput (GB/s)** | **Latency CMD (us)** | **Latency DATA (us)** | +| | 64 | 50 | 5 | 0.9 | 0.06 | * | * | +| | | 25 | 10 | 0.9 | 0.06 | * | * | +| | | | | | | | | +| | 4 * 1024 | 100 | 2.5 | 0.7 | 2.92 | * | * | +| | | 90 | 2.8 | 0.7 | 2.93 | * | * | +| | | 85 | 2.9 | 0.7 | 2.96 | * | * | +| | | 80 | 3.1 | 0.7 | 2.96 | * | * | +| | | 75 | 3.3 | 0.7 | 2.95 | * | * | +| | | 70 | 3.6 | 0.7 | 2.73 | * | * | +| | | 50 | 5.0 | 0.7 | 2.73 | * | * | + +--- + +### MMIO +- repeat_times = 16 +- size = 1 * 1024 * 1024 (size of data a thread should write) + +| theads num | speed | +| ---------- | -------- | +| 1 | 0.263201 | +| 2 | 0.766281 | +| 4 | 1.035202 | +| 8 | 2.051150 | +| 16 | 3.851957 | + + + # QDMA C2H bug 1. When running with more than 8 qs, it will always fail. QDMA C2H data port's ready would be down after receiving @@ -295,3 +426,4 @@ around 8M bridge write, when ops exceeds this, the latency increase a lot. 1. All statistics are calculated at the 250M user clock, (so if your speed is 10.6 GB/s at most, maybe you have used 300M user clock). +2. If GPU speed is significantly lower than that of the CPU, tb/press should be run during testing. diff --git a/reset b/reset new file mode 100755 index 0000000..35d55e2 --- /dev/null +++ b/reset @@ -0,0 +1,97 @@ +#!/usr/bin/env python3 + +import argparse +from os import popen, system +import getpass + +def toInt(i): + try: + ret = int(i) + except ValueError: + ret = 0 + return ret + +if __name__ == "__main__": + + if not getpass.getuser() == "root": + print("\033[41mERROR: Please run this script as root!\033[0m") + quit(1) + + parser = argparse.ArgumentParser() + parser.add_argument("-p", "--module_path", default=None, type=str, help="Path of QDMA driver module file, usually ends with qdma-pf.ko") + parser.add_argument("-b", "--bus_id", default=None, type=str, help="Bus ID of FPGA in hex such as `1a` or `40`.") + parser.add_argument("-q", "--queue_num", default=1, type=int, help="Number of queues you wanna create.") + args = parser.parse_args() + + modulePath = args.module_path + busId = args.bus_id + qNum = args.queue_num + + # Check arguments + + if qNum <= 0: + print("\033[41mERROR: `queue_num` should be larger than 0.\033[0m") + quit(1) + if modulePath == None: + # Manually choose module path. + f = popen("locate qdma-pf.ko") + lines = f.readlines() + if len(lines) == 0: + print("\033[41mERROR: No QDMA driver found.\033[0m") + quit(1) + elif len(lines) == 1: + modulePath = lines[0].strip() + else: + print( + "Multiple possible locations of QDMA driver found.\n" + "Please choose the driver you want to load:" + ) + for idx, line in enumerate(lines): + print(f"[\033[36m{idx+1}\033[0m] {line.strip()}") + choice = toInt(input("Enter number here: ")) + while choice <= 0 or choice > len(lines): + print("Invalid input. Please try again.") + choice = toInt(input("Enter number here: ")) + modulePath = lines[choice-1].strip() + print(f"QDMA driver module path: \033[33m{modulePath}\033[0m") + if busId == None: + # Manually choose FPGA bus ID. + f = popen("lspci | grep Xilinx") + lines = f.readlines() + if len(lines) == 0: + print("\033[41mERROR: No FPGA card found.\033[0m") + quit(1) + elif len(lines) == 1: + busId = lines[0][:2] + else: + print( + "Multiple FPGA card found.\n" + "Please choose the card you want to program:" + ) + for idx, line in enumerate(lines): + print(f"[\033[36m{idx+1}\033[0m] {line.strip()}") + choice = toInt(input("Enter number here: ")) + while choice <= 0 or choice > len(lines): + print("Invalid input. Please try again.") + choice = toInt(input("Enter number here: ")) + busId = lines[choice-1][:2] + print(f"FPGA device: \033[33m{busId}:00.0\033[0m") + + # Main execution + print("Please make sure FPGA is running now and not re-programmed yet.") + print("Otherwise FPGA will not run correctly after programmed.") + system(f"cp -f /sys/bus/pci/devices/0000:{busId}:00.0/config ~/myConfig{busId}") + system("rmmod qdma-pf") + print("Now, program the FPGA.") + _ = input("\033[30;46mPress Enter after programming process finishes.\033[0m") + system(f"cp -f ~/myConfig{busId} /sys/bus/pci/devices/0000:{busId}:00.0/config") + system(f"rm ~/myConfig{busId}") + system(f"insmod {modulePath}") + system(f"chmod 666 /sys/bus/pci/devices/0000:{busId}:00.0/qdma/qmax") + system(f"echo 512 > /sys/bus/pci/devices/0000:{busId}:00.0/qdma/qmax") + qid = 0 + while qid < qNum: + system(f"dma-ctl qdma{busId}000 q add idx {qid} mode st dir bi") + system(f"dma-ctl qdma{busId}000 q start idx {qid} dir bi desc_bypass_en pfetch_bypass_en") + qid += 1 + diff --git a/tb/CMakeLists.txt b/tb/CMakeLists.txt new file mode 100644 index 0000000..6ab4333 --- /dev/null +++ b/tb/CMakeLists.txt @@ -0,0 +1,55 @@ +cmake_minimum_required(VERSION 3.10) +set(CMAKE_CXX_STANDARD 17) + +project(qdma_benchmark LANGUAGES CXX) +add_compile_options(-O3 -mavx512f -march=native) + +add_executable(throughput + throughput.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/utils/dma_tool.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/utils/throughput_tool.cpp +) +target_include_directories(throughput PUBLIC + ${CMAKE_CURRENT_SOURCE_DIR}/.. + ${CMAKE_CURRENT_SOURCE_DIR}/utils +) +target_link_libraries(throughput PRIVATE fmt::fmt-header-only) +target_link_libraries(throughput PRIVATE qdma pthread) + + +add_executable(random + random.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/utils/dma_tool.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/utils/random_tool.cpp +) +target_include_directories(random PUBLIC + ${CMAKE_CURRENT_SOURCE_DIR}/.. + ${CMAKE_CURRENT_SOURCE_DIR}/utils +) +target_link_libraries(random PRIVATE fmt::fmt-header-only) +target_link_libraries(random PRIVATE qdma pthread) + + +add_executable(latency + latency.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/utils/dma_tool.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/utils/latency_tool.cpp +) +target_include_directories(latency PUBLIC + ${CMAKE_CURRENT_SOURCE_DIR}/.. + ${CMAKE_CURRENT_SOURCE_DIR}/utils +) +target_link_libraries(latency PRIVATE fmt::fmt-header-only) +target_link_libraries(latency PRIVATE qdma pthread) + + +add_executable(mmio_bridge + mmio_bridge.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/utils/mmio.cpp +) +target_include_directories(mmio_bridge PUBLIC + ${CMAKE_CURRENT_SOURCE_DIR}/.. + ${CMAKE_CURRENT_SOURCE_DIR}/utils +) +target_link_libraries(mmio_bridge PRIVATE fmt::fmt-header-only) +target_link_libraries(mmio_bridge PRIVATE qdma pthread) \ No newline at end of file diff --git a/tb/empty b/tb/empty new file mode 100644 index 0000000..9f01b84 Binary files /dev/null and b/tb/empty differ diff --git a/tb/empty.cu b/tb/empty.cu new file mode 100644 index 0000000..50262d5 --- /dev/null +++ b/tb/empty.cu @@ -0,0 +1,72 @@ +#include +#include + +#include "cuda_runtime.h" + +__global__ void empty(int* d_a, int* d_b, int* d_c) { + // Keep every thread doing arithmetic work to sustain GPU utilization. + int tid = blockIdx.x * blockDim.x + threadIdx.x; + int idx = tid % 100; + int a = d_a[idx]; + int b = d_b[idx]; + + unsigned int acc = static_cast(tid + 1); + for (int j = 0; j < 1000000; j++) { + acc = acc * 1664525u + 1013904223u + static_cast(a + b + idx + j); + } + + d_c[tid] = static_cast(acc); +} + +static volatile sig_atomic_t g_keep_running = 1; + +static void onSignal(int) { g_keep_running = 0; } + +static void checkCuda(cudaError_t err, const char* msg) { + if (err != cudaSuccess) { + std::fprintf(stderr, "%s: %s\n", msg, cudaGetErrorString(err)); + std::exit(1); + } +} + +int main() { + int *d_a, *d_b, *d_c; + constexpr int kDataSize = 100; + constexpr int kThreads = 1024; + constexpr int kBlocks = 1; + + std::signal(SIGINT, onSignal); + + checkCuda(cudaMalloc(&d_a, kDataSize * sizeof(int)), "cudaMalloc d_a failed"); + checkCuda(cudaMalloc(&d_b, kDataSize * sizeof(int)), "cudaMalloc d_b failed"); + checkCuda(cudaMalloc(&d_c, kBlocks * kThreads * sizeof(int)), "cudaMalloc d_c failed"); + + int h_a[kDataSize]; + int h_b[kDataSize]; + for (int i = 0; i < kDataSize; i++) { + h_a[i] = i; + h_b[i] = 100 - i; + } + checkCuda(cudaMemcpy(d_a, h_a, sizeof(h_a), cudaMemcpyHostToDevice), "Memcpy d_a failed"); + checkCuda(cudaMemcpy(d_b, h_b, sizeof(h_b), cudaMemcpyHostToDevice), "Memcpy d_b failed"); + + unsigned long long launch_count = 0; + while (g_keep_running) { + empty<<>>(d_a, d_b, d_c); + checkCuda(cudaGetLastError(), "Kernel launch failed"); + launch_count++; + + // Periodic sync to surface runtime errors while still keeping it busy. + if ((launch_count % 1000ULL) == 0ULL) { + checkCuda(cudaDeviceSynchronize(), "Kernel execution failed"); + } + } + + checkCuda(cudaDeviceSynchronize(), "Final sync failed"); + std::printf("stopped, total launched kernels: %llu\n", launch_count); + + checkCuda(cudaFree(d_a), "cudaFree d_a failed"); + checkCuda(cudaFree(d_b), "cudaFree d_b failed"); + checkCuda(cudaFree(d_c), "cudaFree d_c failed"); + return 0; +} \ No newline at end of file diff --git a/tb/latency.cpp b/tb/latency.cpp new file mode 100644 index 0000000..9da67ba --- /dev/null +++ b/tb/latency.cpp @@ -0,0 +1,19 @@ +#include + +#include "utils/latency_tool.hpp" +#define pci_bus 0x1a + +using namespace std; + +int main() { + // ============= CPU latency benchmark ============== + // cpu_latency_h2c(pci_bus); + // cpu_latency_c2h(pci_bus); + // concurrent_latency(pci_bus); + + // ============= GPU latency benchmark ============== + // gpu_latency_h2c(pci_bus); + gpu_latency_c2h(pci_bus); + + return 0; +} \ No newline at end of file diff --git a/tb/mmio_bridge.cpp b/tb/mmio_bridge.cpp new file mode 100644 index 0000000..b9ce742 --- /dev/null +++ b/tb/mmio_bridge.cpp @@ -0,0 +1,14 @@ +#include +#include "utils/mmio.hpp" +#define pci_bus 0x1a + +using namespace std; + +int main() +{ + + // benchmark_bridge_write(pci_bus, 0); + benchmark_bridge_write(pci_bus, 0); + + return 0; +} \ No newline at end of file diff --git a/tb/press b/tb/press new file mode 100644 index 0000000..58f731f Binary files /dev/null and b/tb/press differ diff --git a/tb/random.cpp b/tb/random.cpp new file mode 100644 index 0000000..69c9490 --- /dev/null +++ b/tb/random.cpp @@ -0,0 +1,19 @@ +#include + +#include "utils/random_tool.hpp" +#define pci_bus 0x1a + +using namespace std; + +int main() { + // ============= CPU Random benchmark ============== + // cpu_random_h2c(pci_bus); + // cpu_random_c2h(pci_bus); + // concurrent_random(pci_bus); + + // ============= GPU Random benchmark ============== + gpu_random_h2c(pci_bus); + gpu_random_c2h(pci_bus); + + return 0; +} \ No newline at end of file diff --git a/tb/throughput.cpp b/tb/throughput.cpp new file mode 100644 index 0000000..f9ca743 --- /dev/null +++ b/tb/throughput.cpp @@ -0,0 +1,13 @@ +#include "utils/throughput_tool.hpp" +#define pci_bus 0x1a + +int main() { + // ============= CPU throughput benchmark ============== + // cpu_throughput_h2c(pci_bus); + // cpu_throughput_c2h(pci_bus); + + // ============= GPU throughput benchmark ============== + gpu_throughput_h2c(pci_bus); + gpu_throughput_c2h(pci_bus); + return 0; +} \ No newline at end of file diff --git a/tb/utils/dma_tool.cpp b/tb/utils/dma_tool.cpp new file mode 100644 index 0000000..d583a2c --- /dev/null +++ b/tb/utils/dma_tool.cpp @@ -0,0 +1,66 @@ +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +void throughput_benchmark_print_counters(FPGACtl* fpga_ctl) { + // Print counters, copied from code generated by Chisel project. + + // Report width 32: + fmt::println("tlb.io.tlb_miss_count : {}", fpga_ctl->readReg(512 + 400)); + fmt::println("io.c2h_cmd.[fire] : {}", fpga_ctl->readReg(512 + 401)); + fmt::println("io.h2c_cmd.[fire] : {}", fpga_ctl->readReg(512 + 402)); + fmt::println("io.c2h_data.[fire] : {}", fpga_ctl->readReg(512 + 403)); + fmt::println("io.h2c_data.[fire] : {}", fpga_ctl->readReg(512 + 404)); + fmt::println("fifo_c2h_cmd.io.out.[fire] : {}", fpga_ctl->readReg(512 + 405)); + fmt::println("fifo_h2c_cmd.io.out.[fire] : {}", fpga_ctl->readReg(512 + 406)); + fmt::println("fifo_c2h_data.io.out.[fire] : {}", fpga_ctl->readReg(512 + 407)); + fmt::println("fifo_h2c_data.io.in.[fire] : {}", fpga_ctl->readReg(512 + 408)); + + // Report width 1: + fmt::println("fifo_c2h_cmd.io.out.valid : {}", (fpga_ctl->readReg(512 + 409) >> 0) & 1); + fmt::println("fifo_c2h_cmd.io.out.ready : {}", (fpga_ctl->readReg(512 + 409) >> 1) & 1); + fmt::println("fifo_h2c_cmd.io.out.valid : {}", (fpga_ctl->readReg(512 + 409) >> 2) & 1); + fmt::println("fifo_h2c_cmd.io.out.ready : {}", (fpga_ctl->readReg(512 + 409) >> 3) & 1); + fmt::println("fifo_c2h_data.io.out.valid : {}", (fpga_ctl->readReg(512 + 409) >> 4) & 1); + fmt::println("fifo_c2h_data.io.out.ready : {}", (fpga_ctl->readReg(512 + 409) >> 5) & 1); + fmt::println("fifo_h2c_data.io.in.valid : {}", (fpga_ctl->readReg(512 + 409) >> 6) & 1); + fmt::println("fifo_h2c_data.io.in.ready : {}", (fpga_ctl->readReg(512 + 409) >> 7) & 1); +} + +void reset_counters(FPGACtl* fpga_ctl) { + fpga_ctl->writeReg(14, 1); + fpga_ctl->writeReg(14, 0); +} + + +void random_benchmark_print_counters(FPGACtl *fpga_ctl) { + // Print counters, copied from code generated by Chisel project. + + //Report width 32: + fmt::println("tlb.io.tlb_miss_count : {}", fpga_ctl->readReg(512+300)); + fmt::println("io.c2h_cmd.[fire] : {}", fpga_ctl->readReg(512+301)); + fmt::println("io.h2c_cmd.[fire] : {}", fpga_ctl->readReg(512+302)); + fmt::println("io.c2h_data.[fire] : {}", fpga_ctl->readReg(512+303)); + fmt::println("io.h2c_data.[fire] : {}", fpga_ctl->readReg(512+304)); + fmt::println("fifo_c2h_cmd.io.out.[fire] : {}", fpga_ctl->readReg(512+305)); + fmt::println("fifo_h2c_cmd.io.out.[fire] : {}", fpga_ctl->readReg(512+306)); + fmt::println("fifo_c2h_data.io.out.[fire] : {}", fpga_ctl->readReg(512+307)); + fmt::println("fifo_h2c_data.io.in.[fire] : {}", fpga_ctl->readReg(512+308)); + + //Report width 1: + fmt::println("fifo_c2h_cmd.io.out.valid : {}", (fpga_ctl->readReg(512+309) >> 0) & 1); + fmt::println("fifo_c2h_cmd.io.out.ready : {}", (fpga_ctl->readReg(512+309) >> 1) & 1); + fmt::println("fifo_h2c_cmd.io.out.valid : {}", (fpga_ctl->readReg(512+309) >> 2) & 1); + fmt::println("fifo_h2c_cmd.io.out.ready : {}", (fpga_ctl->readReg(512+309) >> 3) & 1); + fmt::println("fifo_c2h_data.io.out.valid : {}", (fpga_ctl->readReg(512+309) >> 4) & 1); + fmt::println("fifo_c2h_data.io.out.ready : {}", (fpga_ctl->readReg(512+309) >> 5) & 1); + fmt::println("fifo_h2c_data.io.in.valid : {}", (fpga_ctl->readReg(512+309) >> 6) & 1); + fmt::println("fifo_h2c_data.io.in.ready : {}", (fpga_ctl->readReg(512+309) >> 7) & 1); +} \ No newline at end of file diff --git a/tb/utils/dma_tool.hpp b/tb/utils/dma_tool.hpp new file mode 100644 index 0000000..d3e62b8 --- /dev/null +++ b/tb/utils/dma_tool.hpp @@ -0,0 +1,7 @@ +#include + +void throughput_benchmark_print_counters(FPGACtl* fpga_ctl); + +void reset_counters(FPGACtl* fpga_ctl); + +void random_benchmark_print_counters(FPGACtl* fpga_ctl); \ No newline at end of file diff --git a/tb/utils/latency_tool.cpp b/tb/utils/latency_tool.cpp new file mode 100644 index 0000000..0c96e54 --- /dev/null +++ b/tb/utils/latency_tool.cpp @@ -0,0 +1,443 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "dma_tool.hpp" + +void cpu_latency_h2c(uint8_t pci_bus) { + fmt::println("=====H2C latency benchmark start====="); + + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + + auto cpu_mem_ctl = CPUMemCtl::getInstance(1UL * 1024 * 1024 * 1024); + + size_t size = 1UL * 1024 * 1024 * 1024; + + cpu_mem_ctl->writeTLB([=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }); + + auto dma_buff = cpu_mem_ctl->alloc(size); + volatile auto p = (uint32_t*)dma_buff; + + uint32_t length = 1 * 4 * 1024; + uint32_t total_cmds = 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + uint32_t wait_cycles = 50; // 100=2.5Mops,when 4K burst, 100=10GB/s + + // initial dma buffer + // FPGA: 512-bit = 64Byte = 16 * uint32_t + for (int i = 0; i < total_words; i++) { + for (int j = 0; j < 16; j++) { + p[i * 16 + j] = i; + } + } + + fpga_ctl->writeReg(100, (uint32_t)((unsigned long)p >> 32)); + fpga_ctl->writeReg(101, (uint32_t)((unsigned long)p)); + fpga_ctl->writeReg(102, length); + fpga_ctl->writeReg(104, total_cmds); + fpga_ctl->writeReg(105, total_words); + fpga_ctl->writeReg(106, wait_cycles); + + // start + fpga_ctl->writeReg(103, 0); + fpga_ctl->writeReg(103, 1); + + sleep(3); + unsigned int cycles = fpga_ctl->readReg(512 + 103); + fmt::print("\n"); + fmt::println("burst length: {}", length); + + fmt::println("count_total_words: 0x{:x}, shoule be 0x{:x}", fpga_ctl->readReg(512 + 100), total_words); + fmt::println("count_send_cmd: 0x{:x}, shoule be 0x{:x}", fpga_ctl->readReg(512 + 101), total_cmds); + fmt::println("count_err_data: 0x{:x}, should be 0x0", fpga_ctl->readReg(512 + 102)); + + fmt::println("Cycles: {}", cycles); + double speed = 1.0 * length * total_cmds / (1.0 * cycles * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + fmt::println("Total length: {}", total_words * 64); + fmt::println("Speed: {:.1f} GB/s", speed); + + size_t count_latency = (((size_t)fpga_ctl->readReg(512 + 105)) << 32) + fpga_ctl->readReg(512 + 104); + fmt::println("count_latency 0x{:x}", count_latency); + fmt::println("wait cycles {}", wait_cycles); + double average_latency = 1.0 * count_latency * 4 / total_cmds / 1000; // us + fmt::println("average_latency {:.2f} us", average_latency); + double ops_limit = 1.0 * 250 * 1024 * 1024 / wait_cycles / 1024 / 1024; // Mps + fmt::println("ops_limit {:.1f} Mops", ops_limit); + double ops = 1.0 * total_cmds / (1.0 * cycles * 4 / 1000 / 1000 / 1000) / 1024 / 1024; // Mps + fmt::println("ops {:.1f} Mops", ops); + fmt::print("\n"); + // printCounters(); + cpu_mem_ctl->free(dma_buff); +} + +void cpu_latency_c2h(uint8_t pci_bus) { + fmt::println("=====C2H latency benchmark start====="); + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + + auto cpu_mem_ctl = CPUMemCtl::getInstance(1UL * 1024 * 1024 * 1024); + + size_t size = 1L * 1024 * 1024 * 1024; + + cpu_mem_ctl->writeTLB([=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }); + + auto dma_buff = cpu_mem_ctl->alloc(size); + volatile auto p = (uint32_t*)dma_buff; + + memset(p, 0, size); + + // the num of Bytes(promise) per command + // uint32_t: define how to explain the data(how many bit per step) + uint32_t length = 1 * 4 * 1024; + uint32_t total_cmds = 1 * 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + uint32_t wait_cycles = 50; // 100=2.5Mops,when 4K burst, 100=10GB/s + + fpga_ctl->writeReg(200, (uint32_t)((unsigned long)p >> 32)); + fpga_ctl->writeReg(201, (uint32_t)((unsigned long)p)); + fpga_ctl->writeReg(202, length); + fpga_ctl->writeReg(204, total_words); + fpga_ctl->writeReg(205, total_cmds); + fpga_ctl->writeReg(207, wait_cycles); + + fpga_ctl->writeConfig(0x1408 / 4, 0); + uint32_t tag = fpga_ctl->readConfig(0x140c / 4); + fpga_ctl->writeReg(206, tag); + + fmt::println("{}", tag & 0x7f); + + // length: Bytes per command + // FPGA: 512bits = 64Bytes + // beats: how many 512bits data in one command + int beats = length / 64; + volatile uint32_t* p_ack = p; + + fpga_ctl->writeReg(203, 0); // start + fpga_ctl->writeReg(203, 1); + + for (int i = 0; i < total_cmds; i++) { + uint32_t verifiy_value = i * beats; + + while (true) { + bool done = p_ack[i * beats * 16] == verifiy_value; + // printf("p_ack[i * beats * 16]: %08x, verify_value: %08x\n", p_ack[i * beats * 16], verifiy_value); + if (done) { + break; + } + } + fpga_ctl->writeBridge(0, {1, 1, 1, 1, 1, 1, 1, 1}); + } + + uint32_t count_cmds = fpga_ctl->readReg(512 + 200); + uint32_t count_words = fpga_ctl->readReg(512 + 201); + uint32_t count_time = fpga_ctl->readReg(512 + 202); + uint32_t count_recv_ack = fpga_ctl->readReg(512 + 207); + + uint32_t right_count = 0; + uint32_t wrong_count = 0; + uint32_t value_verify = 0; + + for (int i = 0; i < 16384; i++, value_verify++) { + uint32_t val = p[i * 16]; + if (value_verify != val) { + wrong_count++; + continue; + } + + bool right = true; + for (int j = 1; j < 16; j++) { + if (p[i * 16 + j] != val) { + wrong_count++; + right = false; + break; + } + } + + if (right) { + right_count++; + } + } + + fmt::println("burst length: {}", length); + fmt::println("count_cmds: 0x{:x},should be: 0x{:x}", count_cmds, total_cmds); + fmt::println("count_recv_ack: 0x{:x},should be: 0x{:x}", count_recv_ack, total_cmds); + fmt::println("count_words: 0x{:x},should be: 0x{:x}", count_words, total_words); + fmt::println("count_error: 0x{:x},shoule be: 0x0", wrong_count); + fmt::println("count time: {}", count_time); + + double speed = 1.0 * length * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + fmt::println("Speed: {:.2f} GB/s", speed); + + size_t count_latency_cmd = (((size_t)fpga_ctl->readReg(512 + 204)) << 32) + fpga_ctl->readReg(512 + 203); + size_t count_latency_data = (((size_t)fpga_ctl->readReg(512 + 206)) << 32) + fpga_ctl->readReg(512 + 205); + fmt::println("count_latency_cmd 0x{:x}", count_latency_cmd); + fmt::println("count_latency_data 0x{:x}", count_latency_data); + + double average_latency_cmd = 1.0 * count_latency_cmd * 4 / total_cmds / 1000; // us + double average_latency_data = 1.0 * count_latency_data * 4 / total_cmds / 1000; // us + double ops_limit = 1.0 * 250 * 1024 * 1024 / wait_cycles / 1024 / 1024; // Mps + double ops = 1.0 * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024; // Mps + fmt::println("wait cycles {}", wait_cycles); + fmt::println("average_latency_cmd {:.1f} us", average_latency_cmd); + fmt::println("average_latency_data {:.1f} us", average_latency_data); + fmt::println("ops_limit {:.1f} Mops", ops_limit); + fmt::println("ops {:.1f} Mops", ops); + // printCounters(); + cpu_mem_ctl->free(dma_buff); +} + +void gpu_latency_h2c(uint8_t pci_bus) { + fmt::println("=====GPU H2C latency benchmark start====="); + + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + + auto gpu_mem_ctl = GPUMemCtl::getInstance(0, 1UL * 1024 * 1024 * 1024); + + size_t pool_size = 1UL * 1024 * 1024 * 1024; + + gpu_mem_ctl->writeTLB( + [=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }, + true); + + // get buff (gpu vaddr) + auto dma_buff = gpu_mem_ctl->alloc(pool_size); + // gpu pool's base in gpu vaddr + auto gpu_base = reinterpret_cast(gpu_mem_ctl->getDevPtr()); + auto host_base = static_cast(gpu_mem_ctl->getMapDevPtr()); + volatile auto p = host_base + (reinterpret_cast(dma_buff) - gpu_base) / sizeof(uint32_t); + + uint32_t length = 1 * 4 * 1024; + uint32_t total_cmds = 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + uint32_t wait_cycles = 50; // 100=2.5Mops,when 4K burst, 100=10GB/s + + // initial dma buffer + // FPGA: 512-bit = 64Byte = 16 * uint32_t + for (int i = 0; i < total_words; i++) { + for (int j = 0; j < 16; j++) { + p[i * 16 + j] = i; + } + } + + auto fpga_vaddr = reinterpret_cast(dma_buff); + fpga_ctl->writeReg(100, (uint32_t)(fpga_vaddr >> 32)); + fpga_ctl->writeReg(101, (uint32_t)(fpga_vaddr)); + fpga_ctl->writeReg(102, length); + fpga_ctl->writeReg(104, total_cmds); + fpga_ctl->writeReg(105, total_words); + fpga_ctl->writeReg(106, wait_cycles); + + // start + fpga_ctl->writeReg(103, 0); + fpga_ctl->writeReg(103, 1); + + sleep(3); + unsigned int cycles = fpga_ctl->readReg(512 + 103); + fmt::print("\n"); + fmt::println("burst length: {}", length); + + fmt::println("count_total_words: 0x{:x}, shoule be 0x{:x}", fpga_ctl->readReg(512 + 100), total_words); + fmt::println("count_send_cmd: 0x{:x}, shoule be 0x{:x}", fpga_ctl->readReg(512 + 101), total_cmds); + fmt::println("count_err_data: 0x{:x}, should be 0x0", fpga_ctl->readReg(512 + 102)); + + fmt::println("Cycles: {}", cycles); + fmt::println("Total length: {}", total_words * 64); + + size_t count_latency = (((size_t)fpga_ctl->readReg(512 + 105)) << 32) + fpga_ctl->readReg(512 + 104); + fmt::println("count_latency 0x{:x}", count_latency); + fmt::println("wait cycles {}", wait_cycles); + + double ops_limit = 1.0 * 250 * 1024 * 1024 / wait_cycles / 1024 / 1024; // Mps + fmt::println("ops_limit {:.1f} Mops", ops_limit); + double ops = 1.0 * total_cmds / (1.0 * cycles * 4 / 1000 / 1000 / 1000) / 1024 / 1024; // Mps + fmt::println("ops {:.1f} Mops", ops); + + double speed = 1.0 * length * total_cmds / (1.0 * cycles * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + fmt::println("Speed: {:.1f} GB/s", speed); + + double average_latency = 1.0 * count_latency * 4 / total_cmds / 1000; // us + fmt::println("average_latency {:.2f} us", average_latency); + + fmt::print("\n"); + // printCounters(); + gpu_mem_ctl->free(dma_buff); + GPUMemCtl::cleanCtx(); +} + +void gpu_latency_c2h(uint8_t pci_bus) { + fmt::println("=====GPU C2H latency benchmark start====="); + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + + auto gpu_mem_ctl = GPUMemCtl::getInstance(0, 1UL * 1024 * 1024 * 1024); + + size_t pool_size = 1L * 1024 * 1024 * 1024; + + gpu_mem_ctl->writeTLB( + [=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }, + true); + + auto dma_buff = gpu_mem_ctl->alloc(pool_size); + auto gpu_base = reinterpret_cast(gpu_mem_ctl->getDevPtr()); + auto host_base = static_cast(gpu_mem_ctl->getMapDevPtr()); + volatile auto p = host_base + (reinterpret_cast(dma_buff) - gpu_base) / sizeof(uint32_t); + + memset(p, 0, pool_size); + + // the num of Bytes(promise) per command + // uint32_t: define how to explain the data(how many bit per step) + uint32_t length = 1 * 4 * 1024; + uint32_t total_cmds = 1 * 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + uint32_t wait_cycles = 100; // 100=2.5Mops,when 4K burst, 100=10GB/s + + auto fpga_vaddr = reinterpret_cast(dma_buff); + fpga_ctl->writeReg(200, (uint32_t)(fpga_vaddr >> 32)); + fpga_ctl->writeReg(201, (uint32_t)(fpga_vaddr)); + fpga_ctl->writeReg(202, length); + fpga_ctl->writeReg(204, total_words); + fpga_ctl->writeReg(205, total_cmds); + fpga_ctl->writeReg(207, wait_cycles); + + fpga_ctl->writeConfig(0x1408 / 4, 0); + uint32_t tag = fpga_ctl->readConfig(0x140c / 4); + fpga_ctl->writeReg(206, tag); + + fmt::println("{}", tag & 0x7f); + + // length: Bytes per command + // FPGA: 512bits = 64Bytes + // beats: how many 512bits data in one command + int beats = length / 64; + volatile uint32_t* p_ack = p; + + fpga_ctl->writeReg(203, 0); // start + fpga_ctl->writeReg(203, 1); + + for (int i = 0; i < total_cmds; i++) { + uint32_t verifiy_value = i * beats; + + while (true) { + bool done = p_ack[i * beats * 16] == verifiy_value; + // printf("p_ack[i * beats * 16]: %08x, verify_value: %08x\n", p_ack[i * beats * 16], verifiy_value); + if (done) { + break; + } + } + fpga_ctl->writeBridge(0, {1, 1, 1, 1, 1, 1, 1, 1}); + } + + uint32_t count_cmds = fpga_ctl->readReg(512 + 200); + uint32_t count_words = fpga_ctl->readReg(512 + 201); + uint32_t count_time = fpga_ctl->readReg(512 + 202); + uint32_t count_recv_ack = fpga_ctl->readReg(512 + 207); + + uint32_t right_count = 0; + uint32_t wrong_count = 0; + uint32_t value_verify = 0; + + for (int i = 0; i < 16384; i++, value_verify++) { + uint32_t val = p[i * 16]; + if (value_verify != val) { + wrong_count++; + continue; + } + + bool right = true; + for (int j = 1; j < 16; j++) { + if (p[i * 16 + j] != val) { + wrong_count++; + right = false; + break; + } + } + + if (right) { + right_count++; + } + } + + // for (int i = 0; i < 4; i++) { + // printf("Packet [%03d]: ", i); + // for (int j = 0; j < 16; j++) { + // printf("%08x ", p[i * 16 + j]); + // } + // printf("\n"); + // } + + fmt::println("burst length: {}", length); + fmt::println("count_cmds: 0x{:x},should be: 0x{:x}", count_cmds, total_cmds); + fmt::println("count_recv_ack: 0x{:x},should be: 0x{:x}", count_recv_ack, total_cmds); + fmt::println("count_words: 0x{:x},should be: 0x{:x}", count_words, total_words); + fmt::println("count_error: 0x{:x},shoule be: 0x0", wrong_count); + fmt::println("count time: {}", count_time); + + size_t count_latency_cmd = (((size_t)fpga_ctl->readReg(512 + 204)) << 32) + fpga_ctl->readReg(512 + 203); + size_t count_latency_data = (((size_t)fpga_ctl->readReg(512 + 206)) << 32) + fpga_ctl->readReg(512 + 205); + fmt::println("count_latency_cmd 0x{:x}", count_latency_cmd); + fmt::println("count_latency_data 0x{:x}", count_latency_data); + + double ops_limit = 1.0 * 250 * 1024 * 1024 / wait_cycles / 1024 / 1024; // Mps + double ops = 1.0 * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024; // Mps + double speed = 1.0 * length * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + double average_latency_cmd = 1.0 * count_latency_cmd * 4 / total_cmds / 1000; // us + double average_latency_data = 1.0 * count_latency_data * 4 / total_cmds / 1000; // us + + fmt::println("wait cycles {}", wait_cycles); + fmt::println("ops_limit {:.1f} Mops", ops_limit); + fmt::println("ops {:.1f} Mops", ops); + fmt::println("Speed: {:.2f} GB/s", speed); + fmt::println("average_latency_cmd {:.1f} us", average_latency_cmd); + fmt::println("average_latency_data {:.1f} us", average_latency_data); + // printCounters(); + gpu_mem_ctl->free(dma_buff); + GPUMemCtl::cleanCtx(); +} \ No newline at end of file diff --git a/tb/utils/latency_tool.hpp b/tb/utils/latency_tool.hpp new file mode 100644 index 0000000..edcd367 --- /dev/null +++ b/tb/utils/latency_tool.hpp @@ -0,0 +1,7 @@ +#include + +void cpu_latency_h2c(uint8_t pci_bus); +void cpu_latency_c2h(uint8_t pci_bus); + +void gpu_latency_h2c(uint8_t pci_bus); +void gpu_latency_c2h(uint8_t pci_bus); \ No newline at end of file diff --git a/tb/utils/mmio.cpp b/tb/utils/mmio.cpp new file mode 100644 index 0000000..a6dd485 --- /dev/null +++ b/tb/utils/mmio.cpp @@ -0,0 +1,84 @@ +#include "mmio.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +void* write_bridge_sub(void* args) { + __m512i data; + for (int i = 0; i < 8; i++) { + data[i] = 1; + } + + size_t addr, size, repeat_times; + addr = ((size_t*)args)[0]; + size = ((size_t*)args)[1]; + repeat_times = ((size_t*)args)[2]; + + for (size_t t = 0; t < repeat_times; t++) { + for (size_t i = addr; i < addr + size; i += 512) { + _mm512_stream_si512((__m512i*)(i + 64 * 0), data); + _mm512_stream_si512((__m512i*)(i + 64 * 1), data); + _mm512_stream_si512((__m512i*)(i + 64 * 2), data); + _mm512_stream_si512((__m512i*)(i + 64 * 3), data); + _mm512_stream_si512((__m512i*)(i + 64 * 4), data); + _mm512_stream_si512((__m512i*)(i + 64 * 5), data); + _mm512_stream_si512((__m512i*)(i + 64 * 6), data); + _mm512_stream_si512((__m512i*)(i + 64 * 7), data); + } + } + + return 0; +} + +void benchmark_bridge_write(uint8_t pci_bus, uint8_t is_wc) { + size_t size = 1 * 1024 * 1024 * 1024; // size of data that a single thread should write (in bytes) + int num_threads = 64, repeat_times = 16; + pthread_t tids[num_threads]; + size_t args[num_threads][3]; + + fmt::println("=====MMIO Bridge Write Benchmark====="); + + FPGACtl::explictInit(pci_bus, size * num_threads, is_wc); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + void* bridge = (void*)(fpga_ctl->getBridgeAddr()); + printf("Bridge: %lx\n", (size_t)bridge); + + // initialize params + for (int i = 0; i < num_threads; i++) { + args[i][0] = (size_t)bridge + i * size; + // args[i][0] = (size_t)bridge; + args[i][1] = size; + args[i][2] = (size_t)repeat_times; + } + + printf("Starting %d threads, repeating %d times with writecombine %s.\n", num_threads, repeat_times, is_wc ? "enabled" : "disabled"); + + // start timer + struct timespec start_timer, end_timer; + clock_gettime(CLOCK_MONOTONIC, &start_timer); + + for (int i = 0; i < num_threads; i++) { + int ret = pthread_create(&tids[i], NULL, write_bridge_sub, args[i]); + if (ret != 0) { + cout << "pthread_create error: error_code=" << ret << endl; + } + } + for (int i = 0; i < num_threads; i++) { + pthread_join(tids[i], NULL); + } + + // end timer + clock_gettime(CLOCK_MONOTONIC, &end_timer); + double time = (end_timer.tv_sec - start_timer.tv_sec) + 1.0 * (end_timer.tv_nsec - start_timer.tv_nsec) / 1e9; + printf("time: %f s\n", time); + printf("speed: %f GB/s\n", size * num_threads * repeat_times / time / 1024 / 1024 / 1024); +} \ No newline at end of file diff --git a/tb/utils/mmio.hpp b/tb/utils/mmio.hpp new file mode 100644 index 0000000..a75cff8 --- /dev/null +++ b/tb/utils/mmio.hpp @@ -0,0 +1,13 @@ +#ifndef mmio_hpp +#define mmio_hpp + +#include +#include +#include +#include + +using namespace std; + +void benchmark_bridge_write(uint8_t, uint8_t); + +#endif \ No newline at end of file diff --git a/tb/utils/random_tool.cpp b/tb/utils/random_tool.cpp new file mode 100644 index 0000000..fe984a4 --- /dev/null +++ b/tb/utils/random_tool.cpp @@ -0,0 +1,327 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "dma_tool.hpp" + +void cpu_random_h2c(uint8_t pci_bus) { + fmt::println("=====CPU H2C random benchmark start====="); + + size_t size = 1UL * 1024 * 1024 * 1024; + + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + + auto cpu_mem_ctl = CPUMemCtl::getInstance(1UL * 1024 * 1024 * 1024); + + cpu_mem_ctl->writeTLB([=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }); + + auto dma_buff = cpu_mem_ctl->alloc(size); + volatile auto p = (uint32_t*)dma_buff; + + uint32_t length = 1024; + uint32_t busrt_length_shift = (uint32_t)log2(length); + uint32_t total_cmds = 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + + for (int i = 0; i < size / 64; i++) { + p[i * 16] = i * 64; + } + + fpga_ctl->writeReg(100, (uint32_t)((unsigned long)p >> 32)); + fpga_ctl->writeReg(101, (uint32_t)((unsigned long)p)); + fpga_ctl->writeReg(102, length); + fpga_ctl->writeReg(103, busrt_length_shift); + fpga_ctl->writeReg(105, total_words); + fpga_ctl->writeReg(106, total_cmds); + + reset_counters(fpga_ctl); + + // start + fpga_ctl->writeReg(104, 0); + fpga_ctl->writeReg(104, 1); + + sleep(1); + + uint32_t count_words = fpga_ctl->readReg(512 + 100); + uint32_t count_cmds = fpga_ctl->readReg(512 + 102); + uint32_t count_time = fpga_ctl->readReg(512 + 103); + + fmt::print("\n"); + fmt::println("count_total_words: 0x{}", count_words); + fmt::println("count_err_data: 0x{}, should be 0x0", fpga_ctl->readReg(512 + 101)); + fmt::println("count_send_cmd: 0x{}", count_cmds); + + fmt::println("Cycles: {}", count_time); + + double ops = 1.0 * count_cmds * 1e9 / (1.0 * count_time * 4.0); + fmt::println("OPS: {:.3f} MOPS", ops / 1e6); + + double speed = 1.0 * length * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + + fmt::println("Total length: {}", total_words * 64); + fmt::println("Speed: {:.2f} GB/s", speed); + fmt::print("\n"); + + cpu_mem_ctl->free(dma_buff); + + random_benchmark_print_counters(fpga_ctl); +} + +void cpu_random_c2h(uint8_t pci_bus) { + fmt::println("=====CPU C2H random benchmark start====="); + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + + auto cpu_mem_ctl = CPUMemCtl::getInstance(1UL * 1024 * 1024 * 1024); + + size_t size = 1L * 1024 * 1024 * 1024; + + cpu_mem_ctl->writeTLB([=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }); + + auto dma_buff = cpu_mem_ctl->alloc(size); + volatile auto p = (uint32_t*)dma_buff; + + memset(p, 0, size); + + uint32_t length = 64; // 32K has ever triggered the horrible bug + uint32_t busrt_length_shift = (uint32_t)log2(length); + uint32_t total_cmds = 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + + fpga_ctl->writeReg(200, (uint32_t)((unsigned long)p >> 32)); + fpga_ctl->writeReg(201, (uint32_t)((unsigned long)p)); + fpga_ctl->writeReg(202, length); + fpga_ctl->writeReg(203, busrt_length_shift); + fpga_ctl->writeReg(205, total_words); + fpga_ctl->writeReg(206, total_cmds); + + reset_counters(fpga_ctl); + + fpga_ctl->writeConfig(0x1408 / 4, 0); + uint32_t tag = fpga_ctl->readConfig(0x140c / 4); + fpga_ctl->writeReg(207, tag); + fmt::println("{}", tag & 0x7f); + + fpga_ctl->writeReg(204, 0); // start + fpga_ctl->writeReg(204, 1); + + sleep(2); + uint32_t count_cmds = fpga_ctl->readReg(512 + 200); + uint32_t count_words = fpga_ctl->readReg(512 + 201); + uint32_t count_time = fpga_ctl->readReg(512 + 202); + + fmt::println("count cmd: {},right: {}", count_cmds, total_cmds); + fmt::println("count word: {},right: {}", count_words, total_words); + fmt::println("count time: {}", count_time); + + double ops = 1.0 * count_cmds * 1e9 / (1.0 * count_time * 4.0); + fmt::println("OPS: {:.3f} MOPS", ops / 1e6); + + double speed = 1.0 * length * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + fmt::println("Speed: {:.2f} GB/s", speed); + + int count_written_right_word = 0; + for (int i = 0; i < size / 64; i++) { + if (p[i * 16] == i * 64) { + count_written_right_word++; + } + } + fmt::println("count_written_right_word:{}, total words:{}", count_written_right_word, total_words); + if (count_time > 2 * 603700) { + random_benchmark_print_counters(fpga_ctl); + } + + GPUMemCtl::cleanCtx(); +} + +void gpu_random_h2c(uint8_t pci_bus) { + fmt::println("=====GPU H2C random benchmark start====="); + + size_t pool_size = 1UL * 1024 * 1024 * 1024; + + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + + auto gpu_mem_ctl = GPUMemCtl::getInstance(0, 1UL * 1024 * 1024 * 1024); + + gpu_mem_ctl->writeTLB( + [=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }, + true); + + // get buff (gpu vaddr) + auto dma_buff = gpu_mem_ctl->alloc(pool_size); + // gpu pool's base in gpu vaddr + auto gpu_base = reinterpret_cast(gpu_mem_ctl->getDevPtr()); + auto host_base = static_cast(gpu_mem_ctl->getMapDevPtr()); + volatile auto p = host_base + (reinterpret_cast(dma_buff) - gpu_base) / sizeof(uint32_t); + + uint32_t length = 1024; + uint32_t busrt_length_shift = (uint32_t)log2(length); + uint32_t total_cmds = 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + + for (int i = 0; i < pool_size / 64; i++) { + p[i * 16] = i * 64; + } + + auto fpga_vaddr = reinterpret_cast(dma_buff); + fpga_ctl->writeReg(100, (uint32_t)(fpga_vaddr >> 32)); + fpga_ctl->writeReg(101, (uint32_t)(fpga_vaddr)); + fpga_ctl->writeReg(102, length); + fpga_ctl->writeReg(103, busrt_length_shift); + fpga_ctl->writeReg(105, total_words); + fpga_ctl->writeReg(106, total_cmds); + + reset_counters(fpga_ctl); + + // start + fpga_ctl->writeReg(104, 0); + fpga_ctl->writeReg(104, 1); + + sleep(1); + + uint32_t count_words = fpga_ctl->readReg(512 + 100); + uint32_t count_cmds = fpga_ctl->readReg(512 + 102); + uint32_t count_time = fpga_ctl->readReg(512 + 103); + + fmt::print("\n"); + fmt::println("count_total_words: 0x{}", count_words); + fmt::println("count_err_data: 0x{}, should be 0x0", fpga_ctl->readReg(512 + 101)); + fmt::println("count_send_cmd: 0x{}", count_cmds); + + fmt::println("Cycles: {}", count_time); + + double ops = 1.0 * count_cmds * 1e9 / (1.0 * count_time * 4.0); + fmt::println("OPS: {:.3f} MOPS", ops / 1e6); + + double speed = 1.0 * length * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + + fmt::println("Total length: {}", total_words * 64); + fmt::println("Speed: {:.2f} GB/s", speed); + fmt::print("\n"); + + gpu_mem_ctl->free(dma_buff); + + random_benchmark_print_counters(fpga_ctl); + GPUMemCtl::cleanCtx(); +} + +void gpu_random_c2h(uint8_t pci_bus) { + fmt::println("=====GPU C2H random benchmark start====="); + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + + auto gpu_mem_ctl = GPUMemCtl::getInstance(0, 1UL * 1024 * 1024 * 1024); + + size_t pool_size = 1L * 1024 * 1024 * 1024; + + gpu_mem_ctl->writeTLB( + [=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }, + true); + + auto dma_buff = gpu_mem_ctl->alloc(pool_size); + auto gpu_base = reinterpret_cast(gpu_mem_ctl->getDevPtr()); + auto host_base = static_cast(gpu_mem_ctl->getMapDevPtr()); + volatile auto p = host_base + (reinterpret_cast(dma_buff) - gpu_base) / sizeof(uint32_t); + + memset(p, 0, pool_size); + + uint32_t length = 1024; // 32K has ever triggered the horrible bug + uint32_t busrt_length_shift = (uint32_t)log2(length); + uint32_t total_cmds = 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + + auto fpga_vaddr = reinterpret_cast(dma_buff); + fpga_ctl->writeReg(200, (uint32_t)(fpga_vaddr >> 32)); + fpga_ctl->writeReg(201, (uint32_t)(fpga_vaddr)); + fpga_ctl->writeReg(202, length); + fpga_ctl->writeReg(203, busrt_length_shift); + fpga_ctl->writeReg(205, total_words); + fpga_ctl->writeReg(206, total_cmds); + + reset_counters(fpga_ctl); + + fpga_ctl->writeConfig(0x1408 / 4, 0); + uint32_t tag = fpga_ctl->readConfig(0x140c / 4); + fpga_ctl->writeReg(207, tag); + fmt::println("{}", tag & 0x7f); + + fpga_ctl->writeReg(204, 0); // start + fpga_ctl->writeReg(204, 1); + + sleep(2); + uint32_t count_cmds = fpga_ctl->readReg(512 + 200); + uint32_t count_words = fpga_ctl->readReg(512 + 201); + uint32_t count_time = fpga_ctl->readReg(512 + 202); + + fmt::println("count cmd: {},right: {}", count_cmds, total_cmds); + fmt::println("count word: {},right: {}", count_words, total_words); + fmt::println("count time: {}", count_time); + + double ops = 1.0 * count_cmds * 1e9 / (1.0 * count_time * 4.0); + fmt::println("OPS: {:.3f} MOPS", ops / 1e6); + + double speed = 1.0 * length * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + fmt::println("Speed: {:.2f} GB/s", speed); + + int count_written_right_word = 0; + for (int i = 0; i < pool_size / 64; i++) { + if (p[i * 16] == i * 64) { + count_written_right_word++; + } + } + fmt::println("count_written_right_word:{}, total words:{}", count_written_right_word, total_words); + if (count_time > 2 * 603700) { + random_benchmark_print_counters(fpga_ctl); + } + + gpu_mem_ctl->free(dma_buff); + GPUMemCtl::cleanCtx(); +} \ No newline at end of file diff --git a/tb/utils/random_tool.hpp b/tb/utils/random_tool.hpp new file mode 100644 index 0000000..e56ecf4 --- /dev/null +++ b/tb/utils/random_tool.hpp @@ -0,0 +1,7 @@ +#include + +void cpu_random_h2c(uint8_t pci_bus); +void cpu_random_c2h(uint8_t pci_bus); + +void gpu_random_h2c(uint8_t pci_bus); +void gpu_random_c2h(uint8_t pci_bus); \ No newline at end of file diff --git a/tb/utils/throughput_tool.cpp b/tb/utils/throughput_tool.cpp new file mode 100644 index 0000000..c4db055 --- /dev/null +++ b/tb/utils/throughput_tool.cpp @@ -0,0 +1,364 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "dma_tool.hpp" + +void cpu_throughput_h2c(uint8_t pci_bus) { + fmt::println("=====CPU H2C throughput benchmark start====="); + + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + auto cpu_mem_ctl = CPUMemCtl::getInstance(1UL * 1024 * 1024 * 1024); + + size_t pool_size = 1UL * 1024 * 1024 * 1024; + + cpu_mem_ctl->writeTLB([=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }); + + auto dma_buff = cpu_mem_ctl->alloc(pool_size); + auto p = (uint32_t*)dma_buff; + + // data length per cmd deliver + uint32_t length = 1 * 1024; + uint32_t total_cmds = 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + + uint32_t range = 1 * 1024 * 1024 * 1024; + uint32_t range_words = range / 64; + + // initial dma buffer + // FPGA: 512-bit = 64Byte = 16 * uint32_t + for (int i = 0; i < range_words; i++) { + for (int j = 0; j < 16; j++) { + p[i * 16 + j] = i; + } + } + + auto fpga_vaddr = reinterpret_cast(dma_buff); + fpga_ctl->writeReg(100, (uint32_t)(fpga_vaddr >> 32)); + fpga_ctl->writeReg(101, (uint32_t)(fpga_vaddr)); + fpga_ctl->writeReg(102, length); + + fpga_ctl->writeReg(104, total_cmds); + fpga_ctl->writeReg(105, total_words); + fpga_ctl->writeReg(106, range); + fpga_ctl->writeReg(107, range_words); + + reset_counters(fpga_ctl); + + // start + fpga_ctl->writeReg(103, 0); + fpga_ctl->writeReg(103, 1); + + sleep(1); + + auto cycles = fpga_ctl->readReg(512 + 102); + + fmt::println("Number of errors: {}", fpga_ctl->readReg(512 + 101)); + fmt::println("Cycles: {}", cycles); + + double speed = 1.0 * length * total_cmds / (1.0 * cycles * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + + fmt::println("Total length: {}", total_words * 64); + fmt::println("Speed: {:.2f} GB/s", speed); + + fmt::print("Real words in q: {}", fpga_ctl->readReg(512 + 100)); + fmt::print("\n"); + + cpu_mem_ctl->free(dma_buff); + + throughput_benchmark_print_counters(fpga_ctl); +} + +void cpu_throughput_c2h(uint8_t pci_bus) { + fmt::println("=====CPU C2H throughput benchmark start====="); + + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + auto cpu_mem_ctl = CPUMemCtl::getInstance(1UL * 1024 * 1024 * 1024); + + size_t pool_size = 1L * 1024 * 1024 * 1024; + + cpu_mem_ctl->writeTLB([=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }); + + auto dma_buff = cpu_mem_ctl->alloc(pool_size); + volatile auto p = (uint32_t*)dma_buff; + + memset(p, 0, pool_size); + + uint32_t length = 1 * 1024; + uint32_t total_cmds = 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + + auto fpga_vaddr = reinterpret_cast(dma_buff); + printf("fpga_vaddr: 0x%lx\n", fpga_vaddr); + fpga_ctl->writeReg(200, (uint32_t)(fpga_vaddr >> 32)); + fpga_ctl->writeReg(201, (uint32_t)(fpga_vaddr)); + fpga_ctl->writeReg(202, length); + fpga_ctl->writeReg(204, total_cmds); + fpga_ctl->writeReg(205, total_words); + + reset_counters(fpga_ctl); + + fpga_ctl->writeConfig(0x1408 / 4, 0); + uint32_t tag = fpga_ctl->readConfig(0x140c / 4); + printf("tag: %d\n", tag); + fmt::println("{}", tag & 0x7f); + fpga_ctl->writeReg(206, tag); + + fpga_ctl->writeReg(203, 0); // start + fpga_ctl->writeReg(203, 1); + + sleep(5); + + uint32_t count_cmd = fpga_ctl->readReg(512 + 200); + uint32_t count_word = fpga_ctl->readReg(512 + 201); + uint32_t count_time = fpga_ctl->readReg(512 + 202); + + fmt::println("count cmd: {},right: {}", count_cmd, total_cmds); + fmt::println("count word: {},right: {}", count_word, total_words); + fmt::println("count time: {}", count_time); + + double speed = 1.0 * length * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + fmt::println("Speed: {:.2f} GB/s", speed); + + uint32_t right_count = 0; + uint32_t wrong_count = 0; + + // for (int i = 0; i < 4; i++) { + // printf("Packet [%03d]: ", i); + // for (int j = 0; j < 16; j++) { + // printf("%08x ", p[i * 16 + j]); + // } + // printf("\n"); + // } + + for (int i = 0; i < total_words; i++) { + bool is_right = true; + + for (int j = 0; j < 16; j++) { + if ((uint32_t)p[i * 16 + j] != i) { + is_right = false; + break; + } + } + + if (is_right) { + right_count++; + } else { + wrong_count++; + } + } + fmt::println("right data count: {}, wrong data count: {}\n", right_count, wrong_count); + + cpu_mem_ctl->free(dma_buff); + + throughput_benchmark_print_counters(fpga_ctl); +} + +void gpu_throughput_h2c(uint8_t pci_bus) { + fmt::println("=====GPU H2C throughput benchmark start====="); + + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + auto gpu_mem_ctl = GPUMemCtl::getInstance(0, 1UL * 1024 * 1024 * 1024); + + size_t pool_size = 1UL * 1024 * 1024 * 1024; + + gpu_mem_ctl->writeTLB( + [=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }, + true); + + // get buff (gpu vaddr) + auto dma_buff = gpu_mem_ctl->alloc(pool_size); + // gpu pool's base in gpu vaddr + auto gpu_base = reinterpret_cast(gpu_mem_ctl->getDevPtr()); + auto host_base = static_cast(gpu_mem_ctl->getMapDevPtr()); + volatile auto p = host_base + (reinterpret_cast(dma_buff) - gpu_base) / sizeof(uint32_t); + + // data length per cmd deliver + uint32_t length = 1 * 1024; + uint32_t total_cmds = 1 * 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + + uint32_t range = 1 * 1024 * 1024 * 1024; + uint32_t range_words = range / 64; + + // initial dma buffer + // FPGA: 512-bit = 64Byte = 16 * uint32_t + for (int i = 0; i < range_words; i++) { + for (int j = 0; j < 16; j++) { + p[i * 16 + j] = i; + } + } + + auto fpga_vaddr = reinterpret_cast(dma_buff); + fpga_ctl->writeReg(100, (uint32_t)(fpga_vaddr >> 32)); + fpga_ctl->writeReg(101, (uint32_t)(fpga_vaddr)); + fpga_ctl->writeReg(102, length); + + fpga_ctl->writeReg(104, total_cmds); + fpga_ctl->writeReg(105, total_words); + fpga_ctl->writeReg(106, range); + fpga_ctl->writeReg(107, range_words); + + reset_counters(fpga_ctl); + + // start + fpga_ctl->writeReg(103, 0); + fpga_ctl->writeReg(103, 1); + + sleep(1); + + auto cycles = fpga_ctl->readReg(512 + 102); + + fmt::println("Number of errors: {}", fpga_ctl->readReg(512 + 101)); + fmt::println("Cycles: {}", cycles); + + double speed = 1.0 * length * total_cmds / (1.0 * cycles * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + + fmt::println("Total length: {}", total_words * 64); + fmt::println("Speed: {:.2f} GB/s", speed); + + fmt::print("Real words in q: {}", fpga_ctl->readReg(512 + 100)); + fmt::print("\n"); + + gpu_mem_ctl->free(dma_buff); + + throughput_benchmark_print_counters(fpga_ctl); + GPUMemCtl::cleanCtx(); +} + +void gpu_throughput_c2h(uint8_t pci_bus) { + fmt::println("=====GPU C2H throughput benchmark start====="); + + FPGACtl::explictInit(pci_bus, 4 * 1024 * 1024); + auto fpga_ctl = FPGACtl::getInstance(pci_bus); + auto gpu_mem_ctl = GPUMemCtl::getInstance(0, 1UL * 1024 * 1024 * 1024); + + size_t pool_size = 1L * 1024 * 1024 * 1024; + + gpu_mem_ctl->writeTLB( + [=](uint32_t page_index, uint32_t page_size, uint64_t vaddr, uint64_t paddr) { + fpga_ctl->writeReg(8, (uint32_t)(vaddr)); + fpga_ctl->writeReg(9, (uint32_t)((vaddr) >> 32)); + fpga_ctl->writeReg(10, (uint32_t)(paddr)); + fpga_ctl->writeReg(11, (uint32_t)((paddr) >> 32)); + fpga_ctl->writeReg(12, (page_index == 0)); + fpga_ctl->writeReg(13, 1); + fpga_ctl->writeReg(13, 0); + }, + true); + + auto dma_buff = gpu_mem_ctl->alloc(pool_size); + auto gpu_base = reinterpret_cast(gpu_mem_ctl->getDevPtr()); + auto host_base = static_cast(gpu_mem_ctl->getMapDevPtr()); + volatile auto p = host_base + (reinterpret_cast(dma_buff) - gpu_base) / sizeof(uint32_t); + + memset(p, 0, pool_size); + + uint32_t length = 1 * 1024; + uint32_t total_cmds = 1 * 256 * 1024; + uint32_t total_words = length / 64 * total_cmds; + + auto fpga_vaddr = reinterpret_cast(dma_buff); + fpga_ctl->writeReg(200, (uint32_t)(fpga_vaddr >> 32)); + fpga_ctl->writeReg(201, (uint32_t)(fpga_vaddr)); + fpga_ctl->writeReg(202, length); + fpga_ctl->writeReg(204, total_cmds); + fpga_ctl->writeReg(205, total_words); + + reset_counters(fpga_ctl); + + fpga_ctl->writeConfig(0x1408 / 4, 0); + uint32_t tag = fpga_ctl->readConfig(0x140c / 4); + printf("tag: %d\n", tag); + fmt::println("{}", tag & 0x7f); + fpga_ctl->writeReg(206, tag); + + fpga_ctl->writeReg(203, 0); // start + fpga_ctl->writeReg(203, 1); + + sleep(3); + + uint32_t count_cmd = fpga_ctl->readReg(512 + 200); + uint32_t count_word = fpga_ctl->readReg(512 + 201); + uint32_t count_time = fpga_ctl->readReg(512 + 202); + + fmt::println("count cmd: {},right: {}", count_cmd, total_cmds); + fmt::println("count word: {},right: {}", count_word, total_words); + fmt::println("count time: {}", count_time); + + double speed = 1.0 * length * total_cmds / (1.0 * count_time * 4 / 1000 / 1000 / 1000) / 1024 / 1024 / 1024; + fmt::println("Speed: {:.2f} GB/s", speed); + + uint32_t right_count = 0; + uint32_t wrong_count = 0; + uint32_t value_verify = 0; + + for (int i = 0; i < 16384; i++, value_verify++) { + uint32_t val = p[i * 16]; + if (value_verify != val) { + wrong_count++; + continue; + } + + bool right = true; + for (int j = 1; j < 16; j++) { + if (p[i * 16 + j] != val) { + wrong_count++; + right = false; + break; + } + } + + if (right) { + right_count++; + } + } + fmt::println("right data count: {}, wrong data count: {}\n", right_count, wrong_count); + + gpu_mem_ctl->free(dma_buff); + fpga_ctl->writeReg(203, 0); + + throughput_benchmark_print_counters(fpga_ctl); + GPUMemCtl::cleanCtx(); +} \ No newline at end of file diff --git a/tb/utils/throughput_tool.hpp b/tb/utils/throughput_tool.hpp new file mode 100644 index 0000000..734e3fc --- /dev/null +++ b/tb/utils/throughput_tool.hpp @@ -0,0 +1,7 @@ +#include + +void cpu_throughput_h2c(uint8_t pci_bus); +void cpu_throughput_c2h(uint8_t pci_bus); + +void gpu_throughput_h2c(uint8_t pci_bus); +void gpu_throughput_c2h(uint8_t pci_bus); \ No newline at end of file