6

CudaとC++を使用してGPUで2つのタスク（2つのカーネルに分離）を実行しようとしています。入力として、NxM行列（ホストのメモリにfloat配列として格納されている）を取得します。次に、この行列に対していくつかの操作を実行するカーネルを使用して、NxMxD行列にします。次に、この3Dマトリックスに対していくつかの操作を実行する2番目のカーネルがあります（値を読み取るだけで、値を書き込む必要はありません）。

テクスチャメモリでの操作は私のタスクでははるかに高速であるように思われるので、私の質問は、カーネル1の後にデバイスのグローバルメモリからデータをコピーし、ホストに戻さずにカーネル2のテクスチャメモリに直接転送できるかどうかです。 ？

アップデート

問題をわかりやすく説明するために、いくつかのコードを追加しました。

これが2つのカーネルです。1つ目は、今のところ単なるプレースホルダーであり、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に正しくバインドされていないようです。

答え

cudaMemcpy3Dを使用して問題を解決しました。これは、メイン関数の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
4

2 に答える 2

3

さまざまなcudaMemcpyルーチンの名前は、この質問が最初に尋ねられたときはやや複雑でしたが、それ以来Nvidiaによってクリーンアップされています。

cudaMemcpy3D()3D配列を操作するには、線形メモリ内の3Dデータから3D配列にコピーする機能を（他の配列間で）使用する必要があります。
cudaMemcpyToArray()以前は線形データを2D配列にコピーするために必要な関数でしたが、より一貫性のある名前が付けられたため、非推奨になりましたcudaMemcpy2D()。

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

于 2012-11-15T17:15:39.567 に答える
2

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

于 2012-11-15T13:50:34.253 に答える