绑定纹理有两种

http://blog.sina.com.cn/s/blog_685768c30100ua2y.html

http://blog.sina.com.cn/s/blog_604afbb00100n853.html

http://tech.it168.com/a2009/1116/809/000000809439_2.shtml

从GPU申请内存

把数据从CPU复制到GPU

绑定纹理

gpu全局内存只支持单下标访问

比如判断素数的程序中， num[bid*bid * THREAD_NUM + tid] 

1.   cudaMalloc((void**) &gpudata, sizeof(long) * TEST);   
2.  cudaMemcpy(gpudata, data, sizeof(long) * TEST,cudaMemcpyHostToDevice); 

然后在核函数中就可以这样访问 num[bid*bid * THREAD_NUM + tid] 

GPU端线性存储器的使用说明

虽然CUDA的显存分配函数包括1D，2D和3D的形式，但均不支持多下标访问。对于1D线性空间采用cudaMalloc进行分配和cudaMemcpy进行数据拷贝，其使用方式与分页内存的方式基本一致，当然添加了数据拷贝的方向的控制。还有个需要注意的地方主机和设备间数据交换会自动同步，而设备与设备间不会，需要使用cudaThreadSynchronize()。 但对于2D和3D则不同，以2D为例，分配的函数为cudaMallocPitch，由于它不支持双下标寻址也不支持二级指针，其实就是cudaMalloc的对齐形式，但数据访问方式有 大的改变，须采用标准访问形式，即： T*pElement = (T*)((char*)BaseAddress + Row * pitch) +Column; 注意指针BaseAddress仍为一级指针，本人测试过若将其声明为二级指针，按道理采用如下访问方式： T*pElement = (T*)((char*)*BaseAddress + Row * pitch) +Column 此方式在模拟条件下能得到正确结果，但实际设备上无法得到正确输出，这也表征了CUDA的线性存储器本质上与内存的不同。 cudaMemcpy2D是用于2D线性存储器的数据拷贝，函数原型为： cudaMemcpy2D(void* dst，size_t dpitch，const void* src，s ize_tspitch ，size_twidth，size_t height，enum cudaMemcpyKind kind ) 这里需要特别注意width与pitch的区别，width是实际需要拷贝的数据宽度而pitch是2D线性存储空间分配时对齐的行宽，而当数据传递发生在设备与主机之间时， 主机 端pitch==width。 综上我们可以看到，CUDA下对二维线性空间的 访问是不提供多下标支持的，访问时依然是通过计算偏移量得到，不同的地方在于使用pitch对齐后 非常利于实 现coalesce访问。

patch的理解：

　　C语言申请2维内存时，一般是连续存放的。a[y][x]存放在第y*widthofx*sizeof(元素)+x*sizeof(元素)个字节。但在cuda的globalmemory访问中，从256字节对齐的地址(addr=0, 256, 512, ...)开始的连续访问是最有效率的。这样，为了提高内存访问的效率，有了cudaMallocPitch函数。cudaMallocPitch函数分配的内存中，数组的每一行的第一个元素的开始地址都保证是对齐的。因为每行有多少个数据是不确定的widthofx*sizeof(元素)不一定是256的倍数。故此，为保证数组的每一行的第一个元素的开始地址对齐，cudaMallocPitch在分配内存时，每行会多分配一些字节，以保证widthofx*sizeof(元素)+多分配的字节是256的倍数(对齐)。这样，y*widthofx*sizeof(元素)+x*sizeof(元素)来计算a[y][x]的地址就不正确了。而应该是y*[widthofx*sizeof(元素)+多分配的字节]+x*sizeof(元素)。而函数中返回的pitch的值就是widthofx*sizeof(元素)+多分配的字节。

cudaMallocPitch两个函数的用法，先看看cudalibrary中如何定义的这两个函数：

cudaError_t cudaMallocPitch	(	void **	devPtr,
		size_t *	pitch,
		size_t	width,
		size_t	height
	)

Allocates at least widthInBytes * height bytes
of linear memory on the device and returns in *devPtr a
pointer to the allocated memory. The function may pad the
allocation to ensure that corresponding pointers in any given row
will continue to meet the alignment requirements for coalescing as
the address is updated from row to row. The pitch returned
in *pitch by cudaMallocPitch() is
the width in bytes of the allocation. The intended usage
of pitch is
as a separate parameter of the allocation, used to compute
addresses within the 2D array. Given the row and column of an array
element of type T,
the address is computed as:

