Skip to content

Commit 10a0384

Browse files
henricasanovahenricasanova
committed
Merge branch 'stress-ng_cpu_benchmark'
2 parents 0366f00 + a58e402 commit 10a0384

30 files changed

Lines changed: 1045 additions & 850 deletions

‎Makefile‎

Lines changed: 0 additions & 40 deletions
This file was deleted.

‎bin/cpu-benchmark.cpp‎

Lines changed: 0 additions & 89 deletions
This file was deleted.

‎bin/cuda/.clang-format‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,3 @@
1+
BasedOnStyle: LLVM
2+
IndentWidth: 2
3+
ColumnLimit: 120

‎bin/cuda/CMakeLists.txt‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
cmake_minimum_required(VERSION 3.21) # HIP language support requires 3.21
2+
cmake_policy(VERSION 3.21.3...3.27)
3+
project(MyProj LANGUAGES CUDA)
4+
add_executable(gpu_benchmark gpu_benchmark.cu kernels.cu)
5+

‎bin/cuda/gpu_benchmark.cu‎

Lines changed: 185 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,185 @@
1+
#include "gpu_benchmark.h"
2+
3+
#include "kernels.h"
4+
#include <cub/cub.cuh>
5+
6+
#include <chrono>
7+
#include <iostream>
8+
9+
// The macro wraps any CUDA API call
10+
#define CUDA_CHECK(ans) \
11+
{ gpuAssert((ans), __FILE__, __LINE__); }
12+
13+
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) {
14+
if (code != cudaSuccess) {
15+
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
16+
if (abort)
17+
exit(code);
18+
}
19+
}
20+
21+
float getElapsedTime(const cudaEvent_t &gpu_start, cudaEvent_t &gpu_stop) {
22+
float gpu_elapsed_time;
23+
CUDA_CHECK(cudaEventRecord(gpu_stop, 0));
24+
CUDA_CHECK(cudaEventSynchronize(gpu_stop));
25+
CUDA_CHECK(cudaEventElapsedTime(&gpu_elapsed_time, gpu_start, gpu_stop));
26+
return gpu_elapsed_time / 1000.0f;
27+
}
28+
29+
// Function to run the GPU benchmark with no time limit
30+
void runBenchmark(long max_work) {
31+
uint32_t n = 256 * 256;
32+
uint64_t m = max_work * 16384 / n;
33+
34+
unsigned long long int *d_count;
35+
curandState *d_state;
36+
CUDA_CHECK(cudaMalloc((void **)&d_count, 256 * sizeof(unsigned long long int)));
37+
CUDA_CHECK(cudaMalloc((void **)&d_state, n * sizeof(curandState)));
38+
CUDA_CHECK(cudaMemset(d_count, 0, 256 * sizeof(unsigned long long int)));
39+
40+
// set up timing stuff
41+
cudaEvent_t gpu_start, gpu_stop;
42+
CUDA_CHECK(cudaEventCreate(&gpu_start));
43+
CUDA_CHECK(cudaEventCreate(&gpu_stop));
44+
45+
// set kernel
46+
dim3 gridSize = 256;
47+
dim3 blockSize = 256;
48+
setup_kernel<<<gridSize, blockSize>>>(d_state);
49+
50+
// monte carlo kernel
51+
CUDA_CHECK(cudaEventRecord(gpu_start, 0));
52+
monte_carlo_kernel<<<gridSize, blockSize>>>(d_state, d_count, m);
53+
CUDA_CHECK(cudaDeviceSynchronize());
54+
55+
float gpu_elapsed_time = getElapsedTime(gpu_start, gpu_stop);
56+
CUDA_CHECK(cudaEventDestroy(gpu_start));
57+
CUDA_CHECK(cudaEventDestroy(gpu_stop));
58+
59+
// Allocate device output array
60+
unsigned long long int *d_out = nullptr;
61+
CUDA_CHECK(cudaMalloc((void **)&d_out, sizeof(unsigned long long int)));
62+
63+
// Request and allocate temporary storage
64+
void *d_temp_storage = nullptr;
65+
size_t temp_storage_bytes = 0;
66+
CUDA_CHECK(cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_count, d_out, 256));
67+
CUDA_CHECK(cudaMalloc((void **)&d_temp_storage, temp_storage_bytes));
68+
69+
// Run
70+
CUDA_CHECK(cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_count, d_out, 256));
71+
72+
// copy results back to the host
73+
unsigned long long int h_count = 0;
74+
CUDA_CHECK(cudaMemcpy(&h_count, d_out, sizeof(unsigned long long int), cudaMemcpyDeviceToHost));
75+
76+
// display results and timings for gpu
77+
float pi = h_count * 4.0 / (n * m);
78+
std::cout << "Approximate pi calculated on GPU is: " << pi << " and calculation took " << gpu_elapsed_time << "s\n";
79+
std::cout << "Benchmark completed!" << std::endl;
80+
81+
CUDA_CHECK(cudaFree(d_count));
82+
CUDA_CHECK(cudaFree(d_state));
83+
CUDA_CHECK(cudaFree(d_out));
84+
CUDA_CHECK(cudaFree(d_temp_storage));
85+
}
86+
87+
// Function to run the GPU benchmark for a specified time
88+
void runBenchmarkTime(long max_work, int runtime_in_seconds) {
89+
90+
uint32_t n = 256 * 256;
91+
uint64_t m = max_work * 16384 / n;
92+
93+
// allocate memory
94+
unsigned long long int *d_count;
95+
curandState *d_state;
96+
CUDA_CHECK(cudaMalloc((void **)&d_count, 256 * sizeof(unsigned long long int)));
97+
CUDA_CHECK(cudaMalloc((void **)&d_state, n * sizeof(curandState)));
98+
CUDA_CHECK(cudaMemset(d_count, 0, 256 * sizeof(unsigned long long int)));
99+
100+
// set up timing stuff
101+
cudaEvent_t gpu_start, gpu_stop;
102+
CUDA_CHECK(cudaEventCreate(&gpu_start));
103+
CUDA_CHECK(cudaEventCreate(&gpu_stop));
104+
105+
// set kernel
106+
dim3 gridSize = 256;
107+
dim3 blockSize = 256;
108+
109+
setup_kernel<<<gridSize, blockSize>>>(d_state);
110+
111+
CUDA_CHECK(cudaEventRecord(gpu_start, 0));
112+
int iteration = 0;
113+
// Run the workload loop until the specified runtime is reached
114+
while (getElapsedTime(gpu_start, gpu_stop) < runtime_in_seconds) {
115+
monte_carlo_kernel<<<gridSize, blockSize>>>(d_state, d_count, m);
116+
CUDA_CHECK(cudaDeviceSynchronize()); // Ensure the kernel has finished executing
117+
iteration++;
118+
}
119+
120+
float gpu_elapsed_time = getElapsedTime(gpu_start, gpu_stop);
121+
CUDA_CHECK(cudaEventDestroy(gpu_start));
122+
CUDA_CHECK(cudaEventDestroy(gpu_stop));
123+
124+
// copy results back to the host
125+
// Allocate device output array
126+
unsigned long long int *d_out = nullptr;
127+
CUDA_CHECK(cudaMalloc((void **)&d_out, sizeof(unsigned long long int)));
128+
129+
// Request and allocate temporary storage
130+
void *d_temp_storage = nullptr;
131+
size_t temp_storage_bytes = 0;
132+
CUDA_CHECK(cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_count, d_out, 256));
133+
CUDA_CHECK(cudaMalloc((void **)&d_temp_storage, temp_storage_bytes));
134+
135+
// Run
136+
CUDA_CHECK(cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_count, d_out, 256));
137+
138+
// copy results back to the host
139+
unsigned long long int h_count = 0;
140+
CUDA_CHECK(cudaMemcpy(&h_count, d_out, sizeof(unsigned long long int), cudaMemcpyDeviceToHost));
141+
142+
// display results and timings for gpu
143+
float pi = h_count * 4.0 / (n * m) / iteration;
144+
std::cout << "Approximate pi calculated on GPU is: " << pi << " and calculation took " << gpu_elapsed_time << "s\n";
145+
146+
CUDA_CHECK(cudaFree(d_count));
147+
CUDA_CHECK(cudaFree(d_state));
148+
CUDA_CHECK(cudaFree(d_out));
149+
CUDA_CHECK(cudaFree(d_temp_storage));
150+
}
151+
152+
int main(int argc, char *argv[]) {
153+
// Check for the correct number of command line arguments
154+
if (argc == 2) {
155+
// Parse the command line arguments
156+
long max_work = std::atol(argv[1]);
157+
158+
// Validate the input arguments
159+
if (max_work <= 0) {
160+
std::cerr << "max_work must be a positive integer." << std::endl;
161+
return 1;
162+
}
163+
164+
runBenchmark(max_work);
165+
166+
} else if (argc == 3) {
167+
// Parse the command line arguments
168+
long max_work = std::atol(argv[1]);
169+
int runtime_in_seconds = std::atoi(argv[2]);
170+
171+
// Validate the input arguments
172+
if (max_work <= 0 || runtime_in_seconds <= 0) {
173+
std::cerr << "Both max_work and runtime_in_seconds must be positive integers." << std::endl;
174+
return 1;
175+
}
176+
177+
runBenchmarkTime(max_work, runtime_in_seconds);
178+
179+
} else {
180+
std::cerr << "Usage: " << argv[0] << " <max_work> [runtime_in_seconds]" << std::endl;
181+
return 1;
182+
}
183+
184+
return 0;
185+
}

