具有READ和WRITE的纹理内存

我正在开发一个CUDA应用程序，其中内核必须多次进入全局内存。该内存由所有CTA随机访问(没有局部性，因此不能使用共享内存)。我需要优化它。我听说纹理内存可以缓解这个问题，但内核可以读写纹理内存吗？1D纹理记忆？2D纹理记忆？还有CUDA阵列呢？

CUDA纹理是只读的。纹理读取被缓存。所以性能增益是概率性的。

CUDA Toolkit 3.1以后的版本也有称为Surfaces的可写纹理，但它们仅适用于计算能力>=2.0的设备。曲面就像纹理一样，但优点是它们也可以由内核编写。

曲面只能绑定到使用标志cudaArraySurfaceLoadStore创建的cudaArray。

这是sgarizvi回答的后续内容。

如今，具有计算能力的卡>=2.0比2012更常见，也就是说，在提出这个问题的时候。

下面是一个关于如何使用CUDA表面内存写入纹理的最小示例。

#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] = %fn", i, h_arr[i]);
return 0;
}

这是Farzad回答的后续内容。

Farzad的观点在CUDA C编程指南：中得到了强调

缓存纹理和曲面内存(请参阅设备内存访问)在同一内核调用中，缓存与关于全局内存写入和表面内存写入，因此纹理提取或表面读取到已写入的地址在同一内核调用中通过全局写入或表面写入返回未定义的数据。换句话说，线程可以安全地读取一些纹理或仅当此内存位置已由以前的内核调用或内存副本更新，但如果有以前由来自相同的内核调用。

这意味着可以修改纹理绑定到的全局内存位置，但这不能发生在操作纹理获取的同一内核中。另一方面；书写纹理"；在上面的意义上，在内核之间是可能的，因为纹理缓存在内核启动时被清除，请参阅cuda内核中的add(a，b，c)，使用纹理对象进行a&b-适用于"增量运算"add(a，b，a)？。

下面，我将提供一个示例，其中修改纹理绑定到的全局内存位置。在这个例子中，我用以下方式调用CUDA内核

median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_out, N);
...
square<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_vec, pitch, N);
...
median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_out, N);

在median_filter_periodic_boundary内核中，操作纹理获取，而在square内核中，修改纹理绑定到的全局内存位置。

这是代码：

#include <stdio.h>
#include "TimingGPU.cuh"
#include "Utilities.cuh"
texture<float, 1, cudaReadModeElementType> signal_texture;
#define BLOCKSIZE 32
/*************************************************/
/* KERNEL FUNCTION FOR MEDIAN FILTER CALCULATION */
/*************************************************/
__global__ void median_filter_periodic_boundary(float * __restrict__ d_out, const unsigned int N){
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N) {
float signal_center = tex1D(signal_texture, (float)(tid + 0.5 - 0) / (float)N);
float signal_before = tex1D(signal_texture, (float)(tid + 0.5 - 1) / (float)N);
float signal_after  = tex1D(signal_texture, (float)(tid + 0.5 + 1) / (float)N);
d_out[tid] = (signal_center + signal_before + signal_after) / 3.f;

}
}

/*************************************************/
/* KERNEL FUNCTION FOR MEDIAN FILTER CALCULATION */
/*************************************************/
__global__ void square(float * __restrict__ d_vec, const size_t pitch, const unsigned int N){
unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N) d_vec[tid] = 2.f * tid;
}
/********/
/* MAIN */
/********/
int main() {

const int N = 10;                                                                                
// --- Input/output host array declaration and initialization
float *h_vec = (float *)malloc(N * sizeof(float));
for (int i = 0; i < N; i++) h_vec[i] = (float)i;
// --- Input/output host and device array vectors
size_t pitch;
float *d_vec;   gpuErrchk(cudaMallocPitch(&d_vec, &pitch, N * sizeof(float), 1));
printf("pitch = %in", pitch);
float *d_out;   gpuErrchk(cudaMalloc(&d_out, N * sizeof(float)));
gpuErrchk(cudaMemcpy(d_vec, h_vec, N * sizeof(float), cudaMemcpyHostToDevice));

// --- CUDA texture memory binding and properties definition
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
//Alternatively
//cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
size_t texture_offset = 0;
gpuErrchk(cudaBindTexture2D(&texture_offset, signal_texture, d_vec, channelDesc, N, 1, pitch)); 
signal_texture.normalized = true; 
signal_texture.addressMode[0] = cudaAddressModeWrap;

// --- Median filter kernel execution
median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_out, N);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(h_vec, d_out, N * sizeof(float), cudaMemcpyDeviceToHost));
printf("nnFirst filteringn");
for (int i=0; i<N; i++) printf("h_vec[%i] = %fn", i, h_vec[i]);
// --- Square kernel execution
square<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_vec, pitch, N);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(h_vec, d_vec, N * sizeof(float), cudaMemcpyDeviceToHost));
printf("nnSquaringn");
for (int i=0; i<N; i++) printf("h_vec[%i] = %fn", i, h_vec[i]);
// --- Median filter kernel execution
median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_out, N);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
printf("nnSecond filteringn");
gpuErrchk(cudaMemcpy(h_vec, d_out, N * sizeof(float), cudaMemcpyDeviceToHost));
for (int i=0; i<N; i++) printf("h_vec[%i] = %fn", i, h_vec[i]);
printf("Test finishedn");

return 0;
}

请注意以下内容：

我没有将纹理绑定到cudaArray，因为cudaArrays不能从内核中修改
我不将纹理绑定到cudaMalloced阵列，因为绑定到cudaMalloced阵列的纹理只能由tex1Dfetch获取，而tex1Dfetch不是cudaAddressModeWrap寻址模式，保证信号周期性扩展到其边界之外
我将纹理绑定到cudaMallocPitched数组，因为这使得可以通过tex1D获取纹理，这允许cudaAddressModeWrap寻址模式
我使用规范化坐标，因为只有这些坐标才能启用cudaAddressModeWrap寻址模式

我需要点#2、#3和#4，因为我从正在编写的代码中提取了这个示例。

我建议将您的内存声明为音调线性内存，并将其与纹理绑定。我还没有试验这种新的无粘合剂质地。有人试过吗？

上面提到的纹理mem是通过缓存只读的。将其视为只读内存。因此，需要注意的是，在内核本身中，您不会写入绑定到纹理的内存，因为它可能不会更新到纹理缓存。

我遇到了这个问题，经过一点搜索，我发现这个问题和这个答案很有用。纹理内存基本上是全局内存。纹理内存是指可以与全局内存读取关联的特殊缓存机制。因此内核可以操作绑定到纹理的全局内存。但正如它在所提供的链接中所显示的，并没有像tex1D(ref, x) = 12.0这样的指令。

相关内容

最新更新

热门标签：