关于机器学习:MindSporeCUDA编程六存储单元

37次阅读

共计 2739 个字符，预计需要花费 7 分钟才能阅读完成。

CUDA 的存储单元蕴含以下类型：

如下表所示：名称地位用处应用办法限度备注 Register 寄存器 GPU 的 SM 上存储局部变量每个 SM 上有成千上万个一个线程最大数量为 256 个须要省着用线程公有，最快线程退出则生效 Shared memoryGPU 芯片上实现 Block 内的线程通信，目前最快的多 Thread 沟通的中央__shared__修饰符须要__syncThreads() 同步分为 32 个 banks 须要省着用，会影响流动 warp 数量可被 1 个 block 所有 thread 拜访，次快高带宽，低提早 Local memory 寄存单线程的大型数组和变量（Register 不够时用它）没有特定的存储单元线程公有，速度较慢，速度与 Global memory 靠近 Constant memory 常量内存驻留在 device memory 中用于同一 warp 的所有 thread 同时拜访同样的常量数据，比方光线追踪__constant__修饰符必须在 host 端应用 cudaMemcpyToSymbol 初始化没有特定的存储单元，然而有独自的缓存只读，全局 Global memory 等同于 GPU 显存驻留在 device memory 中输出数据，写入后果全局，速度较慢 Texture memory 纹理内存用于减速局部性拜访，比方热传导模型只读，全局，速度次于 Shared Memory（提早比 Shared Memory 高，带宽比 hared Memory 小）Host memory：可分页内存主机端内存应用 malloc 拜访应用 free 开释不能够应用 DMA 拜访内存页能够置换到磁盘中另一种 Host memory：又称：Page-locked Memory,Zero-Copy Memory 主机端内存应用 cudaMallocHost 拜访应用 cudaFreeHost 开释属于另一种 Global memory 如何应用 Shared Memory 优化 CUDA 利用呢？

Shared Memory 的特点是快的时候特地快，慢的时候特地慢。什么时候快？同一 warp 中所有线程拜访不同的 banks 或者同一 warp 中所有线程读取同一地址（通过播送）什么时候慢？同一 warp 中多个线程拜访同一个 bank 的不同地址（此时将产生 bank conflict）串行拜访请留神：bank conflict 产生的起因就是 warp 的调配和 bank 的调配重叠了：

如何防止 bank conflict，简略的办法是 Padding 法（如同叫做补边）：

通过减少一个空列，让 bank 强行错位，使得每段间断的数据被调配到不同的 bank 中。具体做法很简略：

就是在设置 Shared Memory 的时候，不设置成方阵 BLOCK_SIZE X BLOCK_SIZE, 而设置成 BLOCK_SIZE X (BLOCK_SIZE+1). 最初，咱们能够应用 Shared Memory 优化 mXn, nXk 的矩阵乘的代码，进步访存的效率。具体方法如下：申请两块 Shared Memory，都是 BLOCK_SIZE X BLOCK_SIZE 大小。一个沿着矩阵 mXn 滑动，一个沿着矩阵 nXk 滑动。将子集的后果累加到目标矩阵中：

具体的代码如下：__global__ void gpu_matrix_mult_shared(int d_a, int d_b, int *d_result, int m, int n, int k)
{

__shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE];
__shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE];

int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
int tmp = 0;
int idx;

for (int sub = 0; sub < gridDim.x; ++sub) 
{
    idx = row * n + sub * BLOCK_SIZE + threadIdx.x;
    tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0;
    idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;
    tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0;
    
    __syncthreads();
    for (int k = 0; k < BLOCK_SIZE; ++k) 
    {tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x];
    }
    __syncthreads();}
if(row < n && col < n)
{d_result[row * n + col] = tmp;
}

} 并将后面代码中调用矩阵乘的中央：gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); 改为 gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); 其余不变。编译，执行：

批改 blocksize，将其别离改为 16，8，4，再进行统计汇总：矩阵 MXN（m) 矩阵 NXK(n) 矩阵 NXK(k)blocksizestop-start(ms)100100100321.83286100100100161.2736510010010081.2329210010010043.528651001001006（补测）2.199910010010012（补测）1.34755 从下面的后果来看，blocksize 为 8，16，32 时如同差别不大，然而 blocksize 为 4 的时候速度降得比拟厉害。blocksize 为 4 时，其实并没有产生 bank conflict！而只是因为 4X4，只有 16 个线程，而一个 warp 须要 32 个线程，所以相当于计算时，有一半算力被节约掉了，进而速度慢了一倍。专家建议，至多应该 NXN>32 比拟好。将矩阵从 100 改为 1000 试试，然而发现一旦改为 1000 后，CPU 计算可能算不过去了，须要将 CPU 那局部代码和前面比拟的代码屏蔽掉。

再从新统计：矩阵 MXN（m) 矩阵 NXK(n) 矩阵 NXK(k)blocksizestop-start(ms)10001000100032265.10610001000100016228.091000100010008202.3821000100010004518.3151000100010006（补测）386.17110001000100012（补测）246.29

正文完