bin/gpu_benchmark.h renamed to bin/cuda/gpu_benchmark.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2,10 +2,9 @@
22
#define GPU_BENCHMARK_H
33

44
#include <cuda_runtime.h>
5+
#include <curand_kernel.h>
56

67
void runBenchmark(int max_work);
78
void runBenchmarkTime(int max_work, int runtime_in_seconds);
89

910
#endif // GPU_BENCHMARK_H
10-
11-

‎bin/cuda/kernels.cu‎

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
#include "kernels.h"
2+
3+
#include <cub/cub.cuh>
4+
5+
__global__ void setup_kernel(curandState *state) {
6+
int index = threadIdx.x + blockDim.x * blockIdx.x;
7+
curand_init(123456789, index, 0, &state[index]);
8+
}
9+
10+
__global__ void monte_carlo_kernel(curandState *state, unsigned long long int *count, int64_t m) {
11+
unsigned int index = threadIdx.x + blockDim.x * blockIdx.x;
12+
13+
unsigned long long int thread_data = 0;
14+
15+
unsigned int temp = 0;
16+
while (temp < m) {
17+
float x = curand_uniform(&state[index]);
18+
float y = curand_uniform(&state[index]);
19+
float r = x * x + y * y;
20+
21+
if (r <= 1) {
22+
thread_data++;
23+
}
24+
temp++;
25+
}
26+
27+
typedef cub::BlockReduce<unsigned long long int, 256> BlockReduceT;
28+
__shared__ typename BlockReduceT::TempStorage temp_storage;
29+
unsigned long long int aggregate = BlockReduceT(temp_storage).Sum(thread_data);
30+
31+
// update to our global variable count
32+
if (threadIdx.x == 0) {
33+
count[blockIdx.x] += aggregate;
34+
}
35+
}

bin/kernels.cuh renamed to bin/cuda/kernels.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,6 @@
44
#include <curand_kernel.h>
55

66
__global__ void setup_kernel(curandState *state);
7-
__global__ void monte_carlo_kernel(curandState *state, int *count, int m);
7+
__global__ void monte_carlo_kernel(curandState *state, unsigned long long int *count, int64_t m);
88

99
#endif
10-

0 commit comments

Comments
 (0)