diff --git a/pi_gpu/Makefile b/pi_gpu/Makefile index 39aa56a40234293d7b3eefcd1851d3e01056b40a..caead4fcb3057616f103d559bbd1edb24523ce65 100644 --- a/pi_gpu/Makefile +++ b/pi_gpu/Makefile @@ -1,2 +1,2 @@ pi_gpu: pi_gpu.cu - nvcc $< -o $@ + nvcc $< -o $@ --use_fast_math diff --git a/pi_gpu/pi_gpu.cu b/pi_gpu/pi_gpu.cu index bc8438761b7ccc99e1ed8aa58bcef170447b1ede..6080b2ef3fbe517c78ce294f088e31918ebfeabc 100644 --- a/pi_gpu/pi_gpu.cu +++ b/pi_gpu/pi_gpu.cu @@ -4,29 +4,30 @@ // based on a series expansion of pi, but ignoring numeric concerns #include <chrono> -#include <cuda_runtime.h> #include <iostream> -#include <mpi.h> -#include "constants.h" -#include "kernels.h" using namespace std; // currently init_kernel assumes n_terms_per_thread is a multiple of 10 -uint64_t const n_terms_per_thread = 100000; -uint64_t const n_threads_per_gpu = 1024 * 2; +uint64_t const n_terms_per_thread = 1000000; +uint64_t const n_threads_per_gpu = 1024 * 1024; uint64_t const n_terms_per_gpu = n_terms_per_thread * n_threads_per_gpu; -uint64_t const n_threads_per_block = 512; +uint64_t const n_threads_per_block = 1024; uint64_t const n_blocks_per_gpu = (n_threads_per_gpu + n_threads_per_block - 1) / n_threads_per_block; -int const n_loops = 1; +int const n_loops = 8; //-------------------------------------------------------------------------------------------------- __global__ -void init_kernel(double *arr, int gpu_idx) { +void init_kernel(double *arr) { uint64_t const thread_idx = blockIdx.x * blockDim.x + threadIdx.x; - uint64_t const start = n_terms_per_gpu * gpu_idx + n_terms_per_thread * thread_idx + 1; - uint64_t const end = n_terms_per_gpu * (gpu_idx + 1) + n_terms_per_thread * thread_idx + 1; + if (thread_idx >= n_threads_per_gpu) { + return; + } + + uint64_t const start = n_terms_per_thread * thread_idx + 1; + uint64_t const end = start + n_terms_per_thread; + double sum = 0.0; uint64_t i = start; while (i < end) { @@ -49,8 +50,8 @@ void reduce_sum_kernel(double *arr, uint64_t stride) { } //-------------------------------------------------------------------------------------------------- -void init(double *arr, int gpu_idx) { - init_kernel<<<n_blocks_per_gpu, n_threads_per_block>>>(arr, gpu_idx); +void init(double *arr) { + init_kernel<<<n_blocks_per_gpu, n_threads_per_block>>>(arr); } //-------------------------------------------------------------------------------------------------- @@ -72,15 +73,16 @@ int main(int argc, char** argv) { double result; // timing data + decltype(std::chrono::high_resolution_clock::now()) global_start; + decltype(std::chrono::high_resolution_clock::now()) global_stop; decltype(std::chrono::high_resolution_clock::now()) start; decltype(std::chrono::high_resolution_clock::now()) stop; + global_start = std::chrono::high_resolution_clock::now(); for (int i = 0; i < n_loops; ++i) { - if (rank == 0) { - start = std::chrono::high_resolution_clock::now(); - } + start = std::chrono::high_resolution_clock::now(); - init(d_arr, gpu_idx); + init(d_arr); reduce(d_arr); cudaDeviceSynchronize(); cudaMemcpy(&result, d_arr, sizeof(double), cudaMemcpyDeviceToHost); @@ -88,16 +90,18 @@ int main(int argc, char** argv) { stop = std::chrono::high_resolution_clock::now(); auto const duration = std::chrono::duration_cast<std::chrono::milliseconds>(stop - start); auto const millis = duration.count(); - auto const n_terms_total = n_tasks * n_terms_per_gpu; - auto const gflops = n_terms_total * 5.0 / (millis * 1e-3) * 1e-9; + auto const gflops = n_terms_per_gpu * 5.0 / (millis * 1e-3) * 1e-9; std::cout << "loop " << i << '\n'; - std::cout << "processes = " << n_tasks << ", terms per GPU = " << n_terms_per_gpu - << ", total terms = " << n_terms_total << '\n'; + std::cout << "processes = " << 1 << ", terms per GPU = " << n_terms_per_gpu << '\n'; std::cout << "time = " << millis * 1e-3 << "s, estimated GFlops = " << gflops << '\n'; std::cout << "pi ~ " << result << '\n'; std::cout << '\n'; } + global_stop = std::chrono::high_resolution_clock::now(); + auto const duration = std::chrono::duration_cast<std::chrono::milliseconds>(global_stop - global_start); + auto const millis = duration.count(); + std::cout << "total time = " << millis * 1e-3 << "s\n"; cudaFree(d_arr); return 0; diff --git a/pi_gpu/pi_gpu.slurm b/pi_gpu/pi_gpu.slurm index 492e0d4406c6e860011ab7d58a1fbf8cfcfee162..07c78d0f822c1c4ab17e2c0b50e153771663fcd4 100644 --- a/pi_gpu/pi_gpu.slurm +++ b/pi_gpu/pi_gpu.slurm @@ -8,8 +8,8 @@ #SBATCH --cpus-per-task=1 #SBATCH --ntasks-per-core=1 #SBATCH --threads-per-core=1 -#SBATCH --mem=100M -#SBATCH --time 00:01:00 +#SBATCH --mem=1G +#SBATCH --time 00:05:00 source ./load_modules.sh srun ./pi_gpu