cudaErrorIllegalAdress on cudaMemcpy

Question

私はcudaの新人で、球面上にランダムな点を生成する少しのコードを書こうとしています。ここにコードがあります。

__global__ 
    void setup_kernel(curandStateMRG32k3a *state) 
    { 
     int id = threadIdx.x + blockIdx.x * blockDim.x; 
     curand_init(0, id, 0, &state[id]); 
    } 

    __global__ 
    void computeRandomVectors(float* x, float* y, float* z, unsigned int numberOfElements,curandStateMRG32k3a *state) 
    { 
     float a,b; 
     unsigned int i = blockDim.x * blockIdx.x + threadIdx.x; 
     curandStateMRG32k3a localState = state[i]; 
     if(i < numberOfElements) 
     { 
      a = curand_uniform(&localState); 
      b = curand_uniform(&localState); 
      while(a * a + b * b > 1.0f) 
     { 
      a = curand_uniform(&localState) * 2.0f - 1.0f; 
      b = curand_uniform(&localState) * 2.0f - 1.0f; 
     } 
      x[i] = 2.0f * a * sqrtf(1.0f - a * a - b * b); 
      y[i] = 2.0f * b * sqrtf(1.0f - a * a - b * b); 
      z[i] = 1.0f - 2.0f * (a * a + b * b); 
     } 
    } 

    void generatePointsOnASphere(thrust::host_vector<float>& h_x,  thrust::host_vector<float>& h_y, thrust::host_vector<float>& h_z) 
    { 
     if(h_x.size() != h_y.size() && h_x.size() != h_z.size()) 
     { 
      std::cout << "The three component vectors have unmatching size()" << std::endl; 
      return; 
     } 

     size_t size = h_x.size() * sizeof(float); 

     float* h_p_x = (float*) calloc(h_x.size(),sizeof(float)); 
     float* h_p_y = (float*) calloc(h_x.size(),sizeof(float)); 
     float* h_p_z = (float*) calloc(h_x.size(),sizeof(float)); 
     if(h_p_x==NULL || h_p_y==NULL || h_p_z==NULL) 
     { 
      std::cout << "Host memory allocation failure" << std::endl; 
      return; 
     } 

     float* d_p_x; 
     float* d_p_y; 
     float* d_p_z; 

    if(cudaMalloc((void **)&d_p_x,size) != cudaSuccess || 
    cudaMalloc((void **)&d_p_y,size) != cudaSuccess || 
    cudaMalloc((void **)&d_p_z,size) != cudaSuccess) 
    { 
     std::string errorString(cudaGetErrorName(cudaGetLastError())); 
     std::cout << errorString << std::endl; 
     std::cout << "Device memory allocation failure" << std::endl; 
     return; 
    } 
    curandStateMRG32k3a *devStates; 
    if(cudaMalloc((void **)&devStates, h_x.size() * sizeof(curandStateMRG32k3a)) != cudaSuccess) 
    { 
     std::string errorString(cudaGetErrorName(cudaGetLastError())); 
     std::cout << errorString << std::endl; 
     std::cout << "Random generator states memory allocation failure" << std::endl; 
     return; 
    } 

    int threads = 256; 
    dim3 grid = size/threads; 
    setup_kernel<<<grid,threads>>>(devStates); 

    if(cudaMemcpy(d_p_x,h_p_x,size,cudaMemcpyHostToDevice) != cudaSuccess || 
    cudaMemcpy(d_p_y,h_p_y,size,cudaMemcpyHostToDevice) != cudaSuccess || 
    cudaMemcpy(d_p_z,h_p_z,size,cudaMemcpyHostToDevice) != cudaSuccess) 
    { 
     std::string errorString(cudaGetErrorName(cudaGetLastError())); 
     std::cout << errorString << std::endl; 
     std::cout << "Host to Device memory copy failure" << std::endl; 
    } 

    computeRandomVectors<<< grid, threads >>>(d_p_x,d_p_y,d_p_z,size/sizeof(float), devStates); 

    if(cudaMemcpy(h_p_x,d_p_x,size,cudaMemcpyDeviceToHost) != cudaSuccess || 
    cudaMemcpy(h_p_y,d_p_y,size,cudaMemcpyDeviceToHost) != cudaSuccess || 
    cudaMemcpy(h_p_z,d_p_z,size,cudaMemcpyDeviceToHost) != cudaSuccess) 
    { 
     std::string errorString(cudaGetErrorName(cudaGetLastError())); 
     std::cout << errorString << std::endl;  
     std::cout << "Device to Host memory copy failure" << std::endl; 
    } 
    for(size_t i = 0; i < h_x.size(); ++i) 
    { 
     h_x[i] = h_p_x[i]; 
     h_y[i] = h_p_y[i]; 
     h_z[i] = h_p_z[i]; 
    } 

    free (h_p_x); 
    free (h_p_y); 
    free (h_p_z); 
    cudaFree (devStates); 
    cudaFree (d_p_x); 
    cudaFree (d_p_y); 
    cudaFree (d_p_z); 
    cudaDeviceReset(); 
} 

