Cuda - vom globalen Speicher des Geräts in den Texturspeicher kopieren

Question

Ich versuche, zwei Aufgaben (getrennt in 2 Kerne) auf der GPU mit Cuda und C++ durchzuführen. Als Eingabe nehme ich eine NxM-Matrix (im Speicher des Hosts als Float-Array gespeichert). Ich werde dann einen Kernel verwenden, der einige Operationen an dieser Matrix ausführt, um daraus eine NxMxD-Matrix zu machen. Ich habe dann einen zweiten Kernel, der einige Operationen an dieser 3D-Matrix durchführt (ich lese nur die Werte, ich muss keine Werte schreiben).

Im Texturspeicher zu arbeiten scheint für meine Aufgabe viel schneller zu sein, also meine Frage ist, ob es möglich ist, meine Daten aus dem globalen Speicher auf dem Gerät nach Kernel 1 zu kopieren und direkt in den Texturspeicher für Kernel 2 zu übertragen zurück zum Gastgeber?

UPDATE

Ich habe einige Code hinzugefügt mein Problem besser zu veranschaulichen.

Hier sind die zwei Kernel. Die erste ist nur ein Platzhalter für den Moment und repliziert die 2D-Matrix 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); 
} 
} 

Die zweite wird diese 3D-Matrix nehmen, jetzt als Textur dargestellt und einige Operationen darauf ausführen. Blank für jetzt.

__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; 
} 

Dann wird der Hauptteil des Codes ist wie folgt:

// 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); 

} 

Im glücklich, dass der erste Kern richtig funktioniert, aber die 3D-Matrix imData3D_dev scheint nicht richtig auf die Textur texImIp3D gebunden zu sein .

ANTWORT

löste ich mein Problem cudaMemcpy3D verwenden. Hier ist der überarbeitete Code für den zweiten Teil der Hauptfunktion. imData3D_dev enthält die 3D-Matrix im globalen Speicher des ersten Kernels.

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

Die Benennung der verschiedenen cudaMemcpy-Routinen ist leider etwas verworren. Um auf einem 3D-Array zu arbeiten, müssen Sie cudaMemcpy3D() verwenden, das (unter anderem) die Möglichkeit bietet, 3D-Daten im linearen Speicher in ein 3D-Array zu kopieren.
cudaMemcpyToArray() ist zum Kopieren von linearen Daten in ein 2D-Array.

Wenn Sie ein Gerät der Rechenleistung 2.0 oder höher verwenden, möchten Sie jedoch keine der Funktionen cudaMemcpy*() verwenden. Verwenden Sie stattdessen eine surface, mit der Sie direkt in die Textur schreiben können, ohne dass zwischen den Kernen Daten kopiert werden müssen. (Sie müssen Lese- und Schreibvorgänge jedoch immer noch in zwei verschiedene Kernel separieren, genau wie Sie es jetzt tun, da der Texturcache nicht mit Oberflächenschreibvorgängen übereinstimmt und nur beim Start des Kernels ungültig wird).

Answer 2

cudaMemcpyToArray() nimmt cudaMemcpyDeviceToDevice als Art Parameter, so sollte es möglich sein.