Skip to content

<feat>: add gpu benchmark #1

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Jac0912 wants to merge 1 commit into
base: main
Choose a base branch
from
Open

<feat>: add gpu benchmark #1

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
132 changes: 132 additions & 0 deletions README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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.
97 changes: 97 additions & 0 deletions reset
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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

55 changes: 55 additions & 0 deletions tb/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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)
Binary file added tb/empty
View file
Open in desktop
Binary file not shown.
72 changes: 72 additions & 0 deletions tb/empty.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,72 @@
#include <csignal>
#include <cstdio>

#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<unsigned int>(tid + 1);
for (int j = 0; j < 1000000; j++) {
acc = acc * 1664525u + 1013904223u + static_cast<unsigned int>(a + b + idx + j);
}

d_c[tid] = static_cast<int>(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<<<kBlocks, kThreads>>>(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;
}
Loading