Использование глобальной и постоянной памяти в CUDA

Question

Привет, у меня есть следующий фрагмент кода:

#if USE_CONST == 1
    __constant__ double PNT[ SIZE ];    
#else
    __device__ double *PNT;
#endif

и чуть позже у меня есть:

#if USE_CONST == 0
    cudaMalloc((void **)&PNT, sizeof(double)*SIZE);
    cudaMemcpy(PNT, point, sizeof(double)*SIZE, cudaMemcpyHostToDevice);
#else
    cudaMemcpyToSymbol(PNT, point, sizeof(double)*SIZE);
#endif

, тогда как point где-то определено в коде ранее,При работе с USE_CONST=1 все работает как положено, но при работе без него - нет.Я получаю доступ к массиву в своей функции ядра через

PNT[ index ]

Где проблема между обоими вариантами?Спасибо!

Answer 1

Правильное использование cudaMemcpyToSymbol до CUDA 4.0:

cudaMemcpyToSymbol("PNT", point, sizeof(double)*SIZE)

или альтернативно:

double *cpnt;
cudaGetSymbolAddress((void **)&cpnt, "PNT");
cudaMemcpy(cpnt, point, sizeof(double)*SIZE, cudaMemcpyHostToDevice);

, что может быть немного быстрее, если вы планируете получить доступ к символу изAPI хоста более одного раза.

РЕДАКТИРОВАТЬ: неправильно понял вопрос.Для версии с глобальной памятью сделайте что-то похожее со второй версией для постоянной памяти

double *gpnt;
cudaGetSymbolAddress((void **)&gpnt, "PNT");
cudaMemcpy(gpnt, point, sizeof(double)*SIZE.  cudaMemcpyHostToDevice););

Answer 2

Хотя это старый вопрос, я добавляю его для будущих googlers:

Проблема здесь:

cudaMalloc((void **)&PNT, sizeof(double)*SIZE);
cudaMemcpy(PNT, point, sizeof(double)*SIZE, cudaMemcpyHostToDevice);

cudaMalloc записывает в хост-версию PNT, котораяна самом деле это переменная устройства, к которой нельзя обращаться с хоста.Поэтому правильным будет выделить память, скопировать адрес в символ устройства и скопировать память в память, указанную этим символом:

void* memPtr;
cudaMalloc(&memPtr, sizeof(double)*SIZE);
cudaMemcpyToSymbol(PNT, &memPtr, sizeof(memPtr));
// In other places you'll need an additional:
// cudaMemcpyFromSymbol(&memPtr, PNT, sizeof(memPtr));
cudaMemcpy(memPtr, point, sizeof(double)*SIZE, cudaMemcpyHostToDevice);

Проще будет:

#if USE_CONST == 1
    __constant__ double PNT[ SIZE ];    
#else
    __device__ double PNT[ SIZE ];
#endif

// No #if required anymore:
cudaMemcpyToSymbol(PNT, point, sizeof(double)*SIZE);