Cuda - 从设备全局内存复制到纹理内存

Posted

技术标签:

【中文标题】Cuda - 从设备全局内存复制到纹理内存【英文标题】:Cuda - copy from device global memory to texture memory 【发布时间】:2012-11-04 02:54:05 【问题描述】:

我正在尝试使用 Cuda 和 C++ 在 GPU 上执行两个任务(分为 2 个内核)。作为输入,我采用 NxM 矩阵(作为浮点数组存储在主机的内存中)。然后,我将使用一个内核对该矩阵执行一些操作,使其成为 NxMxD 矩阵。然后我有第二个内核,它在这个 3D 矩阵上执行一些操作(我只是读取值,我不必向它写入值)。

对我的任务来说,在纹理内存中操作似乎要快得多,所以我的问题是,是否可以在内核 1 之后从设备上的全局内存中复制我的数据,并将其直接传输到内核 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);


第二个将采用这个 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 解决了我的问题。这是主函数第二部分的修改代码。 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

【问题讨论】:

【参考方案1】:

当这个问题第一次被问到时,各种 cudaMemcpy 例程的命名有些令人费解,但后来 Nvidia 已经清理了。

要对 3D 数组进行操作,您需要使用 cudaMemcpy3D(),它(在其他之间)能够将线性内存中的 3D 数据复制到 3D 数组中。cudaMemcpyToArray() 曾经是将线性数据复制到二维数组,但已被弃用,取而代之的是更一致的命名cudaMemcpy2D()

如果您使用的是计算能力为 2.0 或更高版本的设备,则您不想使用任何 cudaMemcpy*() 函数。而是使用surface,它允许您直接写入纹理,而无需在内核之间复制任何数据。 (您仍然需要像现在一样将读取和写入分开到两个不同的内核中,因为纹理缓存与表面写入不一致,并且仅在内核启动时失效)。

【讨论】:

【参考方案2】:

cudaMemcpyToArray() 接受 cudaMemcpyDeviceToDevice 作为其 kind 参数,因此应该可以。

【讨论】:

感谢您的回复。我试过使用 cudaMemcpyToArray() 但它似乎没有为我复制。我已经粘贴了上面的代码。

以上是关于Cuda - 从设备全局内存复制到纹理内存的主要内容,如果未能解决你的问题,请参考以下文章

如何使用推力和 CUDA 流将内存从主机异步复制到设备

如何将嵌套结构的成员复制到 CUDA 设备的内存空间?

分配给设备内存的 CUDA 全局(如 C 语言)动态数组

如何将动态矩阵复制到 CUDA 中的设备内存?

CUDA 通过数组偏移量从设备内存中复制单个元素是不是安全?

将数据从寄存器复制到全局存储器