参考博客:
https://www.cnblogs.com/skyfsm/p/9673960.html
https://blog.csdn.net/bruce_0712/article/details/73656087
1.GPU特点
核心很多;
规则数据结构,高度统一;
相互无依赖的大规模数据;
不需要被打断的纯净的计算环境;
可预测存储模式;
在并行化计算表现极佳;
非常适合处理图像领域的计算,图像在计算机中的表现形式是矩阵,而很多矩阵运算可以实现并行计算;
一般,串行部分跑CPU,并行部分跑GPU,异构计算;
2.CUDA线程模型
Device: GPU
Host: CPU
单位术语解释,从小到大:
1.Thread:线程,并行的基本单位
2.Thread Block:线程块,相互合作的线程组,线程块有如下特点:
1.允许彼此同步;
2.可以通过共享内存快速交换数据
3.以1维,2维,或者3维组织;
3.Grid:一组线程块
1.以一维或者二维组织
2.共享全局内存
每个block(线程块)和每个thread都有自己的ID,我们通过相应的索引找到相应的线程和线程块;
threadIdx:线程的ID,可能是一维,二维或三维(因为上级单位block可以三维的组织形式);
blockIdx:线程块的ID,可能是一维或者二维(因为上级单位Grid可以二维的组织形式);
Kernel:在GPU上执行的核心程序,Kernel函数是运行在某个Grid上的;
Kernel的理解:Kernel在device上执行实际上是启动很多线程,一个kernel所启动的所有线程称为一个网络grid,同一个网络上的线程共享相同的全局内存空间;
grid和block都是定义为dim3类型的变量,dim3可以看作包含三个无符号整数(x,y,z)成员的结构体变量,在定义时,缺省值初始化为1;
因此,grid和block可以灵活地定义为1-dim,2-dim,3-dim;
kernel的调用也必须执行<<<grid,block>>>来指定kernel所使用的网络维度和线程块维度,通过这种方式索引到我们想要的线程;
CUDA的<<<grid,block>>>其实就是一个多级索引的方法,第一级索引就是(grid.xIdx,grid.yIdx),第二级索引(block.xIdx,block.yIdx,block.zIdx)来定位到指定的线程;
SP:最基本的处理单元,也称为CUDA core;最后的指令和认为都是在SP上处理的。GPU进行并行计算,也就是很多个SP同时做处理;
SM:多个SP加上一些资源组成的streaming multiprocessor,也叫做GPU的大核;
每个SM包含的SP数量依据GPU架构而不同;
简而言之,SP是线程执行的硬件单位,SM中包含多个SP,一个GPU可以有多个SM(比如16个),最终一个GPU可能包含有上千个SP。这么多核心“同时运行”,速度可想而知,这个引号只是想表明实际上,软件逻辑上是所有SP是并行的,但是物理上并不是所有SP都能同时执行计算(比如我们只有8个SM却有1024个线程块需要调度处理),因为有些会处于挂起,就绪等其他状态,这有关GPU的线程调度。
软件 硬件
Thread------------------------------SP
Thread Block------------------------SM
Grid--------------------------------GPU Device
#每个线程由线程处理器(SP)执行
#线程块由多核处理器(SM)来执行
#一个kernel实际上由一个grid来执行,一个kernel一次只能在一个GPU上执行
一个block只会由一个sm调度,通过设定block的属性,设置线程组织方式;block一旦被分配好SM,该BLOCK就会一直驻留在SM中,直到执行结束。一个SM可以拥有多个blocks,但需要序列执行序列;
3.CUDA线程模型
CUDA中的内存模型分为以下几个层次:
每个线程都用自己的registers(寄存器)
每个线程都有自己的local memory(局部内存)
每个线程块内都有自己的shared memory(共享内存),所有线程块内的所有线程共享这段内存资源
每个grid都有自己的global memory(全局内存),不同线程块的线程都可使用
每个grid都有自己的constant memory(常量内存)和texture memory(纹理内存),不同线程块的线程都可使用
线程访问这几类存储器的速度是register > local memory >shared memory > global memory
4.CUDA编程模型
4.1 编写kernel函数格式如下:
<前缀> <返回值> <函数名>(参数列表···)
例如:__global__ float Func(float *a,float *b):
除了前缀之后,其他部分与c++相似;
一般会有如下三个形式的前缀:
__device__ //表示在GPU上调用,在GPU上执行
__global__ //表示在CPU上调用,在GPU上执行
__host__ //表示在CPU上调用,在CPU上执行
4.2在GPU内存分配回收内存的函数接口:
cudaMalloc(): 在设备端分配global memory
函数原型:
cudaError_t cudaMalloc (void **devPtr, size_t size );
第一个参数:参数传递的是指针的地址
第二个参数:分配的空间尺寸
e.x.
float *device_data = NULL;
size_t size = 1024*sizeof(float);
cudaMalloc((void**)&device_data,size);
cudaFree(): 释放存储空间
e.x.
cudaFree(device_data);
4.3 CPU的数据和GPU端数据做数据传输的函数接口 :
函数原型:
cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count,cudaMemcpyKind kind )
其中cudaMemcpykind的可选类型有:
cudaMemcpyHostToHost //CPU数据传给CPU
cudaMemcpyHossToDevice //CPU数据传给GPU
cudaMemcpyDeviceToHost //GPU数据传给CPU
cudaMemcpuDeviceToDevice //GPU数据传给GPU
e.x.
int *d_dataA,A;
cudaMalloc((void**)&d_dataA, sizeof(int) *1024*1024);//分配GPU内存
for (int i = 0; i < Row*Col; i++)
A[i] = 90;
cudaMemcpy(d_dataA, A, sizeof(int)*1024*1024, cudaMemcpyHostToDevice);//将CPU的数据复制给GPU
4.4 线程组织模型:
dim3 dimGrid(3,4); //定义一个dim3类型的变量,内置参数(3,4),组织为2维;
dim3 dimBlock(1,2,3); //定义一个dim3类型的变量,内置参数(1,2,3),组织为3维;
matrix_mul_gpu <<<dimGrid,dimBlock>>>(a,b,c)
//调用一个名为matrix_mul_gpu 的kernel(需先定义),然后dimGrid变量作为kernel的Grid,dimBlock作为kernel的Block;
4.5 线程号的计算:
特别注意:
threadIdx,blockIdx,blcokDim都是CUDA的内置参数,即本身已经定义过了,无需再定义,直接调用即可:
threadIdx:表示kernel中当前线程所在block内的地址,因为Block是3维的,所以threadIdx.x,threadIdx.y,threadIdx.z分别表示所在Block内三个维度的坐标;
blockIdx:表示kernel中当前线程所在的Block在Grid的地址,因为Grid是2维的,所以blockIdx.x表示所属Block在Grid的x坐标,blockIdx.y表示所属Block在Grid的y坐标;
blcokDim:表示kernel中Grid的维度,blockDim.x表示Grid的横向维度,blockDim.y表示Grid的纵向维度;
情况1:n个block,1个thread;
dim3 dimGrid(N);
dim3 dimBlock(1);
线程号计算方式:
threadId = blockIdx.x
情况2:m*n个block,1个thread
dim3 dimGrid(M,N);
dim3 dimBlock(1);
线程号计算方式:
threadId = blockIdx.y*blockDim.x+blockIdx.x;
情况3:1个block,n个thread
dim3 dimGrid(1);
dim3 dimBlock(N);
线程号计算方式:
threadId = threadIdx.x
情况4:M个block,N个thread
dim3 dimGrid(M);
dim3 dimBlock(N);
线程号计算方式:
threadId = threadIdx.x + blockIdx.x * blockDim.x
情况5:M * N个Block,P * Q个Thread;
最为常见,常用于图像处理领域;
dim3 dimGrid(M*N);
dim3 dimBlock(P*Q);
线程号计算方式:
threadId.x = blockIdx.x*blockDIm.x + threadIdx.x
threadid.y = blockIdx.y*blockDim.y + threadIdx.y
上述公式将线程和线程块索引映射为图像像素坐标的计算方法;
5.CUDA编程示例
5.1:显示cuda的一些配置参数:
#include "device_launch_parameters.h"
#include <iostream>
int main()
{
int deviceCount;
cudaGetDeviceCount(&deviceCount);
for (int i = 0; i<deviceCount; i++)
{
cudaDeviceProp devProp;
cudaGetDeviceProperties(&devProp, i);
std::cout << "使用GPU device " << i << ": " << devProp.name << std::endl;
std::cout << "设备全局内存总量: " << devProp.totalGlobalMem / 1024 / 1024 << "MB" << std::endl;
std::cout << "SM的数量:" << devProp.multiProcessorCount << std::endl;
std::cout << "每个线程块的共享内存大小:" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
std::cout << "每个线程块的最大线程数:" << devProp.maxThreadsPerBlock << std::endl;
std::cout << "设备上一个线程块(Block)种可用的32位寄存器数量: " << devProp.regsPerBlock << std::endl;
std::cout << "每个EM的最大线程数:" << devProp.maxThreadsPerMultiProcessor << std::endl;
std::cout << "每个EM的最大线程束数:" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;
std::cout << "设备上多处理器的数量: " << devProp.multiProcessorCount << std::endl;
std::cout << "======================================================" << std::endl;
}
}
5.2 矩阵加法 CPU与GPU的性能对比:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <math.h>
#define Row 1024
#define Col 1024
long long g_cpu_calc_count;
//定义的kernel函数
__global__ void addKernel(int **C, int **A, int ** B)
{
int idx = threadIdx.x + blockDim.x * blockIdx.x;
int idy = threadIdx.y + blockDim.y * blockIdx.y;
if (idx < Col && idy < Row) {
C[idy][idx] = A[idy][idx] + B[idy][idx];
}
}
void matrix_add_cpu(int** A_ptr, int** B_ptr, int** C_ptr, int width)
{
g_cpu_calc_count = 0;
for (int i = 0; i<width; i++)
for (int j = 0; j<width; j++)
{
C_ptr[i][j] = A_ptr[i][j] + B_ptr[i][j];
g_cpu_calc_count++;
}
}
int main()
{
int *A, **A_ptr, *B, **B_ptr, *C, **C_ptr, **d_A_ptr, **d_B_ptr, **d_C_ptr, *d_A,*d_B,*d_C;
int total_size = Row*Col * sizeof(int);
//在CPU上分配内存
A = (int*)malloc(total_size);
B = (int*)malloc(total_size);
C = (int*)malloc(total_size);
A_ptr = (int**)malloc(Row * sizeof(int*));
B_ptr = (int**)malloc(Row * sizeof(int*));
C_ptr = (int**)malloc(Row * sizeof(int*));
//CPU一维数组初始化
for (int i = 0; i<Row*Col; i++)
{
A[i] = 80;
B[i] = 20;
}
for (int i = 0; i<Row; i++)
{
A_ptr[i] = A + Col*i;
B_ptr[i] = B + Col*i;
C_ptr[i] = C + Col*i;
}
const clock_t cpu_begin_time_2 = clock(); // 开始计时
matrix_add_cpu(A_ptr, B_ptr, C_ptr, Col); // CPU计算
float ms = float(clock() - cpu_begin_time_2);
std::cout << "矩阵加法运算CPU单核总运算次数:" << g_cpu_calc_count << std::endl;
printf("CPU cost_time: %.2f ms \n", ms);
//GPU计算
//set value
for (int i = 0; i < Row*Col; i++) {
A[i] = 90;
B[i] = 10;
}
//将主机指针A指向设备数据位置,目的是让设备二级指针能够指向设备数据一级指针
for (int i = 0; i < Row; i++) {
A_ptr[i] = A + Col * i;
B_ptr[i] = B + Col * i;
C_ptr[i] = C + Col * i;
}
//set value
for (int i = 0; i < Row*Col; i++) {
A[i] = 90;
B[i] = 10;
}
const clock_t gpu_begin_time_2 = clock(); // 开始计时
//malloc device memory
cudaMalloc((void**)&d_A_ptr, sizeof(int **) * Row);
cudaMalloc((void**)&d_B_ptr, sizeof(int **) * Row);
cudaMalloc((void**)&d_C_ptr, sizeof(int **) * Row);
cudaMalloc((void**)&d_A, sizeof(int) *Row*Col);
cudaMalloc((void**)&d_B, sizeof(int) *Row*Col);
cudaMalloc((void**)&d_C, sizeof(int) *Row*Col);
//memcpy host to device
cudaMemcpy(d_A_ptr, A_ptr, sizeof(int*) * Row, cudaMemcpyHostToDevice);
cudaMemcpy(d_B_ptr, B_ptr, sizeof(int*) * Row, cudaMemcpyHostToDevice);
cudaMemcpy(d_C_ptr, C_ptr, sizeof(int*) * Row, cudaMemcpyHostToDevice);
cudaMemcpy(d_A, A, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, B, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
dim3 threadPerBlock_2(16, 16); //定义变量作为kernel的Grid
dim3 blockNumber_2((Col + threadPerBlock_2.x - 1) / threadPerBlock_2.x, (Row + threadPerBlock_2.y - 1) / threadPerBlock_2.y);//定义变量作为kernel的Block
printf("Block(%d,%d) Grid(%d,%d).\n", threadPerBlock_2.x, threadPerBlock_2.y, blockNumber_2.x, blockNumber_2.y);
addKernel << <blockNumber_2, threadPerBlock_2 >> > (d_C_ptr, d_A_ptr, d_B_ptr);
//memcpy device to host
cudaMemcpy(C_ptr, d_C_ptr, sizeof(int) * Row * Col, cudaMemcpyDeviceToHost);
ms = float(clock() - gpu_begin_time_2);
std::cout << "矩阵加法运算所有线程数:" << threadPerBlock_2.x*threadPerBlock_2.y * blockNumber_2.x * blockNumber_2.y << std::endl;
std::cout << "矩阵加法运算GPU单线程运算次数:1" << std::endl;
std::cout << "矩阵加法运算CPU拷贝到GPU数据字节数:" << sizeof(int*) * Row * 3 + sizeof(int) * Row * Col * 2 << std::endl;
std::cout << "矩阵加法运算GPU拷贝到CPU数据字节数:" << sizeof(int) * Row * Col << std::endl;
printf("GPU cost_time: %.2f ms \n", ms);
//释放内存
free(A);
free(B);
free(C);
free(A_ptr );
free(B_ptr );
free(C_ptr );
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
cudaFree(d_A_ptr);
cudaFree(d_B_ptr);
cudaFree(d_C_ptr);
}
结果是CPU的速度远优于GPU的速度;
5.3 矩阵乘法 CPU与GPU的性能对比:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <math.h>
#define Row 1024
#define Col 1024
long long g_cpu_calc_count;
__global__ void matrix_mul_gpu(int *M, int* N, int* P, int width)
{
int i = threadIdx.x + blockDim.x * blockIdx.x;
int j = threadIdx.y + blockDim.y * blockIdx.y;
int sum = 0;
for (int k = 0; k<width; k++)
{
int a = M[j*width + k];
int b = N[k*width + i];
sum += a*b;
}
P[j*width + i] = sum;
}
void matrix_mul_cpu(int* M, int* N, int* P, int width)
{
g_cpu_calc_count = 0;
for (int i = 0; i < width; i++) {
for (int j = 0; j<width; j++)
{
int sum = 0;
for (int k = 0; k<width; k++)
{
int a = M[i*width + k];
int b = N[k*width + j];
sum += a*b;
g_cpu_calc_count++;
}
P[i*width + j] = sum;
}
}
}
int main()
{
//malloc host memory
int *A = (int *)malloc(sizeof(int) * Row * Col);
int *B = (int *)malloc(sizeof(int) * Row * Col);
int *C = (int *)malloc(sizeof(int) * Row * Col);
//malloc device memory
int *d_dataA, *d_dataB, *d_dataC;
cudaMalloc((void**)&d_dataA, sizeof(int) *Row*Col);
cudaMalloc((void**)&d_dataB, sizeof(int) *Row*Col);
cudaMalloc((void**)&d_dataC, sizeof(int) *Row*Col);
//set value
for (int i = 0; i < Row*Col; i++) {
A[i] = 90;
B[i] = 10;
}
// CPU计算
const clock_t cpu_begin_time = clock();
matrix_mul_cpu(A, B, C, Col);
float ms = float(clock() - cpu_begin_time);
std::cout << "矩阵乘法运算CPU单核总运算次数:" << g_cpu_calc_count << std::endl;
printf("CPU cost_time: %.2f ms \n", ms);
//GPU计算
//set value
for (int i = 0; i < Row*Col; i++) {
A[i] = 90;
B[i] = 10;
}
const clock_t gpu_begin_time = clock();
//memcpy host to device
cudaMemcpy(d_dataA, A, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
cudaMemcpy(d_dataB, B, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
dim3 threadPerBlock(16, 16);
dim3 blockNumber((Col + threadPerBlock.x - 1) / threadPerBlock.x, (Row + threadPerBlock.y - 1) / threadPerBlock.y);
printf("Block(%d,%d) Grid(%d,%d).\n", threadPerBlock.x, threadPerBlock.y, blockNumber.x, blockNumber.y);
// gpu start calc
matrix_mul_gpu << <blockNumber, threadPerBlock >> > (d_dataA, d_dataB, d_dataC, Col);
//拷贝数据:GPU->CPU
cudaMemcpy(C, d_dataC, sizeof(int) * Row * Col, cudaMemcpyDeviceToHost);
ms = float(clock() - gpu_begin_time);
std::cout << "矩阵乘法运算所有线程数:" << threadPerBlock.x*threadPerBlock.y * blockNumber.x * blockNumber.y << std::endl;
std::cout << "矩阵乘法运算GPU单线程运算次数:" << Col << std::endl;
std::cout << "矩阵乘法运算CPU拷贝到GPU数据字节数:" << sizeof(int) * Row * Col * 2 << std::endl;
std::cout << "矩阵乘法运算GPU拷贝到CPU数据字节数:" << sizeof(int) * Row * Col << std::endl;
printf("GPU cost_time: %.2f ms \n", ms);
//释放内存
free(A);
free(B);
free(C);
cudaFree(d_dataA);
cudaFree(d_dataB);
cudaFree(d_dataC);
}
结果是GPU的速度远优于CPU的速度;
6 CUDA编程个人心得
1.CUDA编程调用GPU运算,会增加CPU与GPU传输数据的开销,也就是说使用CUDA编程GPU加速,本身就会出现一部分额外开销;若CPU与GPU交互的数据量一定,则在GPU上执行的计算量越大,则使用GPU加速的效果越明显。因此不可盲目地使用CUDA的GPU加速。
2.