GPU程式设计(4) -- 记忆体管理

GPU记忆体类别

GPU记忆体类别有非常多种,各有所长,如果善用可进一步提升执行效能,参考下图:
http://img2.58codes.com/2024/200019761XFtXtkcRJ.png
图一. GPU记忆体类别,图片来源:『CUDA Tutorial by Jonathan Hui』

比较说明如下表:
http://img2.58codes.com/2024/20001976G7sObYGuVy.png
图二. GPU记忆体类别与比较

测试

区域(Local)变数:可能储存于暂存器(Register)或区域(Local)记忆体
__global__ void gpu_function(int* d_a, int* d_b, int* d_c) {// 区域(Local)变数 int tid = threadIdx.x;    ...}
全局(Global)变数
int* d_a;...cudaMalloc((void**)&d_a, N * sizeof(int));
常数(Constant)
// 宣告常数__constant__ int d_f;__global__ void gpu_function(float *d_in, float *d_out) {// 使用常数d_out[tid] = d_f * d_in[tid];}int main(void) {    // 设定常数    int h_f = 2;    // 複製 h_f 至 d_fcudaMemcpyToSymbol(d_f, &h_f, sizeof(int),0,cudaMemcpyHostToDevice);    gpu_function << <1, N >> >(d_in, d_out);}
共享(Shared)记忆体较常用,下一段详细说明。纹理(Texture)记忆体用于2D/3D图形处理,后续详加说明。

共享记忆体

当多执行绪运算时,会使用到同一个变数时,我们就可以利用共享记忆体,例如『上一篇』矩阵运算,由于全局(Global)记忆体较共享(Shared)记忆体存取较慢,故可将来源矩阵(a、b)複製到共享变数,再进行计算。

__global__ void gpu_inner_product_shared(float* d_a, float* d_b, float* d_c){// 输出所在格子的座标int row = threadIdx.x;int col = threadIdx.y;//Defining Shared Memory__shared__ float shared_a[A_ROW_SIZE][A_COLUMN_SIZE];__shared__ float shared_b[B_ROW_SIZE][B_COLUMN_SIZE];for (int k = 0; k < A_COLUMN_SIZE; k++){shared_a[row][k] = d_a[row * A_COLUMN_SIZE + k];}for (int k = 0; k < B_ROW_SIZE; k++){shared_b[k][col] = d_b[k * B_COLUMN_SIZE + col];//__syncthreads();}// // 点积// printf("row=%d, col=%d\n", row, col);for (int k = 0; k < A_COLUMN_SIZE; k++){// 第一个输入矩阵的【列】与第二个输入矩阵的【行】相乘d_c[row * B_COLUMN_SIZE + col] += shared_a[row][k] * shared_b[k][col];//printf("result=%d, A=%d, B=%d\n", row * B_COLUMN_SIZE + col, row * A_COLUMN_SIZE + k, k * B_COLUMN_SIZE + col);//__syncthreads();}}

宣告与複製动作都在多执行绪完成,不会有重複宣告的困扰,真是太神奇了。

结语

完整程式放在『GitHub』的InnerProduct目录。


关于作者: 网站小编

码农网专注IT技术教程资源分享平台,学习资源下载网站,58码农网包含计算机技术、网站程序源码下载、编程技术论坛、互联网资源下载等产品服务,提供原创、优质、完整内容的专业码农交流分享平台。

热门文章