cuda: использование глобальных переменных для памяти устройства

Question

У меня вопрос о том, как использовать правильные переменные в коде cuda. В моей программе много массивов, к которым нужно обращаться в разных функциях, и я хочу избежать их передачи и хочу использовать глобальные переменные и массивы 2D mallocpitch вместо сплющенного массива 1D. Итак, я думаю о чем-то вроде этого:

__device__ double * dataPtr ; 
__device__ size_t dataPitch; 
.... 
int main() 
{ 
double * dataPtrLoc; size_t dataPitchLoc; 
cudaMallocPitch((void**) &dataPtrLoc, &dataPitchLoc, width*sizeof(double), height); 
cudaMemcpyToSymbol(dataPtr, &dataPtrLoc, sizeof(dataPtrLoc)); 
cudaMemcpyToSymbol(dataPitch, &dataPitchLoc, sizeof(dataPitchLoc)); 
... 
} 

Это похоже на хороший способ получить данные глобального 2D-устройства? Можете ли вы дать предложения?

Edit: я сделал эту программу, и она собирает и работает нормально:

#include <stdio.h> 
__device__ int *d_gridPtr; 
__device__ size_t d_gridPitch; 

__device__ int valij(int ii, int jj) 
{ 
    int* row = (int*)((char*)d_gridPtr + ii * d_gridPitch); 
    return (row[jj]); 
} 

__global__ void printval() 
{ 
    int val0, val1, val2, val3; 
    val0= valij(0,0); 
    val1= valij(0,1); 
    val2= valij(1,0); 
    val3= valij(1,1); 
    printf("%d %d %d %d \n", val0, val1, val2, val3); 
} 

int main() 
{ 
    size_t d_gridPitchLoc; 
    int * d_gridPtrLoc; 
    cudaMallocPitch((void**)&d_gridPtrLoc, &d_gridPitchLoc, 2 * sizeof(int), 2); 
    cudaMemcpyToSymbol(d_gridPtr, & d_gridPtrLoc, sizeof(d_gridPtrLoc)); 
    cudaMemcpyToSymbol(d_gridPitch, &d_gridPitchLoc, sizeof(float)); 

    int h_mem[2*2]={0,1,100,4}; 
    size_t hostpitch = 2* sizeof(int); 
    cudaMemcpy2D(d_gridPtrLoc,d_gridPitchLoc,h_mem,hostpitch,2*sizeof(int),2,cudaMemcpyHostToDevice); 

    printval<<<1,1>>>(); 
    cudaDeviceReset(); 
} 

Answer 1

Если все потоки перекоса или блокировать доступ одинакового (например, индекс массива) только для чтения адреса глобальной памяти в то же время , тогда рассмотрите возможность хранения этих глобальных данных только для чтения в массиве памяти __constant__. Если вы записываете данные, вы не можете использовать __constant__.

Если ваш массив доступен только для чтения, а ваш шаблон доступа имеет сильную 2D-локальность (в пределах основы и/или блока), вместо этого используйте текстуры.