CUDA cudaMemcpyFromSymbol «недопустимый символ устройства» ошибка?

Question

Итак, я вижу родительский вопрос о том, как скопировать с хоста в постоянную память на GPU, используя cudaMemcpyToSymbol.

Мой вопрос заключается в том, как сделать обратное, скопировав из постоянной памяти устройства на хост, используя cudaMemcpyFromSymbol.

. В следующем минимальном воспроизводимом примере я либо получил

• 1) invalid device symbol ошибка при использовании cudaMemcpyFromSymbol(const_d_a, b, size); или
• 2) получено segmentation fault, если я использую cudaMemcpyFromSymbol(&b, const_d_a, size, cudaMemcpyDeviceToHost).

Я проконсультировался с manual , который предлагает, чтобы я кодировал как в 1), и этот SO вопрос , который предлагает, чтобы я кодировал как в 2).Ни один из них не работает здесь.

Может ли кто-нибудь любезно помочь предложить решение этой проблемы?Должно быть, я что-то неправильно понял ... Спасибо!

Вот код:

// a basic CUDA function to test working with device constant memory
#include <stdio.h>
#include <cuda.h>

const unsigned int N = 10;    // size of vectors

__constant__ float const_d_a[N * sizeof(float)];  

int main()
{
    float * a, * b;  // a and b are vectors. c is the result
    a = (float *)calloc(N, sizeof(float));
    b = (float *)calloc(N, sizeof(float));

    /**************************** Exp 1: sequential ***************************/
    int i;
    int size = N * sizeof(float);
    for (i = 0; i < N; i++){
        a[i] = (float)i / 0.23 + 1;
    }


    // 1. copy a to constant memory
    cudaError_t err = cudaMemcpyToSymbol(const_d_a, a, size);
    if (err != cudaSuccess){
        printf("%s in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);
        exit(EXIT_FAILURE);
    }

    cudaError_t err2 = cudaMemcpyFromSymbol(const_d_a, b, size);
    if (err2 != cudaSuccess){
        printf("%s in %s at line %d\n", cudaGetErrorString(err2), __FILE__, __LINE__);
        exit(EXIT_FAILURE);
    }

    double checksum0, checksum1;
    for (i = 0; i < N; i++){
        checksum0 += a[i];
        checksum1 += b[i];
    }

    printf("Checksum for elements in host memory is %f\n.", checksum0);
    printf("Checksum for elements in constant memory is %f\n.", checksum1);

    return 0;
}

Answer 1

В CUDA различные cudaMemcpy* операции смоделированы после стандартной библиотеки C memcpy.В этой функции первый указатель всегда указатель назначения, а второй указатель всегда указатель источника.Это также верно для всех cudaMemcpy* функций.

Поэтому, если вы хотите сделать cudaMemcpyToSymbol, символ должен быть первым (целевым) аргументом, передаваемым функции (вторым аргументом будетуказатель хоста).Если вы хотите сделать cudaMemcpyFromSymbol, символ должен быть вторым аргументом (исходной позицией), а указатель хоста является первым аргументом.Это не то, что у вас есть здесь:

cudaError_t err2 = cudaMemcpyFromSymbol(const_d_a, b, size);
                                          ^        ^
                                          |       This should be the symbol.
                                          |  
                                   This is supposed to be the host destination pointer.

Это можно выяснить, просмотрев документацию API .

Если мы изменим порядок этих двух аргументов вэта строка кода:

cudaError_t err2 = cudaMemcpyFromSymbol(b, const_d_a, size);

Ваш код будет работать без ошибок, и напечатанные окончательные результаты будут совпадать.

Нет необходимости использовать амперсанд с любым из aили b указатели на эти функции.a и b уже указатели.В приведенном вами примере pi_gpu_h не является указателем.Это обычная переменная.Чтобы скопировать что-то в него, используя cudaMemcpyFromSymbol, необходимо взять адрес этой обычной переменной, потому что функция ожидает указатель (приемника).

В целом, это выглядит неправильно:

__constant__ float const_d_a[N * sizeof(float)];  

Это фактически объявление статического массива, и кроме декоратора __constant__ это должно быть сделано эквивалентно тому, как вы это сделали бы в C или C ++.Нет необходимости умножать N на sizeof(float) здесь, если вы хотите сохранить количество N float.Просто N сам сделает это:

__constant__ float const_d_a[N];

, однако, оставив все как есть, не создаст проблем для размещенного вами кода.