Segmentation fault when using cudaMemcpy

Question

I'm trying to use cudaMemcpy to a std::vector::data to an array for a device kernel and it gives set fault error. The way I do it is:

  cudaMemcpy(d_x, vx.data(), N*sizeof(float), cudaMemcpyHostToDevice);

where vx is vector. The following is the complete example. Any hints on where the problem are would be appreciated.

#include <iostream>
#include <math.h>
#include <vector>

using namespace std;

// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    if(i < n) {
        y[i] = x[i] + y[i];
    }
}


int main(void)
{
    int N = 1<<10;
    float *d_x = NULL, *d_y = NULL;
    cudaMalloc((void **)&d_x, sizeof(float)*N);
    cudaMalloc((void **)&d_y, sizeof(float)*N);

    // Allocate Unified Memory – accessible from CPU or GPU
    vector<float> vx;
    vector<float> vy;

    // initialize x and y arrays on the host
    for (int i = 0; i < N; i++) {
        vx.push_back(1.0f);
        vy.push_back(2.0f);
    }

    cudaMemcpy(d_x, vx.data(), N*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, vy.data(), N*sizeof(float), cudaMemcpyHostToDevice);
    //
    int blockSize;   // The launch configurator returned block size
    int minGridSize; // The minimum grid size needed to achieve the
    // maximum occupancy for a full device launch
    int gridSize;    // The actual grid size needed, based on input size

    cudaOccupancyMaxPotentialBlockSize( &minGridSize, &blockSize, add, 0, N);
    // Round up according to array size
    gridSize = (N + blockSize - 1) / blockSize;

    cout<<"blockSize: "<<blockSize<<" minGridSize: "<<minGridSize<<" gridSize: "<<gridSize<<endl;

    // calculate theoretical occupancy
    int maxActiveBlocks;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor( &maxActiveBlocks, add, blockSize, 0);

    int device;
    cudaDeviceProp props;
    cudaGetDevice(&device);
    cudaGetDeviceProperties(&props, device);

    float occupancy = (maxActiveBlocks * blockSize / props.warpSize) /
        (float)(props.maxThreadsPerMultiProcessor /
                props.warpSize);

    printf("Launched blocks of size %d. Theoretical occupancy: %f\n",
            blockSize, occupancy);


    // Run kernel on 1M elements on the GPU
    add<<<gridSize, blockSize>>>(N, d_x, d_y);

    // Wait for GPU to finish before accessing on host
    cudaDeviceSynchronize();

    // Check for errors (all values should be 3.0f)
    float maxError = 0.0f;

    for (int i = 0; i < N; i++) {
        maxError = fmax(maxError, fabs(d_y[i]-3.0f));
    }
    std::cout << "Max error: " << maxError << std::endl;

    // Free memory
    cudaFree(d_x);
    cudaFree(d_y);

    return 0;
}


blockSize: 1024 minGridSize: 16 gridSize: 1
Launched blocks of size 1024. Theoretical occupancy: 1.000000
Segmentation fault (core dumped)

Oblivion Oblivion · Accepted Answer · 2019-10-15T20:36:31

The problem is here:

for (int i = 0; i < N; i++) {
   maxError = fmax(maxError, fabs(d_y[i]-3.0f));
                                  ^^^^^^
}

And the reason is you cannot dereference device pointer on host.

The solution is copy device memory to host similar to what you did for host to device.

Segmentation fault when using cudaMemcpy

1 Answers