Cuda - copia dalla memoria globale del dispositivo alla memoria texture

Question

Sto provando a eseguire due attività (separate in 2 kernel) sulla GPU utilizzando Cuda e C++. Come input prendo una matrice NxM (memorizzata in memoria sull'host come array float). Utilizzerò quindi un kernel che esegue alcune operazioni su questa matrice per trasformarlo in una matrice NxMxD. Poi ho un secondo kernel che esegue alcune operazioni su questa matrice 3D (ho appena letto i valori, non devo scrivere valori su di esso).

Operare nella memoria texture sembra essere molto più veloce per il mio compito quindi la mia domanda è se è possibile copiare i miei dati dalla memoria globale sul dispositivo dopo il kernel 1 e trasferirlo direttamente alla memoria texture per il kernel 2 senza portarlo di nuovo all'ospite?

UPDATE

ho aggiunto un po 'di codice per illustrare il mio problema in maniera migliore.

Ecco i due kernel. Il primo è solo un segnaposto per ora e replica la matrice 2D in 3D.

__global__ void computeFeatureVector(float* imData3D_dev, int imX, int imY, int imZ) { 

//calculate each thread global index 
int xindex=blockIdx.x*blockDim.x+threadIdx.x; 
int yindex=blockIdx.y*blockDim.y+threadIdx.y;  

#pragma unroll 
for (int z=0; z<imZ; z++) { 
    imData3D_dev[xindex+yindex*imX + z*imX*imY] = tex2D(texImIp,xindex,yindex); 
} 
} 

Il secondo prenderà questa matrice 3D, ora rappresentata come una trama ed eseguire alcune operazioni su di essa. Vuoto per ora.

__global__ void kernel2(float* resData_dev, int imX) { 
//calculate each thread global index 
int xindex=blockIdx.x*blockDim.x+threadIdx.x; 
int yindex=blockIdx.y*blockDim.y+threadIdx.y;  

resData_dev[xindex+yindex*imX] = tex3D(texImIp3D,xindex,yindex, 0); 

return; 
} 

Poi il corpo principale del codice è il seguente:

// declare textures 
texture<float,2,cudaReadModeElementType> texImIp; 
texture<float,3,cudaReadModeElementType> texImIp3D; 

void main_fun() { 

// constants 
int imX = 1024; 
int imY = 768; 
int imZ = 16; 

// input data 
float* imData2D = new float[sizeof(float)*imX*imY];   
for(int x=0; x<imX*imY; x++) 
    imData2D[x] = (float) rand()/RAND_MAX; 

//create channel to describe data type 
cudaArray* carrayImIp; 
cudaChannelFormatDesc channel; 
channel=cudaCreateChannelDesc<float>(); 

//allocate device memory for cuda array 
cudaMallocArray(&carrayImIp,&channel,imX,imY); 

//copy matrix from host to device memory 
cudaMemcpyToArray(carrayImIp,0,0,imData2D,sizeof(float)*imX*imY,cudaMemcpyHostToDevice); 

// Set texture properties 
texImIp.filterMode=cudaFilterModePoint; 
texImIp.addressMode[0]=cudaAddressModeClamp; 
texImIp.addressMode[1]=cudaAddressModeClamp; 

// bind texture reference with cuda array 
cudaBindTextureToArray(texImIp,carrayImIp); 

// kernel params 
dim3 blocknum; 
dim3 blocksize; 
blocksize.x=16; blocksize.y=16; blocksize.z=1; 
blocknum.x=(int)ceil((float)imX/16); 
blocknum.y=(int)ceil((float)imY/16);  

// store output here 
float* imData3D_dev;   
cudaMalloc((void**)&imData3D_dev,sizeof(float)*imX*imY*imZ); 

// execute kernel 
computeFeatureVector<<<blocknum,blocksize>>>(imData3D_dev, imX, imY, imZ); 

//unbind texture reference to free resource 
cudaUnbindTexture(texImIp); 

// check copied ok 
float* imData3D = new float[sizeof(float)*imX*imY*imZ]; 
cudaMemcpy(imData3D,imData3D_dev,sizeof(float)*imX*imY*imZ,cudaMemcpyDeviceToHost);  
cout << " kernel 1" << endl; 
for (int x=0; x<10;x++) 
    cout << imData3D[x] << " "; 
cout << endl; 
delete [] imData3D; 


// 
// kernel 2 
// 


// copy data on device to 3d array 
cudaArray* carrayImIp3D; 
cudaExtent volumesize; 
volumesize = make_cudaExtent(imX, imY, imZ); 
cudaMalloc3DArray(&carrayImIp3D,&channel,volumesize); 
cudaMemcpyToArray(carrayImIp3D,0,0,imData3D_dev,sizeof(float)*imX*imY*imZ,cudaMemcpyDeviceToDevice); 

// texture params and bind 
texImIp3D.filterMode=cudaFilterModePoint; 
texImIp3D.addressMode[0]=cudaAddressModeClamp; 
texImIp3D.addressMode[1]=cudaAddressModeClamp; 
texImIp3D.addressMode[2]=cudaAddressModeClamp; 
cudaBindTextureToArray(texImIp3D,carrayImIp3D,channel); 

// store output here 
float* resData_dev; 
cudaMalloc((void**)&resData_dev,sizeof(float)*imX*imY); 

// kernel 2 
kernel2<<<blocknum,blocksize>>>(resData_dev, imX); 
cudaUnbindTexture(texImIp3D); 

//copy result matrix from device to host memory 
float* resData = new float[sizeof(float)*imX*imY]; 
cudaMemcpy(resData,resData_dev,sizeof(float)*imX*imY,cudaMemcpyDeviceToHost); 

// check copied ok 
cout << " kernel 2" << endl; 
for (int x=0; x<10;x++) 
    cout << resData[x] << " "; 
cout << endl; 


delete [] imData2D; 
delete [] resData; 
cudaFree(imData3D_dev); 
cudaFree(resData_dev); 
cudaFreeArray(carrayImIp); 
cudaFreeArray(carrayImIp3D); 

} 

sono felice che il primo kernel funziona correttamente, ma la matrice imData3D_dev 3D non sembra essere legato correttamente al texImIp3D tessitura .

RISPOSTA

ho risolto il mio problema con cudaMemcpy3D. Ecco un codice rivisto per la seconda parte della funzione principale. imData3D_dev contiene la matrice 3D nella memoria globale dal primo kernel.

cudaArray* carrayImIp3D; 
cudaExtent volumesize; 
volumesize = make_cudaExtent(imX, imY, imZ); 
cudaMalloc3DArray(&carrayImIp3D,&channel,volumesize); 
cudaMemcpy3DParms copyparms={0}; 

copyparms.extent = volumesize; 
copyparms.dstArray = carrayImIp3D; 
copyparms.kind = cudaMemcpyDeviceToDevice; 
copyparms.srcPtr = make_cudaPitchedPtr((void*)imData3D_dev, sizeof(float)*imX,imX,imY); 
cudaMemcpy3D(&copyparms); 

// texture params and bind 
texImIp3D.filterMode=cudaFilterModePoint; 
texImIp3D.addressMode[0]=cudaAddressModeClamp; 
texImIp3D.addressMode[1]=cudaAddressModeClamp; 
texImIp3D.addressMode[2]=cudaAddressModeClamp; 

cudaBindTextureToArray(texImIp3D,carrayImIp3D,channel); 

// store output here 
float* resData_dev; 
cudaMalloc((void**)&resData_dev,sizeof(float)*imX*imY); 

kernel2<<<blocknum,blocksize>>>(resData_dev, imX); 

    // ... clean up 

Answer 1

La denominazione delle varie routine cudaMemcpy è sfortunatamente un po 'contorta. Per il funzionamento su un array 3D è necessario utilizzare cudaMemcpy3D() che (tra gli altri) ha la capacità di copiare da dati 3D in memoria lineare in un array 3D.
cudaMemcpyToArray() è per copiare dati lineari su un array 2D.

Se si utilizza un dispositivo di capacità di calcolo 2.0 o superiore, non si desidera tuttavia utilizzare alcuna delle funzioni cudaMemcpy*(). Utilizzare invece un surface che consente di scrivere direttamente sulla trama senza la necessità di copiare i dati tra i kernel. (Oue ancora bisogno di separare la lettura e la scrittura in due diversi kernel, proprio come ora, dato che la cache di texture non è coerente con le scritture di superficie e viene invalidata solo all'avvio del kernel).

Answer 2

cudaMemcpyToArray() accetta cudaMemcpyDeviceToDevice come parametro genere, quindi dovrebbe essere possibile.