参加CUDA线上训练营-day02

GPU存储单元

Register寄存器 ：位于GPU的SM上，最快的memory，线程私有，退出失效。存储局部变量。

Shared Memory : 位于GPU芯片上，线程同步。实现Block内的线程通信，目前最快的多Thread沟通的地方。

Local memory ：无特定存储单元，线程私有，速度较慢。存放单线程的大型数组和变量（Register不够时用它）。

Constant memory ：无特定存储单元，固定内存空间驻留在内存，并固定缓存中。用于同一warp的所有thread同时访问同样的常量数据，比如光线追踪。

Global memory：GPU最基础的memory，空间最大。输入数据，写入结果。

Texture memory:驻留在device memory中，使用一个只读cache。用于加速局部性访问。

速度比较（由快到慢）：Register > Shared Memory > Constant memory > Texture memory > 

Local memory 、Global memory

分配与释放

1. 申请存储单元

当我们要为一个方阵M(m * k)申请GPU的存储单元时，使用函数cudaMalloc((void **) &d_a, sizeof(int)*m*k);

d_a:指向存储在Device端数据的地址的指针。(使用双重指针原因：由于cudaMalloc函数返回类型是错误类型cudaError_t，返回值不是单一值。因此需要一个指针指向这个地址，传址改参，而分配内存时候返回一个指针，指向分配地址，因此该参数为指向指针的指针)

szeof(int)*m*k：存储在Device端空间的大小=int型矩阵大小为m*k。

2. GPU与CPU 的数值传递

cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice);

d_a:传输目的地GPU存储单元

h_a:CPU数据

sizeof(int)*m*n：数据传输大小

cudaMemcpyHostToDevice： 数据传输方向

3. 释放申请的存储单元

使用函数：cudaFree(d_a) / cudaFreeHost(h_a) ;

cudaFree（d_a) : 释放基本数据类型值

d_a:指向存储在Device端数据的地址的值

cudaFreeHost(h_a) ：释放数组等

h_a:指向存储在Device端数据的地址的数组指针

矩阵相乘示例

#include <stdio.h>
#include <math.h>

#define BLOCK_SIZE 16

__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{ 
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int sum = 0;
    if( col < k && row < m) 
    {
        for(int i = 0; i < n; i++) 
        {
            sum += a[row * n + i] * b[i * k + col];
        }
        c[row * k + col] = sum;
    }
} 

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0.0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

int main(int argc, char const *argv[])
{
    int m=100;
    int n=100;
    int k=100;

    int *h_a, *h_b, *h_c, *h_cc;
    cudaMallocHost((void **) &h_a, sizeof(int)*m*n);
    cudaMallocHost((void **) &h_b, sizeof(int)*n*k);
    cudaMallocHost((void **) &h_c, sizeof(int)*m*k);
    cudaMallocHost((void **) &h_cc, sizeof(int)*m*k);

    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = rand() % 1024;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = rand() % 1024;
        }
    }

    int *d_a, *d_b, *d_c;
    cudaMalloc((void **) &d_a, sizeof(int)*m*n);
    cudaMalloc((void **) &d_b, sizeof(int)*n*k);
    cudaMalloc((void **) &d_c, sizeof(int)*m*k);


    cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice);

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
   
    gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k);    

    cudaMemcpy(h_c, d_c, sizeof(int)*m*k, cudaMemcpyDeviceToHost);

    cpu_matrix_mult(h_a, h_b, h_cc, m, n, k);

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    cudaFreeHost(h_a);
    cudaFreeHost(h_b);
    cudaFreeHost(h_c);
    cudaFreeHost(h_cc);
    return 0;
}