Skip to content

Commit 5aefaba

Browse files
quantumstevequantumsteve
committed
hipify kernel
Signed-off-by: Steven Hahn <hahnse@ornl.gov>
1 parent 0ff4364 commit 5aefaba

4 files changed

Lines changed: 155 additions & 3 deletions

File tree

‎bin/hip/CMakeLists.txt‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
cmake_minimum_required(VERSION 3.21) # HIP language support requires 3.21
22
cmake_policy(VERSION 3.21.3...3.27)
33
project(MyProj LANGUAGES HIP)
4-
add_executable(gpu_benchmark gpu_benchmark.hip kernels.hip)
4+
add_executable(gpu_benchmark gpu_benchmark.hip kernels.hip zero_one.hip)
55

‎bin/hip/gpu_benchmark.hip‎

Lines changed: 87 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
11
#include "kernels.h"
2+
#include "zero_one.h"
23
#include <hipcub/hipcub.hpp>
34

45
#include <chrono>
@@ -100,6 +101,90 @@ void runBenchmark(long max_work) {
100101
runBenchmarkTime(max_work, std::nullopt);
101102
}
102103

104+
// Function to run the GPU benchmark for a specified time
105+
void memoryBenchmarkTime(long max_work, std::optional<int> runtime_in_seconds) {
106+
uint32_t n = 512;
107+
uint64_t m = (max_work + n - 1) / n;
108+
109+
// allocate memory
110+
unsigned long long int *d_tmp;
111+
HIP_CHECK(hipMalloc((void **)&d_tmp, max_work * sizeof(unsigned long long int)));
112+
113+
// set up timing stuff
114+
hipEvent_t gpu_start, gpu_stop;
115+
HIP_CHECK(hipEventCreate(&gpu_start));
116+
HIP_CHECK(hipEventCreate(&gpu_stop));
117+
118+
// set kernel
119+
dim3 gridSize = m;
120+
dim3 blockSize = n;
121+
122+
std::cout << n << " " << m << std::endl;
123+
124+
// allocate memory
125+
unsigned long long int *d_count;
126+
HIP_CHECK(hipMalloc((void **)&d_count, m * sizeof(unsigned long long int)));
127+
HIP_CHECK(hipMemset(d_count, 0, m * sizeof(unsigned long long int)));
128+
129+
// monte carlo kernel
130+
HIP_CHECK(hipEventRecord(gpu_start, 0));
131+
132+
int iteration = 0;
133+
if (!runtime_in_seconds.has_value()) {
134+
HIP_CHECK(hipMemset(d_tmp, 0, max_work * sizeof(unsigned long long int)));
135+
stress_vm_zero<<<gridSize, blockSize>>>(d_tmp, d_count, max_work);
136+
HIP_CHECK(hipMemset(d_tmp, 0xff, max_work * sizeof(unsigned long long int)));
137+
stress_vm_one<<<gridSize, blockSize>>>(d_tmp, d_count, max_work);
138+
HIP_CHECK(hipDeviceSynchronize());
139+
iteration++;
140+
} else {
141+
// Run the workload loop until the specified runtime is reached
142+
while (getElapsedTime(gpu_start, gpu_stop) < runtime_in_seconds) {
143+
HIP_CHECK(hipMemset(d_tmp, 0, max_work * sizeof(unsigned long long int)));
144+
stress_vm_zero<<<gridSize, blockSize>>>(d_tmp, d_count, max_work);
145+
HIP_CHECK(hipMemset(d_tmp, 0xff, max_work * sizeof(unsigned long long int)));
146+
stress_vm_one<<<gridSize, blockSize>>>(d_tmp, d_count, max_work);
147+
HIP_CHECK(hipDeviceSynchronize()); // Ensure the kernel has finished executing
148+
iteration++;
149+
}
150+
}
151+
152+
float gpu_elapsed_time = getElapsedTime(gpu_start, gpu_stop);
153+
HIP_CHECK(hipEventDestroy(gpu_start));
154+
HIP_CHECK(hipEventDestroy(gpu_stop));
155+
156+
// copy results back to the host
157+
// Allocate device output array
158+
unsigned long long int *d_out = nullptr;
159+
HIP_CHECK(hipMalloc((void **)&d_out, sizeof(unsigned long long int)));
160+
161+
// Request and allocate temporary storage
162+
void *d_temp_storage = nullptr;
163+
size_t temp_storage_bytes = 0;
164+
HIP_CHECK(hipcub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_count, d_out, m));
165+
HIP_CHECK(hipMalloc((void **)&d_temp_storage, temp_storage_bytes));
166+
167+
// Run
168+
HIP_CHECK(hipcub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_count, d_out, m));
169+
170+
// copy results back to the host
171+
unsigned long long int h_count = 0;
172+
HIP_CHECK(hipMemcpy(&h_count, d_count, sizeof(unsigned long long int), hipMemcpyDeviceToHost));
173+
174+
// display results and timings for gpu
175+
float pi = (float)h_count;
176+
std::cout << "Approximate pi calculated on GPU is: " << pi << " and calculation took " << gpu_elapsed_time << "s\n";
177+
178+
HIP_CHECK(hipFree(d_tmp));
179+
HIP_CHECK(hipFree(d_count));
180+
HIP_CHECK(hipFree(d_temp_storage));
181+
}
182+
183+
void memoryBenchmark(long max_work) {
184+
memoryBenchmarkTime(max_work, std::nullopt);
185+
}
186+
187+
103188
int main(int argc, char *argv[]) {
104189
// Check for the correct number of command line arguments
105190
if (argc == 2) {
@@ -112,7 +197,7 @@ int main(int argc, char *argv[]) {
112197
return 1;
113198
}
114199

115-
runBenchmark(max_work);
200+
memoryBenchmark(max_work);
116201

117202
} else if (argc == 3) {
118203
// Parse the command line arguments
@@ -125,7 +210,7 @@ int main(int argc, char *argv[]) {
125210
return 1;
126211
}
127212

128-
runBenchmarkTime(max_work, runtime_in_seconds);
213+
memoryBenchmarkTime(max_work, runtime_in_seconds);
129214

130215
} else {
131216
std::cerr << "Usage: " << argv[0] << " <max_work> [runtime_in_seconds]" << std::endl;

‎bin/hip/zero_one.h‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,11 @@
1+
#include "hip/hip_runtime.h"
2+
#ifndef __ZERO_ONE_CUH__
3+
#define __ZERO_ONE_CUH__
4+
5+
__global__ void stress_vm_zero(unsigned long long int *buf,
6+
unsigned long long int *count,
7+
const size_t sz);
8+
__global__ void stress_vm_one(unsigned long long int *buf,
9+
unsigned long long int *count,
10+
const size_t sz);
11+
#endif

‎bin/hip/zero_one.hip‎

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
1+
#include "hip/hip_runtime.h"
2+
#include "zero_one.h"
3+
4+
#include <hipcub/hipcub.hpp>
5+
6+
using uint_type = unsigned long long int;
7+
8+
/*
9+
* stress_vm_zero()
10+
* set all memory to zero and see if any bits are stuck at one and
11+
*/
12+
__global__ void stress_vm_zero(
13+
uint_type *buf,
14+
uint_type *count,
15+
const size_t sz)
16+
{
17+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
18+
if (tid >= sz)
19+
return;
20+
21+
uint_type bit_errors = __popcll(buf[tid]);
22+
23+
typedef hipcub::BlockReduce<unsigned long long int, 512> BlockReduceT;
24+
__shared__ typename BlockReduceT::TempStorage temp_storage;
25+
uint_type aggregate = BlockReduceT(temp_storage).Sum(bit_errors);
26+
27+
// update to our global variable count
28+
if (threadIdx.x == 0) {
29+
count[blockIdx.x] += aggregate;
30+
}
31+
}
32+
33+
/*
34+
* stress_vm_one()
35+
* set all memory to one and see if any bits are stuck at zero
36+
*/
37+
__global__ void stress_vm_one(
38+
uint_type *buf,
39+
uint_type *count,
40+
const size_t sz)
41+
{
42+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
43+
if (tid >= sz)
44+
return;
45+
46+
uint_type bit_errors = __popcll(~buf[tid]);
47+
48+
typedef hipcub::BlockReduce<unsigned long long int, 512> BlockReduceT;
49+
__shared__ typename BlockReduceT::TempStorage temp_storage;
50+
uint_type aggregate = BlockReduceT(temp_storage).Sum(bit_errors);
51+
52+
// update to our global variable count
53+
if (threadIdx.x == 0) {
54+
count[blockIdx.x] += aggregate;
55+
}
56+
}

0 commit comments

Comments
 (0)