Cuda - デバイスのグローバルメモリからテクスチャメモリへコピー

Question

私は、CudaとC++を使用してGPU上で2つのタスク（2つのカーネルに分かれている）を実行しようとしています。入力として、NxM行列（メモリ上に浮動小数点配列として格納されます）を取ります。次に、この行列にいくつかの演算を実行してNxMxD行列にするカーネルを使用します。私はこの3Dマトリックス上でいくつかの操作を実行する2番目のカーネルを持っています（ちょうど値を読み込んで、値を書き込む必要はありません）。

テクスチャメモリでの操作は私の仕事にとってはるかに高速に思えるので、カーネル1の後にデバイスのグローバルメモリからデータをコピーし、カーネル2のテクスチャメモリに直接持ち込むことはできませんホストに戻って？

UPDATE

私はより良い私の問題を説明するためにいくつかのコードを追加しました。

ここに2つのカーネルがあります。最初のものは現在のプレースホルダーであり、2Dマトリックスを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); 
} 
} 

2番目のテクスチャはこの3Dマトリクスを取り、テクスチャとして表現され、いくつかの操作を実行します。今のところ空白。

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

次のように続いて、コードの本体は次のとおりです。

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

} 

イム最初のカーネルが正常に動作しているが、3DマトリックスimData3D_devが正しくテクスチャtexImIp3Dにバインドしていないように見えることを幸せに。

ANSWER

私はcudaMemcpy3Dを使用して私の問題を解決しました。 main関数の2番目の部分の改訂コードです。 imData3D_devには、最初のカーネルからのグローバルメモリ内の3Dマトリックスが含まれています。

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

さまざまなcudaMemcpyルーチンの名前は、残念ながら少し畳み込まれています。 3Dアレイを操作するには、cudaMemcpy3D()を使用する必要があります（他のものとの間で）3Dデータを3Dメモリから3Dアレイにコピーする機能があります。
cudaMemcpyToArray()は、線形データを2D配列にコピーするためのものです。

コンピューティング機能が2.0以上のデバイスを使用している場合は、cudaMemcpy*()機能を使用したくない場合。代わりにsurfaceを使用すると、カーネル間のデータコピーを必要とせずにテクスチャに直接書き込むことができます。 （テクスチャキャッシュはサーフェスの書き込みと一貫性がなく、カーネルの起動時には無効になるため、今のように読み書きを2つの異なるカーネルに分ける必要があります）。

Answer 2

cudaMemcpyToArray()その種類パラメータとしてcudaMemcpyDeviceToDeviceを受け入れ、それは可能なはずです。