乐趣区

关于机器学习:MindSporeCUDA编程四Global-Memory

jiezi

在 GPU 上，on-board memory 蕴含以下类型：local memory 每个 thread 一个。线程公有。global memory 每个 grid 一个。每个 thread 都能够读。constant memory 每个 grid 一个。只读。每个 thread 都能够读。texture memory 每个 grid 一个。只读。每个 thread 都能够读。on-chip memory 蕴含以下类型：registers 每个 thread 一个。线程公有。shared memory 每个 block 一个，一个 block 下所有线程都能够拜访。HOST 内存函数 malloc 申请 memset 初始化 free 开释 DEVICE 内存函数 cudaMalloc 申请 cudaMemset 初始化 cudaFree 开释请留神，这里函数只返回状态。所以调配的内存地址作为函数参数。HOST《-》DEVICE 相互拷贝 cudaMemcpy(目标内存地址，源内存地址，内存大小，cudaMemcpyHostToDevice/cudaMemcpyDeviceToHost/cudaMemcpyDeviceToDevice/cudaMemcpyHostToHost) 以矩阵乘为例:

CPU 的做法是嵌套循环，如上图所示。GPU 的做法应该是应用 index(blockIdx 和 threadIdx 的组合公式）替换原来的下标 i,j。这也是个别 CUDA 程序的套路——把 for loop 开展成每个线程解决其中的一步。那么，如何应用 CUDA 将坐标拆开呢？将二维坐标（矩阵）改为 在全局中的索引：须要找到每个线程须要解决元素的地位。

ty= 线程在 y 方向的坐标 tx= 线程在 x 方向的坐标 ty=blockIdx.yblockDim.y + threadIdx.ytx=blockIdx.xblockDim.x + threadIdx.xnx= x 方向有多少数据。index = ty * nx + tx 目标是将高维降为低维。

矩阵乘的每个核函数的算法如下：

典型的核函数算法代码如下：

须要留神：矩阵乘 矩阵 M 是 mXn，矩阵 N 是 nXk，这外面须要 矩阵 M 和矩阵 N 都有 n。否则无奈相乘。上代码：matrix_mul.cu#include <stdio.h>

include <math.h>

define BLOCK_SIZE 16

// 应用 GPU 进行矩阵计算
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;
}

}

// 应用 CPU 进行矩阵计算
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[])
{

/* 矩阵 A mXn，矩阵 B nXk --》矩阵乘计算的后果是 mXk */
int m=3;
int n=4;
int k=5;

int *h_a, *h_b, *h_c, *h_cc;

// 调配原矩阵的内存 h 是 host memory
cudaMallocHost((void **) &h_a, sizeof(int)*m*n);
cudaMallocHost((void **) &h_b, sizeof(int)*n*k);

// 调配 CPU 后果内存
cudaMallocHost((void **) &h_c, sizeof(int)*m*k);

// 调配 GPU 后果内存
cudaMallocHost((void **) &h_cc, sizeof(int)*m*k);


// 初始化矩阵 A（mxn）srand(time(0));
printf("---------------h_a------------------\n");
for (int i = 0; i < m; ++i) {for (int j = 0; j < n; ++j) {h_a[i * n + j] = rand() % 1024;
        printf("%d",  h_a[i * n + j] );
        printf(" ");
    }
    printf("\n");
}


// 初始化矩阵 B（nxk）printf("---------------h_b------------------\n");
for (int i = 0; i < n; ++i) {for (int j = 0; j < k; ++j) {h_b[i * k + j] = rand() % 1024;
        printf("%d",  h_b[i * k + j] );
        printf(" ");
    }
    printf("\n");
}

int *d_a, *d_b, *d_c;

// 调配 原矩阵的 GPU 内存 d 是 device memory
cudaMalloc((void **) &d_a, sizeof(int)*m*n);
cudaMalloc((void **) &d_b, sizeof(int)*n*k);

// 调配 目标矩阵的 GPU 内存
cudaMalloc((void **) &d_c, sizeof(int)*m*k);

// copy matrix A and B from host to device memory
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 计算, 后果放入 h_c
gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k);    

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

//CPU 计算，后果间接放入 h_cc
cpu_matrix_mult(h_a, h_b, h_cc, m, n, k);

int ok = 1;
for (int i = 0; i < m; ++i)
{for (int j = 0; j < k; ++j)
    {
        // 比拟大小的时候应用 a-b<0.0000000001 
        if(fabs(h_cc[i*k + j] - h_c[i*k + j])>(1.0e-10))
        {ok = 0;}
    }
}

 printf("---------------h_c  cpu result------------------\n");
 for(int i=0;i<m;i++)
    {for(int j=0;j<k;j++)
        {
            // 矩阵小的时候还能够打印，大的时候就别打了
            printf("%d",h_c[i*k + j] );
            printf(" ");
        }
       printf("\n");
    }
  

 printf("---------------h_cc gpu result----------------\n");
 for(int i=0;i<m;i++)
    {for(int j=0;j<k;j++)
        {
            // 矩阵小的时候还能够打印，大的时候就别打了
            printf("%d",h_cc[i*k + j] );
            printf(" ");
        }
       printf("\n");
    }
  

if(ok)
{printf("Pass!!!\n");
}
else
{printf("Error!!!\n");
}

// free memory
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
cudaFreeHost(h_a);
cudaFreeHost(h_b);
cudaFreeHost(h_c);
return 0;

} 代码中张小白加上了正文，曾经介绍得比较清楚了。咱们执行下看看：

代码以 3X4 和 4X5 的矩阵相乘，失去了 3X5 的矩阵后果。这个后果跟 CPU 计算的后果做了比照。显示 Pass 示意后果是统一的（其实张小白把两个后果都打印的进去，当然也是统一的）这外面有个小 TIPS，就是在调用 rand() 生成随机数的时候，能够应用 srand(time(0)) 做随机数种子，这样下次调用的时候跟这次生成的内容就会不一样。如果去掉这句话，每次执行的后果都是一样的。当然，如果在同一秒同时执行，srand(time(0)) 也会导致同时生成的随机数是一样的。这点须要留神。

退出移动版