GPU记忆体类别
GPU记忆体类别有非常多种,各有所长,如果善用可进一步提升执行效能,参考下图:
图一. GPU记忆体类别,图片来源:『CUDA Tutorial by Jonathan Hui』
比较说明如下表:
图二. 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目录。