纹理内存:
和常量内存类似的是,纹理内存同样缓存在芯片上,某些情况下可以减少对内存的请求并提供更高效的内存带宽,纹理缓存是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序而设计的
使用纹理内存,首先需要将输入的数据声明为 texture 类型的引用
//这些变量将位于 GPU 上
texture<float> texConstSrc;
texture<float> texInx;
texture<float> texOut;
之后要将这些变量绑定到内存缓冲区,这相当于告诉 CUDA 运行时两件事情:
- 希望将指定的缓冲区作为纹理来使用
- 将纹理引用作为纹理的“名字”
分配了内存之后需要将内存和之前声明的纹理引用进行绑定
cudaMalloc((void**)&dev_inSrc,imageSize);
cudaMalloc((void**)&dev_outSrc,imageSize);
cudaMalloc((void**)&dev_constSrc,imageSize);
cudaBindTexture(Null,texConstSrc,data.dev_constSrc,imageSize);
cudaBindTexture(Null,texInx,data.dev_inSrc,imageSize);
cudaBindTexture(Null,texOut,data.dev_outSrc,imageSize);
编译器需要知道tex1Dfetch() 应该对哪些纹理进行采样,使用 布尔型dstOut 传递给 blend_kernel() ,这个标志会告诉我们使用哪个缓冲区作为输入,哪个作为输出
最后除了要释放全局缓冲区,还要清楚与纹理的绑定
cudaUnbindTexture( texIn);
cudaUnbindTexture( texOut);
cudaUnbindTexture( texConstSrc);
cudaFree(d->dev_inSrc);
cudaFree(d->dev_outSrc);
cudaFree(d->dev_constSrc);