T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column;

For allocations of 2D arrays, it is recommended that programmers
consider performing pitch allocations using cudaMallocPitch().
Due to pitch alignment restrictions in the hardware, this is
especially true if the application will be performing 2D memory
copies between different regions of device memory (whether linear
memory or CUDA arrays).

 

Parameters:

	devPtr	- Pointer to allocated pitched device memory
	pitch	- Pitch for allocation
	width	- Requested pitched allocation width
	height	- Requested pitched allocation height

 

cudaError_t cudaMemcpy2D	(	void *	dst,
		size_t	dpitch,
		const void *	src,
		size_t	spitch,
		size_t	width,
		size_t	height,
		enum cudaMemcpyKind	kind
	)

 

Copies a matrix (height rows
of width bytes
each) from the memory area pointed to by src to
the memory area pointed to by dst,
where kind is
one ofcudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost,
or cudaMemcpyDeviceToDevice,
and specifies the direction of the copy. dpitch and spitch are
the widths in memory in bytes of the 2D arrays pointed to
by dst and src,
including any padding added to the end of each row. The memory
areas may not overlap. Calling cudaMemcpy2D() with dst and src pointers
that do not match the direction of the copy results in an undefined
behavior. cudaMemcpy2D() returns
an error if dpitch or spitch is
greater than the maximum allowed.

 

 

Parameters:

	dst	- Destination memory address
	dpitch	- Pitch of destination memory
	src	- Source memory address
	spitch	- Pitch of source memory
	width	- Width of matrix transfer (columns in bytes)
	height	- Height of matrix transfer (rows)
	kind	- Type of transfer

 

由此，可以对这两个函数有个充分的认识。此外，cudaMallocPitch和cudaMemcpy2D，一般用于二维数组各维度size不是2的幂次方的问题。使用cudaMallocPitch()那么该数组的对齐、大小、起始址等就自动做好了，其返回的pitch就是真正分配给数组的size(往往大于其真正申请的大小)。

cudaMallocPitch((void**)(&dev_features),
&fea_pitch, sizeof(unsigned char) * sfeaturesw,
sfeaturesh);

cudaChannelFormatDesc feaDesc =
cudaCreateChannelDesc<unsigned
char>();

cudaMemcpy2D(dev_features, fea_pitch, sfeatures,
sizeof(unsigned char) * sfeaturesw, sizeof(unsigned char) *
sfeaturesw, sfeaturesh, cudaMemcpyHostToDevice);

cudaBindTexture2D(NULL, features2D, dev_features, feaDesc,
sfeaturesw, sfeaturesh, fea_pitch);

--------------------------------------------------------------------------------

int sfeatures_size = sizeof(unsigned char) * sfeaturesw *
sfeaturesh;

cudaChannelFormatDesc chDesc2 =
cudaCreateChannelDesc<unsigned
char>();

cudaMallocArray(&featuresArray,
&chDesc2, sfeaturesw, sfeaturesh);

cudaMemcpyToArray( featuresArray, 0, 0, sfeatures,
sfeatures_size, cudaMemcpyHostToDevice );

cudaBindTextureToArray( features2D, featuresArray);
   

-------------------------------------------------------------------------------------

int grid_data_size = sizeof(float) * gridl;

cudaMalloc((void**)&dev_grid,grid_data_size);

cudaMemcpy(dev_grid,sgrid,grid_data_size,cudaMemcpyHostToDevice);

cudaBindTexture(0,gridData1D,dev_grid);

----------------------------------------------------------------------------------

cudaError_t cudaMemcpy2D	(	void *	dst,
		size_t	dpitch,
		const void *	src,
		size_t	spitch,
		size_t	width,
		size_t	height,
		enum cudaMemcpyKind	kind
	)

cudaError_t cudaMallocPitch	(	void **	devPtr,
		size_t *	pitch,
		size_t	width,
		size_t	height
	)

对于一维纹理，不管是 Linear Memory还是使用 cudaMallocPitch的，都可以使用tex1Dfetch和tex1D

而对于二维纹理，不管是cudaArray还是 cudaMallocPitch都是使用tex2D