为指针数组double *host_array[MN][4]的每一个指针变量分配了空间
现为使编程方便，希望实现以下形式的绑定：（伪代码，以下代码为示意代码，声明方法错误）
texture<int2,1> tx_array[MN][4];
然后用cudaBindTexture();绑定host_array(i)(j)与tx_array(i)(j) （怎么打不出二维数组的形式。。⊙﹏⊙b汗）
其实目标就是为了可以方便的利用 i , j 这样的数组索引的方式组织数据，
请各位不吝赐教，最好有个简单的代码示意，谢谢！
楼主您好：
通过仔细阅读您的文字，（排版真乱），我成功猜测了您的意图。
您试图使用4MN个1维纹理引用，每个都和一段device上的缓冲区绑定。而这些缓冲区的大小可能一样，也可能不一样。
然后您试图直接用tx_array[j][i]的形式，访问到其中的某一个具体的纹理引用。
那么这样做是可以的。但是需要注意3点：
（1）对纹理引用的声明不能写成数组。您可以连续写4MN个。
（2）您可以在device code里使用纹理引用的数组。
（3-1）如果您对(2)中的纹理引用的数组的下表访问是在编译时刻可以确定的，那么没有任何额外代价；
（3-2）如果您对(3)中的纹理引用的数组的下表访问是在编译时刻不能确定的，那么您会有额外的local memory访问的代码（请责备编译器吧）。
根据您的要求，给予您示意代码：
声明：
texture<int,1> tx,tx1,tx2,tx3,tx4,tx5;
在kernel中的访问：
global void 脱泥(…)
{
texture<int,1> s[2][3] = {{tx, tx1, tx2},{tx3,tx4,tx5}};
…;
int a = tex1Dfetch(s[Y坐标][X坐标], 偏移量);
…;
}
请注意如果是Y和X是可以在编译时刻确定的，那么直接编译出来的是无额外代价的。
如果不能，则会引入local traffic的额外代价。
本文不讨论是否可以手工消除额外的local memory traffic代价，也不鼓励尝试消除。
此外，需要注意的是，您不能绑定一个在host memory上的数组到一个纹理(本文不讨论是否真的不能)。您应该考虑将其放置到device memory上。
此外，作为一个特例，如果每个缓冲区的大小都一致，您应当考虑1D Layered Texture + 坐标变换，也许这样更方便点。
感谢您莅临CUDAZone China,
祝您春节愉快！
谢谢，也祝您春节愉快。
能够服务您是我的荣幸！
Kepler与CUDA5上新的texture object允许部分texture数组的使用。如下面的例子：
编译：nvcc -arch=sm_35 texture_array.cu
#include <stdio.h>
#define SIZE 32
#define N 2
global void init(float *a, float b)
{
a[threadIdx.x] = (float)threadIdx.x + b;
}
global void kernel(cudaTextureObject_t *tex)
{
for (int i = 0; i < N; i++) {
float x = tex1Dfetch(tex[i], threadIdx.x);
printf(“(%d,%d,%f)\n”, threadIdx.x, i, x);
}
}
int main() {
// declare and allocate memory
float buffer[N];
for (int i = 0; i < N; i++) {
cudaMalloc(&buffer[i], SIZEsizeof(float));
init<<<1, SIZE>>>(buffer[i], i);
}
// create texture object
cudaTextureObject_t tex[N];
cudaTextureObject_t d_tex;
cudaMalloc(&d_tex, Nsizeof(cudaTextureObject_t));
for (int i = 0; i < N; i++) {
cudaResourceDesc resDesc;
cudaTextureDesc texDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = buffer[i];
resDesc.res.linear.desc.f = cudaChannelFormatKindFloat;
resDesc.res.linear.desc.x = 32; // bits per channel
resDesc.res.linear.sizeInBytes = SIZE*sizeof(float);
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = cudaReadModeElementType;
cudaCreateTextureObject(&tex[i], &resDesc, &texDesc, NULL);
cudaMemcpy(d_tex+i, &tex[i], sizeof(cudaTextureObject_t), cudaMemcpyHostToDevice);
}
kernel<<<1,SIZE>>>(d_tex);
// destroy texture object
for (int i = 0; i < N; i++)
cudaDestroyTextureObject(tex[i]);
for (int i = 0; i < N; i++)
cudaFree(buffer[i]);
cudaFree(d_tex);
printf(“%s\n”, cudaGetErrorString(cudaGetLastError()));
cudaDeviceReset();
}
这个例子，的确可以。3.5的texture object可以作为__global__参数的特性真心不错。 WangPeng的例子真心不错，辛苦了。
除了WangPeng同学的使用texture object来传递给__global__函数外，如果需要，还可以在__device__中传递，并且无需3.x卡：
升级到5.0的toolkit, 然后直接用texture<…>作为你的__device__参数。然后你可以在2.x上编译，运行，为你的__device__来dispatch多个texture reference.
这是CUDA 5.0提供的新型“软件升级”。也许对楼主有用。
需要更多信息，请搜索"CUDA 5.0 + 透明类型指针"等相关关键字，或者在本帖继续回复。
此外，WP同学的一点排版错误，建议您修正。或者大家自动脑补。
忘了提一下，另一个Kepler+CUDA5下更简单解决lz问题的方案是kernel里直接用__ldg来访问全局内存数组。完全可以避免texture数组而达到用texture访问的目的。详见编程指南。
是的。但是ldg不是所有的Kepler卡都支持的。目前普通的消费级别的3.0无法使用。3.5可以。
当然，楼主如果用3.5，请直接用ldg或者constant __restricted__来提示编译器。