Хост и устройство CUDA, использующие одну и ту же __constant__ память

Question

У меня есть функция устройства / хоста, которая использует постоянную память. Он работает нормально на устройстве, но на хосте кажется, что эта память остается неинициализированной.

#include <iostream>
#include <stdio.h>


const __constant__ double vals[2] = { 0.0, 1000.0 };

__device__ __host__ double f(size_t i)
{
    return vals[i];
}

__global__ void kern()
{
    printf("vals[%d] = %lf\n", threadIdx.x, vals[threadIdx.x]);
}

int main() {
    std::cerr << f(0) << " " << f(1) << std::endl;
    kern<<<1, 2>>>();
    cudaThreadSynchronize();
}

Печать (требуется CC 2.0 или выше)

0 0
vals[0] = 0.000000
vals[1] = 1000.000000

В чем проблема и как можно одновременно инициализировать константы памяти устройства и хоста?

Answer 1

Поскольку CygnusX1 неправильно понял, что я имел в виду в своем комментарии к ответу MurphEngineer, возможно, я должен опубликовать свой собственный ответ.То, что я имел в виду , было таким:

__constant__ double dc_vals[2] = { 0.0, 1000.0 };
       const double hc_vals[2] = { 0.0, 1000.0 };

__device__ __host__ double f(size_t i)
{
#ifdef __CUDA_ARCH__
    return dc_vals[i];
#else
    return hc_vals[i];
#endif
}

Это имеет тот же результат, что и Cygnus ', но он более гибок перед лицом реального кода: он позволяет вам иметь значения, определенные во время выполненияв ваших постоянных массивах, например, и позволяет вам использовать функции API CUDA, такие как cudaMemcpyToSymbol / cudsaMemcpyFromSymbol в массиве __constant__.

Более реалистичный полный пример:

#include <iostream>
#include <stdio.h>

__constant__ double dc_vals[2];
       const double hc_vals[2];

__device__ __host__ double f(size_t i)
{
#ifdef __CUDA_ARCH__
    return dc_vals[i];
#else
    return hc_vals[i];
#endif
}

__global__ void kern()
{
    printf("vals[%d] = %lf\n", threadIdx.x, vals[threadIdx.x]);
}

int main() {
    hc_vals[0] = 0.0;
    hc_vals[1] = 1000.0;

    cudaMemcpyToSymbol(dc_vals, hc_vals, 2 * sizeof(double), 0, cudaMemcpyHostToDevice);

    std::cerr << f(0) << " " << f(1) << std::endl;
    kern<<<1, 2>>>();
    cudaThreadSynchronize();
}

Answer 2

Использование спецификатора __constant__ явно выделяет эту память на устройстве.Невозможно получить доступ к этой памяти с хоста - даже с помощью нового компонента CUDA Unified Addressing (который работает только для памяти, выделенной с помощью cudaMalloc () и его друзей).Квалификация переменной с помощью const просто говорит: «это постоянный указатель на (...)».

Правильный способ сделать это, действительно, иметь два массива: один на хосте, а другой наУстройство.Инициализируйте ваш хост-массив, затем используйте cudaMemcpyToSymbol () для копирования данных в массив устройств во время выполнения.Для получения дополнительной информации о том, как это сделать, см. Эту ветку: http://forums.nvidia.com/index.php?showtopic=69724

Answer 3

Я думаю, что MurphEngineer хорошо объяснил почему не работает.

Чтобы быстро решить эту проблему, вы можете следовать идее гарнизма, примерно так:

#ifdef __CUDA_ARCH__
#define CONSTANT __constant__
#else
#define CONSTANT
#endif

const CONSTANT double vals[2] = { 0.0, 1000.0 };

Таким образом, компиляция хоста создаст обычный константный массив хоста, а компиляция устройства создаст __constant__ компиляцию устройства.

Обратите внимание, что с этим приемом может быть сложнее использовать CUDA API для доступа к массиву устройств с функциями, такими как cudaMemcpyToSymbol(), если вы когда-нибудь решите это сделать.

Answer 4

Абсолютно отлично. Я боролся с той же проблемой, и это дает решение. Однако код, предложенный harrism, выдает ошибки при компиляции. Вот исправленный код, который правильно компилируется с nvcc:

#include <iostream>
#include <stdio.h>

__constant__ double dc_vals[2];
       const double hc_vals[2] = {0.0, 1000.0};

__device__ __host__ double f(size_t i)
{
#ifdef __CUDA_ARCH__
    return dc_vals[i];
#else
    return hc_vals[i];
#endif
}

__global__ void kern()
{

    printf("Device: vals[%d] = %lf\n", threadIdx.x, f(threadIdx.x));
}

int main() {

    cudaMemcpyToSymbol(dc_vals, hc_vals, 2 * sizeof(double), 0, cudaMemcpyHostToDevice);

    std::cerr << "Host: " << f(0) << " " << f(1) << std::endl;
    kern<<<1, 2>>>();
    cudaThreadSynchronize();
}