我正在开发一个 CUDA 应用程序，其中内核必须多次进入全局内存。此内存由所有 CTA 随机访问（无局部性，因此无法使用共享内存）。我需要优化它。听说纹理内存可以缓解这个问题，但是内核可以读写纹理内存吗？一维纹理记忆？2D纹理记忆？还有 CUDA 数组呢？
问问题
8089 次
5 回答
10
CUDA 纹理是只读的。纹理读取被缓存。所以性能增益是概率性的。
CUDA Toolkit 3.1 及更高版本还具有称为表面的可写纹理，但它们仅适用于计算能力 >=2.0 的设备。表面就像纹理，但优点是它们也可以由内核编写。
曲面只能cudaArray使用 flag 绑定到创建cudaArraySurfaceLoadStore。
于 2012-09-20T09:44:14.290 回答
5
这是 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] = %f\n", i, h_arr[i]);
return 0;
}
于 2015-09-01T08:22:43.533 回答
2
这是法扎德回答的后续。
Farzad 的观点在 CUDA C 编程指南中得到了强调：
纹理和表面内存被缓存（请参阅设备内存访问），并且在同一个内核调用中，缓存在全局内存写入和表面内存写入方面并不保持一致，因此任何纹理获取或表面读取到一个地址在同一个内核调用中通过全局写入或表面写入写入会返回未定义的数据。换句话说，线程可以安全地读取某个纹理或表面内存位置，前提是该内存位置已被先前的内核调用或内存副本更新，但如果它先前已由同一个线程或来自同一线程的另一个线程更新，则不能内核调用。
这意味着可以修改纹理绑定到的全局内存位置，但这不能发生在操作纹理提取的同一个内核中。另一方面，由于纹理缓存在内核启动时被清除，因此上述意义上的“写入纹理”是可能的，请参见cuda kernel for add(a,b,c) using texture objects for a & b -对于“增量操作”添加（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 = %i\n", 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("\n\nFirst filtering\n");
for (int i=0; i<N; i++) printf("h_vec[%i] = %f\n", 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("\n\nSquaring\n");
for (int i=0; i<N; i++) printf("h_vec[%i] = %f\n", i, h_vec[i]);
// --- Median filter kernel execution
median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_out, N);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
printf("\n\nSecond filtering\n");
gpuErrchk(cudaMemcpy(h_vec, d_out, N * sizeof(float), cudaMemcpyDeviceToHost));
for (int i=0; i<N; i++) printf("h_vec[%i] = %f\n", i, h_vec[i]);
printf("Test finished\n");
return 0;
}
请注意以下事项：
- 我没有将纹理绑定到 a
cudaArray，因为cudaArrays 不能从内核中修改； - 我没有将
cudaMalloc纹理绑定到 ed 数组，因为绑定到ed 数组的纹理cudaMalloc只能被获取，tex1Dfetch并且tex1Dfetch不能cudaAddressModeWrap保证信号在其边界之外的周期性扩展的寻址模式； - 我将纹理绑定到
cudaMallocPitched 数组，因为这样可以通过 获取纹理tex1D，这允许cudaAddressModeWrap寻址模式； - 我正在使用归一化坐标，因为只有那些才能启用
cudaAddressModeWrap寻址模式。
我需要点和#2，因为我从正在编写的代码中提取了这个示例。#3#4
于 2015-09-03T15:04:55.480 回答
1
我建议将您的记忆声明为倾斜的线性记忆并将其与纹理绑定。我还没有尝试新的无绑定纹理。有人试过吗？
如前所述，纹理内存是通过缓存只读的。将其视为只读存储器。因此，重要的是要注意在内核本身内，您不要写入绑定到纹理的内存，因为它可能不会更新到纹理缓存。
于 2012-09-21T07:34:33.330 回答