Ошибка при доступе к массиву CUDA из cudaMalloc

Question

Я использую VS2019 и у меня есть графический процессор NVIDIA GeForce. Я попробовал код по этой ссылке: https://towardsdatascience.com/writing-lightning-fast-code-with-cuda-c18677dcdd5f

Однако я хочу попробовать использовать cudaMalloc вместо использования управляемой памяти с cudaMallocManaged

Я пытался код ниже:

__global__
void add(int n, float* x, float* y)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride)
        y[i] = x[i] + y[i];
}


int main()
{
    int N = 1 << 20;   

    float* x, * y;
    cudaMalloc(&x, N * sizeof(float));
    cudaMalloc(&y, N * sizeof(float));
    cudaMemset(x,1.0, N * sizeof(float)); //want to set x as an array of 1.0s
    cudaMemset(y,2.0, N * sizeof(float)); //want to set y as an array of 2.0s

    int device = -1;
    cudaGetDevice(&device);

    int blockSize = 1024;
    int numBlocks = (N + blockSize - 1) / blockSize;
    auto t1 = std::chrono::high_resolution_clock::now();
    add << <numBlocks, blockSize >> > (N, x, y);

    cudaDeviceSynchronize();
    auto t2 = std::chrono::high_resolution_clock::now(); 
    auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(t2 - t1).count(); 

    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(y[i] - 3.0f));
    std::cout << "Max error: " << maxError << std::endl;
    std::cout << "duration CUDA: "<<duration; 

    cudaFree(x);
    cudaFree(y);

    return 0;
}

Но я получаю необработанную ошибку исключения в maxError = fmax(maxError, fabs(y[i] - 3.0f));, я предполагаю, потому что я не использовал cudaMemset правильно? Как мне его изменить?

Answer 1

В произвольном порядке:

1. Память устройства (т. Е. Память, выделенная с помощью cudaMalloc) не может быть доступна напрямую на хосте, поэтому ваши вычисления maxError недопустимы, поскольку y указатель в памяти устройства Чтобы выполнить проверку на наличие ошибок, вам требуется копия y на локальную копию хоста памяти перед запуском l oop
2. cudaMemset устанавливает байты , а не слова ( так же, как обычный memset). Вам нужно либо установить значения на хосте и скопировать их на устройство, либо в другое ядро, либо использовать что-то вроде thrust::fill_n.
3. В духе вашего предыдущего вопрос , как правило, задержка установки при первом вызове ядра имеет место, поэтому выполните разминку перед синхронизацией

Выполнение этих трех вещей дает мне следующее:

int main()
{
    int N = 1 << 20;   

    std::vector<float> xh(N, 1.0f);
    std::vector<float> yh(N, 2.0f);

    float* x, * y;
    cudaMalloc(&x, N * sizeof(float));
    cudaMemcpy(x, &xh[0], N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMalloc(&y, N * sizeof(float));
    cudaMemcpy(y, &yh[0], N * sizeof(float), cudaMemcpyHostToDevice);

    int blockSize, numBlocks;
    cudaOccupancyMaxPotentialBlockSize(&numBlocks, &blockSize, add);

    for(int rep=0; rep<10; rep++) {
        auto t1 = std::chrono::high_resolution_clock::now();
        add << <numBlocks, blockSize >> > (N, x, y);
        cudaDeviceSynchronize();
        auto t2 = std::chrono::high_resolution_clock::now(); 
        auto duration = std::chrono::duration_cast<std::chrono::microseconds>(t2 - t1).count(); 
        std::cout << rep << " duration CUDA: " << duration <<std::endl; 
    }

    cudaMemcpy(&yh[0], y, N * sizeof(float), cudaMemcpyDeviceToHost);

    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(yh[i] - 12.0f));
    std::cout << "Max error: " << maxError << std::endl;

    cudaFree(x);
    cudaFree(y);

    cudaDeviceReset();

    return 0;
}

И его компиляция и запуск дает мне следующее:

$ nvcc -arch=sm_52 -std=c++11 -o devmem devmem.cu
$ ./devmem 
0 duration CUDA: 155
1 duration CUDA: 94
2 duration CUDA: 95
3 duration CUDA: 94
4 duration CUDA: 94
5 duration CUDA: 93
6 duration CUDA: 93
7 duration CUDA: 99
8 duration CUDA: 92
9 duration CUDA: 93
Max error: 0

По сравнению с таймингами в моем последнем ответе вам, вы можете видеть, что использование памяти устройства обеспечивает ускорение по сравнению с управляемой памятью в моей системе. Как всегда, ваши результаты могут отличаться.