このコードは、ベクトルの要素数が4000未満であれば（私は1K、2K、3K及び4Kを試みた）動作します。それは私に最初のcudaMemcpyの不正なアドレスの不正なアドレスを与えます。私はメモリが不足しているとは思わないが、gtx 980（4GBのグローバルメモリ）を使って作業している。どのようにこれを修正するための任意のアイデア？

---

EDIT：提案の修正後のコード以下の通りです：

__global__ 
void setup_kernel(curandStateMRG32k3a *state, unsigned int numberOfElements) 
{ 
    int id = threadIdx.x + blockIdx.x * blockDim.x; 
    if(id < numberOfElements) curand_init(0, id, 0, &state[id]); 
} 

__global__ 
void computeRandomVectors(float* x, float* y, float* z, unsigned int numberOfElements,curandStateMRG32k3a *state) 
{ 
    float a,b; 
    unsigned int i = blockDim.x * blockIdx.x + threadIdx.x; 
    curandStateMRG32k3a localState = state[i]; 
    if(i < numberOfElements) 
    { 
     a = curand_uniform(&localState); 
     b = curand_uniform(&localState); 
     while(a * a + b * b > 1.0f) 
    { 
     a = curand_uniform(&localState) * 2.0f - 1.0f; 
     b = curand_uniform(&localState) * 2.0f - 1.0f; 
    } 
     x[i] = 2.0f * a * sqrtf(1.0f - a * a - b * b); 
     y[i] = 2.0f * b * sqrtf(1.0f - a * a - b * b); 
     z[i] = 1.0f - 2.0f * (a * a + b * b); 
    } 
} 

