Cuda - kopia z pamięci globalnej urządzenia do pamięci tekstury

Question

Próbuję wykonać dwa zadania (rozdzielone na 2 jądra) na GPU przy użyciu Cuda i C++. Jako wejście biorę macierz NxM (przechowywaną w pamięci na hoście jako tablicę float). Następnie użyję jądra, które wykonuje niektóre operacje na tej macierzy, aby utworzyć macierz NxMxD. Następnie mam drugie jądro, które wykonuje pewne operacje na tej macierzy 3D (po prostu odczytałem wartości, nie muszę zapisywać do nich wartości).

Operowanie w pamięci tekstury wydaje się być znacznie szybsze dla mojego zadania, więc moje pytanie brzmi, czy możliwe jest skopiowanie moich danych z pamięci globalnej na urządzeniu po jądrze 1 i przeniesienie go bezpośrednio do pamięci tekstury dla jądra 2 bez dostarczania go wróć do hosta?

UPDATE

dodałem trochę kodu do zilustrowania mój problem lepiej.

Oto dwa jądra. Pierwszy to po prostu uchwyt na miejsce i replikuje matrycę 2D w 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); 
} 
} 

Drugi zabierze tę macierz 3D, obecnie przedstawioną jako tekstura i wykona na niej pewne operacje. Blank na razie.

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

Następnie główny korpus kodu jest następujący:

// 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 zadowolony, że pierwszy jądro działa poprawnie, ale imData3D_dev matryca 3D nie wydaje się być związany z teksturą texImIp3D poprawnie .

ODPOWIEDŹ

mi rozwiązać mój problem za pomocą cudaMemcpy3D. Oto poprawiony kod dla drugiej części głównej funkcji. imData3D_dev zawiera macierz 3D w pamięci globalnej z pierwszego jądra.

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

Nazewnictwo różnych procedur cudaMemcpy jest niestety dość skomplikowane. Do pracy na macierzy 3D należy użyć cudaMemcpy3D(), która (między innymi) ma możliwość kopiowania danych 3D z pamięci liniowej do tablicy 3D.
cudaMemcpyToArray() służy do kopiowania danych liniowych do tablicy 2D.

Jeśli używasz urządzenia o mocy obliczeniowej 2.0 lub wyższej, nie chcesz jednak używać żadnej z funkcji cudaMemcpy*(). Zamiast tego użyj surface, który pozwala bezpośrednio pisać do tekstury bez potrzeby kopiowania danych między jądrami. (nadal musisz oddzielić czytanie i pisanie na dwa różne jądra, tak jak teraz, ponieważ pamięć podręczna tekstur nie jest spójna z zapisami na powierzchni i jest tylko unieważniana podczas uruchamiania jądra).

Answer 2

cudaMemcpyToArray() przyjmuje cudaMemcpyDeviceToDevice jako jego rodzaju parametrów, więc powinno być możliwe.