CUDA has supported multiple GPUs since the beginning, but originally, each GPU had to be controlled by a separate CPU thread. For workloads that required a lot of CPU power, that requirement was never very onerous because the full power of modern multicore processors can be unlocked only through multithreading.
The multithreaded implementation of multi-GPU N-Body creates one CPU thread per GPU to dispatch and synchronize the work for a given N-body pass. The parent thread splits the work evenly between GPUs, hands each child thread its share, then joins the threads to wait for all of them to finish before proceeding. As the number of GPUs grows, synchronization overhead starts to chip away at the benefits from parallelism.
This implementation uses C++ std::thread, like the
multithreaded CPU implementation of N-body described in Section 14.7
(see also Section A.2): the application thread launches one worker per
GPU and joins them to synchronize on their completion.
Listing 9-5 gives the host code that initializes the GPUs. The global
g_numGPUs holds the GPU count; each GPU is set up by
calling initializeGPU(), which calls
cudaSetDevice() and creates that device’s context. Rather
than dedicating a persistent CPU thread to each GPU, the worker threads
launched for each N-body pass select their device with
cudaSetDevice() when they run.
int g_numCPUCores; int g_numGPUs; struct gpuInit_struct{ int iGPU; cudaError_t status;}; voidinitializeGPU( void *_p ){ cudaError_t status; gpuInit_struct *p = (gpuInit_struct *) _p; cuda(SetDevice( p->iGPU ) ); cuda(SetDeviceFlags( cudaDeviceMapHost ) ); cuda(Free(0) );Error: p->status = status; } // ... below is from main() if ( g_numGPUs ) { // optionally override GPU count from command line chCommandLineGet( &g_numGPUs, "numgpus", argc, argv ); for ( int i = 0; i < g_numGPUs; i++ ) { gpuInit_struct initGPU = {i}; initializeGPU( &initGPU ); if ( cudaSuccess != initGPU.status ) { fprintf( stderr, "Initializing GPU %d failed " " with %d (%s)\n", i, initGPU.status, cudaGetErrorString( initGPU.status ) ); return 1; } } }Listing 9-6 shows the host code that runs on each worker thread: the
gpuDelegation structure encapsulates the work that a given GPU must do,
and gpuWorkerThread() selects that GPU with cudaSetDevice()
before performing its share. The application thread code, shown in
Listing 9-7, creates a gpuDelegation structure for each GPU, launches a
std::thread running gpuWorkerThread() for each, and joins
them to wait until all have finished.
struct gpuDelegation { size_t i; // base offset for this thread to process size_t n; // size of this thread's problem size_t N; // total number of bodies int device; // CUDA device this worker runs on float *hostPosMass; float *hostForce; float softeningSquared; cudaError_t status;}; voidgpuWorkerThread( void *_p ){ cudaError_t status_cudart; gpuDelegation *p = (gpuDelegation *) _p; float *dptrPosMass = 0; float *dptrForce = 0; cuda(SetDevice( p->device ) ); // // Each GPU has its own device pointer to the host pointer. // cuda(Malloc( &dptrPosMass, 4*p->N*sizeof(float) ) ); cuda(Malloc( &dptrForce, 3*p->n*sizeof(float) ) ); cuda(MemcpyAsync( dptrPosMass, p->hostPosMass, 4*p->N*sizeof(float), cudaMemcpyHostToDevice ) ); ComputeNBodyGravitation_multiGPU<<<300,256,256*sizeof(float4)>>>( dptrForce, dptrPosMass, p->softeningSquared, p->i, p->n, p->N ); // NOTE: synchronous memcpy, so no need for further // synchronization with device cuda(Memcpy( p->hostForce+3*p->i, dptrForce, 3*p->n*sizeof(float), cudaMemcpyDeviceToHost ) );Error_cudart: cudaFree( dptrPosMass ); cudaFree( dptrForce ); p->status = status_cudart;}floatComputeGravitation_multiGPU_threaded( float *force, float *posMass, float softeningSquared, size_t N){ std::chrono::steady_clock::time_point start, end; start = std::chrono::steady_clock::now(); { gpuDelegation *pgpu = new gpuDelegation[g_numGPUs]; std::vector<std::thread> threads; size_t bodiesPerGPU = N / g_numGPUs; if ( N % g_numGPUs ) { return 0.0f; } size_t i; for ( i = 0; i < g_numGPUs; i++ ) { pgpu[i].hostPosMass = g_hostAOS_PosMass; pgpu[i].hostForce = g_hostAOS_Force; pgpu[i].softeningSquared = softeningSquared; pgpu[i].i = bodiesPerGPU*i; pgpu[i].n = bodiesPerGPU; pgpu[i].N = N; pgpu[i].device = (int) i; threads.emplace_back( gpuWorkerThread, &pgpu[i] ); } for ( auto &t : threads ) t.join(); delete[] pgpu; } end = std::chrono::steady_clock::now(); return std::chrono::duration<double>(end - start).count() * 1000.0f;}The performance and scaling results of the single-threaded and multithreaded versions of multi-GPU N-body are summarized in Section 14.6.