具有读写功能的纹理存储器
具有读写功能的纹理存储器
提问于 2012-09-20 17:03:08
我正在开发一个CUDA应用程序，其中内核必须多次访问全局内存。此内存由所有CTA随机访问(无局部性，因此不能使用共享内存)。我需要优化它。我听说纹理内存可以缓解这个问题，但是内核可以读写纹理内存吗?一维纹理内存?二维纹理内存？另外，CUDA阵列又如何呢？
回答 5
Stack Overflow用户
发布于 2012-09-20 17:44:14
CUDA纹理是只读的。纹理读取被缓存。因此，性能增益是概率的。
CUDA Toolkit3.1及更高版本也有可写纹理，称为表面，但它们仅适用于具有计算能力>=2.0的设备。表面就像纹理一样，但优点是它们也可以由内核写入。
曲面只能绑定到使用标志cudaArraySurfaceLoadStore创建的cudaArray。
Stack Overflow用户
发布于 2015-09-01 16:22:44
这是sgarizvi答案的后续。
如今，与2012相比，具有计算能力>=2.0的卡要普遍得多，也就是说，在提出这个问题的时候。
下面是一个关于如何使用CUDA surface memory写入texture的最小示例。
#include <stdio.h>
#include "TimingGPU.cuh"
#include "Utilities.cuh"
surface<void, cudaSurfaceType1D> surfD;
/*******************/
/* KERNEL FUNCTION */
/*******************/
__global__ void SurfaceMemoryWrite(const int N) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
surf1Dwrite((float)tid, surfD, tid * sizeof(float), cudaBoundaryModeTrap);
}
/********/
/* MAIN */
/********/
int main() {
const int N = 10;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
//Alternatively
//cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
cudaArray *d_arr; gpuErrchk(cudaMallocArray(&d_arr, &channelDesc, N, 1, cudaArraySurfaceLoadStore));
gpuErrchk(cudaBindSurfaceToArray(surfD, d_arr));
SurfaceMemoryWrite<<<1, N>>>(N);
float *h_arr = new float[N];
gpuErrchk(cudaMemcpyFromArray(h_arr, d_arr, 0, 0, N * sizeof(float), cudaMemcpyDeviceToHost));
for (int i=0; i<N; i++) printf("h_arr[%i] = %f\n", i, h_arr[i]);
return 0;
}Stack Overflow用户
发布于 2012-09-21 15:34:33
我建议将你的记忆声明为音调线性记忆，并绑定纹理。我还没有尝试新的无绑定纹理。有人试过吗？
如上所述，纹理mem通过缓存是只读的。将其视为只读存储器。因此，请务必注意，在内核本身内，不要写入绑定到纹理的内存，因为它可能不会更新到纹理缓存。
页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持
原文链接：
https://stackoverflow.com/questions/12509346
复制相关文章
相似问题
腾讯云开发者
