Cuda-从设备全局内存复制到纹理内存
Cuda - copy from device global memory to texture memory
我正试图使用Cuda和C++在GPU上执行两项任务(分为两个内核)。作为输入,我取一个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(©parms);
// 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
当第一次提出这个问题时,各种cudaMemcpy例程的命名曾经有些复杂,但从那以后,Nvidia已经清理掉了。
对于在3D阵列上操作,您需要使用cudaMemcpy3D()
,它(在其他阵列之间)能够将线性存储器中的3D数据复制到3D阵列中cudaMemcpyToArray()
曾经是将线性数据复制到2D阵列所需的函数,但已被弃用,取而代之的是名称更一致的cudaMemcpy2D()
。
但是,如果您使用的设备具有2.0或更高的计算能力,则不希望使用任何cudaMemcpy*()
函数。相反,使用一个曲面,它允许您直接写入纹理,而不需要在内核之间复制任何数据。(您仍然需要将读取和写入两个不同的内核分开,尽管就像现在一样,因为纹理缓存与表面写入不一致,只有在内核启动时才无效)。
cudaMemcpyToArray()
接受cudaMemcpyDeviceToDevice
作为其类型参数,因此这应该是可能的。
- C#中等价的C++内存复制
- 坚持将双指针管理的内存复制到二维数组
- OpenCL 将字符从全局内存复制到本地内存
- 内存C++复制是否将内存地址复制到另一个阵列
- 将内存复制到结构时,exc_bad_access
- 动态内存复制是如何工作的
- Cuda-从设备全局内存复制到纹理内存
- 使用动态内存复制类
- 将内存复制到标准::复制
- 通过指针的内存复制有时会丢失数据
- 内存访问与内存复制
- 使用汇编代码从内存复制到c++中的寄存器
- 执行 SSE 内存复制导致的堆损坏 - CRT 检查找不到任何内容
- 内存复制导致偏移量
- 如何将图像从全局内存复制到openCL的本地内存
- 从cuda 3D内存复制到线性内存:复制的数据不是我所期望的
- 使用内存复制对象时出现双自由或损坏错误
- c++内存复制结构体到字节数组中
- 使用内存复制对象数组
- 将原始数据从内存复制到文件(cuda)