void generatePointsOnASphere(thrust::host_vector<float>& h_x, thrust::host_vector<float>& h_y, thrust::host_vector<float>& h_z) 
{ 
    if(h_x.size() != h_y.size() && h_x.size() != h_z.size()) 
    { 
     std::cout << "The three component vectors have unmatching size()" << std::endl; 
     return; 
    } 

    size_t size = h_x.size() * sizeof(float); 

    float* h_p_x = (float*) calloc(h_x.size(),sizeof(float)); 
    float* h_p_y = (float*) calloc(h_x.size(),sizeof(float)); 
    float* h_p_z = (float*) calloc(h_x.size(),sizeof(float)); 
    if(h_p_x==NULL || h_p_y==NULL || h_p_z==NULL) 
    { 
     std::cout << "Host memory allocation failure" << std::endl; 
     return; 
    } 

    float* d_p_x; 
    float* d_p_y; 
    float* d_p_z; 

    if(cudaMalloc((void **)&d_p_x,size) != cudaSuccess || 
cudaMalloc((void **)&d_p_y,size) != cudaSuccess || 
cudaMalloc((void **)&d_p_z,size) != cudaSuccess) 
    { 
     std::string errorString(cudaGetErrorName(cudaGetLastError())); 
     std::cout << errorString << std::endl; 
     std::cout << "Device memory allocation failure" << std::endl; 
     return; 
    } 
    curandStateMRG32k3a *devStates; 
    if(cudaMalloc((void **)&devStates, h_x.size() * sizeof(curandStateMRG32k3a)) != cudaSuccess) 
    { 
     std::string errorString(cudaGetErrorName(cudaGetLastError())); 
     std::cout << errorString << std::endl; 
     std::cout << "Random generator states memory allocation failure" << std::endl; 
     return; 
    } 

    if(cudaMemcpy(d_p_x,h_p_x,size,cudaMemcpyHostToDevice) != cudaSuccess || 
cudaMemcpy(d_p_y,h_p_y,size,cudaMemcpyHostToDevice) != cudaSuccess || 
cudaMemcpy(d_p_z,h_p_z,size,cudaMemcpyHostToDevice) != cudaSuccess) 
    { 
     std::string errorString(cudaGetErrorName(cudaGetLastError())); 
     std::cout << errorString << std::endl; 
     std::cout << "Host to Device memory copy failure" << std::endl; 
    } 

    int threads = 512; 
    dim3 grid = (h_x.size() + threads - 1)/threads; 
    setup_kernel<<<grid,threads>>>(devStates, size/sizeof(float)); 
    computeRandomVectors<<< grid, threads >>>(d_p_x,d_p_y,d_p_z,size/sizeof(float), devStates); 
    cudaDeviceSynchronize(); 
    if(cudaMemcpy(h_p_x,d_p_x,size,cudaMemcpyDeviceToHost) != cudaSuccess || 
cudaMemcpy(h_p_y,d_p_y,size,cudaMemcpyDeviceToHost) != cudaSuccess || 
cudaMemcpy(h_p_z,d_p_z,size,cudaMemcpyDeviceToHost) != cudaSuccess) 
    { 
     std::string errorString(cudaGetErrorName(cudaGetLastError())); 
     std::cout << errorString << std::endl;  
     std::cout << "Device to Host memory copy failure" << std::endl; 
    } 
    for(size_t i = 0; i < h_x.size(); ++i) 
    { 
     h_x[i] = h_p_x[i]; 
     h_y[i] = h_p_y[i]; 
     h_z[i] = h_p_z[i]; 
    } 

    free (h_p_x); 
    free (h_p_y); 
    free (h_p_z); 
    cudaFree (devStates); 
    cudaFree (d_p_x); 
    cudaFree (d_p_y); 
    cudaFree (d_p_z); 
    cudaDeviceReset(); 
} 

私はここに掲載保つために残念に感じるが、私は私のミスが今あるものを理解することによって、私は私が得るかもしれないと思うと思いますクーダのより良い理解。 これで、h_x.size（）が20kのときに、cudaMemcpyデバイス - >ホストでerrorIllegalAdressが取得されました。私はまだコードが小さなものではなく、大きなものではないことを理解していません。

Answer 1

問題はここにある：

size_t size = h_x.size() * sizeof(float); 

    ... 
    int threads = 256; 
    dim3 grid = size/threads; 

あなたsize変数はバイトの数によってスケーリングされます。したがって、グリッドサイズに使用する正しい変数ではありません。グリッドサイズは、

dim3 grid = h_x.size()/threads; 

などと計算する必要があります。また、ベクトルの長さ（h_x.size()）がthreadsつまり256で割り切れる場合を除いて、この構造体はすべてのcurand状態を適切に初期化しないことにも注意してください。これに対処する方法は、あなたのsetup_kernelにスレッドチェックを入れますカーネル：

__global__ 
void setup_kernel(curandStateMRG32k3a *state, int size) 
{ 
    int id = threadIdx.x + blockIdx.x * blockDim.x; 
    if (id < size) 
    curand_init(0, id, 0, &state[id]); 
} 

とベクトルの大きさをカバーするのに十分なスレッドが起動します。

dim3 grid = (h_x.size()+threads-1)/threads;