I'm applying UVA and OpenMP in my algorithm to make it powerful.
The thing is that when I launch a parallel kernel, that is for example, 3 CPU threads launch one kernel at the same time. One thread has nan values.
It seems that GPU X cannot read a variable from GPU0.
That is weird taking into account that I grant access to every GPU to 0 (In this case 1 and 2).
Is there a problem to use UVA and OpenMP together? Or is a problem of the code?
Here is the code and the results.
I've created a MCVE to demonstrate the error here:
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <math.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "math_constants.h"
#include <omp.h>
#include <cufft.h>
inline bool IsGPUCapableP2P(cudaDeviceProp *pProp)
{
#ifdef _WIN32
return (bool)(pProp->tccDriver ? true : false);
#else
return (bool)(pProp->major >= 2);
#endif
}
inline bool IsAppBuiltAs64()
{
#if defined(__x86_64) || defined(AMD64) || defined(_M_AMD64)
return 1;
#else
return 0;
#endif
}
__global__ void kernelFunction(cufftComplex *I, int i, int N)
{
int j = threadIdx.x + blockDim.x * blockIdx.x;
int k = threadIdx.y + blockDim.y * blockIdx.y;
if(j==0 & k==0){
printf("I'm thread %d and I'm reading device_I[0] = %f\n", i, I[N*j+k].x);
}
}
__host__ int main(int argc, char **argv) {
int num_gpus;
cudaGetDeviceCount(&num_gpus);
if(num_gpus < 1){
printf("No CUDA capable devices were detected\n");
return 1;
}
if (!IsAppBuiltAs64()){
printf("%s is only supported with on 64-bit OSs and the application must be built as a 64-bit target. Test is being waived.\n", argv[0]);
exit(EXIT_SUCCESS);
}
printf("Number of host CPUs:\t%d\n", omp_get_num_procs());
printf("Number of CUDA devices:\t%d\n", num_gpus);
for(int i = 0; i < num_gpus; i++){
cudaDeviceProp dprop;
cudaGetDeviceProperties(&dprop, i);
printf("> GPU%d = \"%15s\" %s capable of Peer-to-Peer (P2P)\n", i, dprop.name, (IsGPUCapableP2P(&dprop) ? "IS " : "NOT"));
//printf(" %d: %s\n", i, dprop.name);
}
printf("---------------------------\n");
num_gpus = 3; //The case that fails
omp_set_num_threads(num_gpus);
if(num_gpus > 1){
for(int i=1; i<num_gpus; i++){
cudaDeviceProp dprop0, dpropX;
cudaGetDeviceProperties(&dprop0, 0);
cudaGetDeviceProperties(&dpropX, i);
int canAccessPeer0_x, canAccessPeerx_0;
cudaDeviceCanAccessPeer(&canAccessPeer0_x, 0, i);
cudaDeviceCanAccessPeer(&canAccessPeerx_0 , i, 0);
printf("> Peer-to-Peer (P2P) access from %s (GPU%d) -> %s (GPU%d) : %s\n", dprop0.name, 0, dpropX.name, i, canAccessPeer0_x ? "Yes" : "No");
printf("> Peer-to-Peer (P2P) access from %s (GPU%d) -> %s (GPU%d) : %s\n", dpropX.name, i, dprop0.name, 0, canAccessPeerx_0 ? "Yes" : "No");
if(canAccessPeer0_x == 0 || canAccessPeerx_0 == 0){
printf("Two or more SM 2.0 class GPUs are required for %s to run.\n", argv[0]);
printf("Support for UVA requires a GPU with SM 2.0 capabilities.\n");
printf("Peer to Peer access is not available between GPU%d <-> GPU%d, waiving test.\n", 0, i);
exit(EXIT_SUCCESS);
}else{
cudaSetDevice(0);
printf("Granting access from 0 to %d...\n", i);
cudaDeviceEnablePeerAccess(i,0);
cudaSetDevice(i);
printf("Granting access from %d to 0...\n", i);
cudaDeviceEnablePeerAccess(0,0);
printf("Checking GPU%d and GPU%d for UVA capabilities...\n", 0, 1);
const bool has_uva = (dprop0.unifiedAddressing && dpropX.unifiedAddressing);
printf("> %s (GPU%d) supports UVA: %s\n", dprop0.name, 0, (dprop0.unifiedAddressing ? "Yes" : "No"));
printf("> %s (GPU%d) supports UVA: %s\n", dpropX.name, i, (dpropX.unifiedAddressing ? "Yes" : "No"));
if (has_uva){
printf("Both GPUs can support UVA, enabling...\n");
}
else{
printf("At least one of the two GPUs does NOT support UVA, waiving test.\n");
exit(EXIT_SUCCESS);
}
}
}
}
int M = 512;
int N = 512;
cufftComplex *host_I = (cufftComplex*)malloc(M*N*sizeof(cufftComplex));
for(int i=0;i<M;i++){
for(int j=0;j<N;j++){
host_I[N*i+j].x = 0.001;
host_I[N*i+j].y = 0;
}
}
cufftComplex *device_I;
cudaSetDevice(0);
cudaMalloc((void**)&device_I, sizeof(cufftComplex)*M*N);
cudaMemset(device_I, 0, sizeof(cufftComplex)*M*N);
cudaMemcpy2D(device_I, sizeof(cufftComplex), host_I, sizeof(cufftComplex), sizeof(cufftComplex), M*N, cudaMemcpyHostToDevice);
dim3 threads(32,32);
dim3 blocks(M/threads.x, N/threads.y);
dim3 threadsPerBlockNN = threads;
dim3 numBlocksNN = blocks;
#pragma omp parallel
{
unsigned int i = omp_get_thread_num();
unsigned int num_cpu_threads = omp_get_num_threads();
// set and check the CUDA device for this CPU thread
int gpu_id = -1;
cudaSetDevice(i % num_gpus); // "% num_gpus" allows more CPU threads than GPU devices
cudaGetDevice(&gpu_id);
//printf("CPU thread %d (of %d) uses CUDA device %d\n", cpu_thread_id, num_cpu_threads, gpu_id);
kernelFunction<<<numBlocksNN, threadsPerBlockNN>>>(device_I, i, N);
cudaDeviceSynchronize();
}
cudaFree(device_I);
for(int i=1; i<num_gpus; i++){
cudaSetDevice(0);
cudaDeviceDisablePeerAccess(i);
cudaSetDevice(i);
cudaDeviceDisablePeerAccess(0);
}
for(int i=0; i<num_gpus; i++ ){
cudaSetDevice(i);
cudaDeviceReset();
}
free(host_I);
}
The results are:
Both GPUs can support UVA, enabling...
I'm thread 0 and I'm reading device_I[0] = 0.001000
I'm thread 2 and I'm reading device_I[0] = 0.001000
I'm thread 1 and I'm reading device_I[0] = -nan
The command line to compile is:
nvcc -Xcompiler -fopenmp -lgomp -arch=sm_37 main.cu -lcufft
Here is the result of simpleP2P:
[miguel.carcamo@belka simpleP2P]$ ./simpleP2P
[./simpleP2P] - Starting...
Checking for multiple GPUs...
CUDA-capable device count: 8
> GPU0 = " Tesla K80" IS capable of Peer-to-Peer (P2P)
> GPU1 = " Tesla K80" IS capable of Peer-to-Peer (P2P)
> GPU2 = " Tesla K80" IS capable of Peer-to-Peer (P2P)
> GPU3 = " Tesla K80" IS capable of Peer-to-Peer (P2P)
> GPU4 = " Tesla K80" IS capable of Peer-to-Peer (P2P)
> GPU5 = " Tesla K80" IS capable of Peer-to-Peer (P2P)
> GPU6 = " Tesla K80" IS capable of Peer-to-Peer (P2P)
> GPU7 = " Tesla K80" IS capable of Peer-to-Peer (P2P)
Checking GPU(s) for support of peer to peer memory access...
> Peer-to-Peer (P2P) access from Tesla K80 (GPU0) -> Tesla K80 (GPU1) : Yes
> Peer-to-Peer (P2P) access from Tesla K80 (GPU1) -> Tesla K80 (GPU0) : Yes
Enabling peer access between GPU0 and GPU1...
Checking GPU0 and GPU1 for UVA capabilities...
> Tesla K80 (GPU0) supports UVA: Yes
> Tesla K80 (GPU1) supports UVA: Yes
Both GPUs can support UVA, enabling...
Allocating buffers (64MB on GPU0, GPU1 and CPU Host)...
Creating event handles...
cudaMemcpyPeer / cudaMemcpy between GPU0 and GPU1: 0.79GB/s
Preparing host buffer and memcpy to GPU0...
Run kernel on GPU1, taking source data from GPU0 and writing to GPU1...
Run kernel on GPU0, taking source data from GPU1 and writing to GPU0...
Copy data back to host from GPU0 and verify results...
Verification error @ element 0: val = nan, ref = 0.000000
Verification error @ element 1: val = nan, ref = 4.000000
Verification error @ element 2: val = nan, ref = 8.000000
Verification error @ element 3: val = nan, ref = 12.000000
Verification error @ element 4: val = nan, ref = 16.000000
Verification error @ element 5: val = nan, ref = 20.000000
Verification error @ element 6: val = nan, ref = 24.000000
Verification error @ element 7: val = nan, ref = 28.000000
Verification error @ element 8: val = nan, ref = 32.000000
Verification error @ element 9: val = nan, ref = 36.000000
Verification error @ element 10: val = nan, ref = 40.000000
Verification error @ element 11: val = nan, ref = 44.000000
Enabling peer access...
Shutting down...
Test failed!
