CUDA编程入门
我们用host指代CPU及其内存,而用device指代GPU及其内存。
- global:在device上执行,从host中调用(一些特定的GPU也可以从device上调用),返回类型必须是void,不支持可变参数参数,不能成为类成员函数。注意用__global__定义的kernel是异步的,这意味着host不会等待kernel执行完就执行下一步。
- device:在device上执行,单仅可以从device中调用,不可以和__global__同时用。
- host:在host上执行,仅可以从host上调用,一般省略不写,不可以和__global__同时用,但可和__device__,此时函数会在device和host都编译。
1.CUDA的整体结构
kernel是在device上线程中并行执行的函数,核函数用__global__符号声明,在调用时需要用<<<grid, block>>>来指定kernel要运行的线程,在CUDA中,每一个线程都要执行核函数。kernel在device上执行时实际上是启动很多线程,一个kernel所启动的所有线程称为一个网格(grid),同一个网格上的线程共享相同的全局内存空间。每一个grid由多个block组成,每一个block由多个线程组成。
从上图可以看出每一个block可以组织成三维的,但其实block可以1维、2维或3维组织。Grid可以1维、2维组织。
2.CUDA内存模型
如下图所示。可以看到,每个线程有自己的私有本地内存(Local Memory),而每个线程块有包含共享内存(Shared Memory),可以被线程块中所有线程共享,其生命周期与线程块一致。此外,所有的线程都可以访问全局内存(Global Memory)。还可以访问一些只读内存块:常量内存(Constant Memory)和纹理内存(Texture Memory)。
3. Streaming Multiprocessor,SM
SM是GPU的处理器,SM可以并发地执行数百个线程。
- 当一个kernel被执行时,它的gird中的线程块被分配到SM上,一个线程块只能在一个SM上被调度。
- SM一般可以调度多个线程块,一个kernel的各个线程块可能被分配多个SM。
- 当线程块被划分到某个SM上时,它将被进一步划分为多个线程束(一个线程束包含32个线程),因为这才是SM的基本执行单元,但是一个SM同时并发的线程束数是有限的
- 由于SM的基本执行单元是包含32个线程的线程束,所以block大小一般要设置为32的倍数。
SM中包含多个SP,一个GPU可以有多个SM(比如16个),最终一个GPU可能包含有上千个SP。
每个线程由每个线程处理器(SP)执行
线程块由多核处理器(SM)执行
一个kernel其实由一个grid来执行,一个kernel一次只能在一个GPU上执行
block是软件概念,一个block只会由一个sm调度,程序员在开发时,通过设定block的属性,告诉GPU硬件,我有多少个线程,线程怎么组织。而具体怎么调度由sm的warps scheduler负责,block一旦被分配好SM,该block就会一直驻留在该SM中,直到执行结束。一个SM可以同时拥有多个blocks,但需要序列执行
4.cuda编程的相关函数
- 在device上申请一定字节大小的显存:
cudaMalloc(void** devPtr, size_t size); // 在device上申请一定字节大小的显存
// devPtr是指向所分配内存的指针
// 与cudaFree函数配合使用
- 内存和显存之间的数据拷贝:
cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);
src指向数据源,而dst是目标区域,count是复制的字节数,其中kind控制复制的方向:cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost及cudaMemcpyDeviceToDevice,如cudaMemcpyHostToDevice将host上数据拷贝到device上。
- 获取线程的id
我们可以用dim3类来表示网格和线程块的组织方式,网格grid可以表示为一维和二维格式,线程块block可以表示为一维、二维和三维的数据格式。
dim3 DimGrid(100, 50); //5000个线程块,维度是100*50
dim3 DimBlock(4, 8, 8); //每个线层块内包含256个线程,线程块内的维度是4*8*8**
最常见的组织方式如下:
dim3 dimGrid(M, N);
dim3 dimBlock(P, Q);
threadId.x = blockIdx.x*blockDim.x+threadIdx.x; // x轴方向上的id号
threadId.y = blockIdx.y*blockDim.y+threadIdx.y; // y轴方向上的id号
5.实现矩阵乘法
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<ctime>
#include<iostream>
#include<ctime>
#include <stdio.h>
#include <math.h>
#include "malloc.h"
using namespace std;
const int Row = 512; // 行数
const int Col = 512; // 列数
__global__
void matrix_mul_gpu(int *M, int* N, int* P, int width) // width代表列数
{
int i = threadIdx.x + blockDim.x * blockIdx.x; // 第i列的线程
int j = threadIdx.y + blockDim.y * blockIdx.y; // 第j行的线程
int sum = 0;
for (int k = 0; k < width; k++)
{
int a = M[j*width + k]; // 第j行的某一个值
int b = N[k*width + i]; // 第i列的某一个值
sum += a * b;
}
P[j*width + i] = sum;
}
void matrix_mul_cpu(int *M, int* N, int* P, int width)
{
for (int i = 0; i < width; i++)
for (int j = 0; j < width; j++)
{
int sum = 0.0;
for (int k = 0; k < width; k++)
{
int a = M[i*width + k];
int b = N[k*width + j];
sum += a * b;
}
P[i*width + j] = sum;
}
}
int main()
{
clock_t GPUstart, GPUend, GPUresult;
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;
}
GPUstart = clock();
cudaMemcpy(d_dataA, A, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
cudaMemcpy(d_dataB, B, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
dim3 threadPerBlock(16, 16);
// (Col + threadPerBlock.x - 1)/threadPerBlock.x=Col/threadPerBlock.x+1,即多拿一个block来装不能整除的部分
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);
// 每一个线程进行某行乘某列的计算,得到结果中的一个元素。也就是d_dataC中的每一个计算结果都和GPU中线程的布局<blockNumber, threadPerBlock >一致
matrix_mul_gpu << <blockNumber, threadPerBlock >> > (d_dataA, d_dataB, d_dataC, Col);
//拷贝计算数据-一级数据指针
cudaMemcpy(C, d_dataC, sizeof(int) * Row * Col, cudaMemcpyDeviceToHost);
//释放内存
free(A);
free(B);
free(C);
cudaFree(d_dataA);
cudaFree(d_dataB);
cudaFree(d_dataC);
GPUend = clock();
int GPUtime = GPUend - GPUstart;
printf("GPU运行时间:%d\n", GPUtime );
// CPU计算
clock_t CPUstart, CPUend, CPUresult;
int *A2 = (int *)malloc(sizeof(int) * Row * Col);
int *B2 = (int *)malloc(sizeof(int) * Row * Col);
int *C2 = (int *)malloc(sizeof(int) * Row * Col);
//set value
for (int i = 0; i < Row*Col; i++) {
A2[i] = 90;
B2[i] = 10;
}
CPUstart = clock();
matrix_mul_cpu(A2, B2, C2, Col);
CPUend = clock();
int CPUtime = CPUend - CPUstart;
printf("CPU运行时间:%d\n", CPUtime);
printf("加速比为:%lf\n",double(CPUtime)/GPUtime);
return 0;
}
【注】.cu文件好像只能以英文命名,不然可能出错