当前位置:   article > 正文

【CUDA编程系列】CUDA编程基本入门学习笔记

cuda编程

在VS2015E上配置CUDA10.0

参考博客:
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维组织;
  • 1
  • 2
  • 3

3.Grid:一组线程块

1.以一维或者二维组织
2.共享全局内存
  • 1
  • 2

每个block(线程块)和每个thread都有自己的ID,我们通过相应的索引找到相应的线程和线程块;

threadIdx:线程的ID,可能是一维,二维或三维(因为上级单位block可以三维的组织形式);
blockIdx:线程块的ID,可能是一维或者二维(因为上级单位Grid可以二维的组织形式);
  • 1
  • 2

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上执行
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7

一个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++相似;
  • 1
  • 2
  • 3

一般会有如下三个形式的前缀:

__device__    //表示在GPU上调用,在GPU上执行
__global__    //表示在CPU上调用,在GPU上执行
__host__  	  //表示在CPU上调用,在CPU上执行
  • 1
  • 2
  • 3

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);
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8

cudaFree(): 释放存储空间

e.x.
cudaFree(device_data);
  • 1
  • 2

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
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13

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;
  • 1
  • 2
  • 3
  • 4

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);
  • 1
  • 2

线程号计算方式:

threadId = blockIdx.x
  • 1

情况2:m*n个block,1个thread

dim3  dimGrid(M,N);	
dim3 dimBlock(1);
  • 1
  • 2

线程号计算方式:

threadId = blockIdx.y*blockDim.x+blockIdx.x;
  • 1

情况3:1个block,n个thread

dim3 dimGrid(1);
dim3 dimBlock(N);
  • 1
  • 2

线程号计算方式:

threadId  = threadIdx.x
  • 1

情况4:M个block,N个thread

dim3 dimGrid(M);
dim3 dimBlock(N);
  • 1
  • 2

线程号计算方式:

threadId  = threadIdx.x + blockIdx.x * blockDim.x
  • 1

情况5:M * N个Block,P * Q个Thread;
最为常见,常用于图像处理领域;

dim3 dimGrid(M*N);
dim3 dimBlock(P*Q);
  • 1
  • 2

线程号计算方式:

threadId.x = blockIdx.x*blockDIm.x + threadIdx.x
threadid.y = blockIdx.y*blockDim.y + threadIdx.y
  • 1
  • 2

上述公式将线程和线程块索引映射为图像像素坐标的计算方法;

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;
		
	}
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 19
  • 20
  • 21
  • 22
  • 23
  • 24

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);
}

  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 19
  • 20
  • 21
  • 22
  • 23
  • 24
  • 25
  • 26
  • 27
  • 28
  • 29
  • 30
  • 31
  • 32
  • 33
  • 34
  • 35
  • 36
  • 37
  • 38
  • 39
  • 40
  • 41
  • 42
  • 43
  • 44
  • 45
  • 46
  • 47
  • 48
  • 49
  • 50
  • 51
  • 52
  • 53
  • 54
  • 55
  • 56
  • 57
  • 58
  • 59
  • 60
  • 61
  • 62
  • 63
  • 64
  • 65
  • 66
  • 67
  • 68
  • 69
  • 70
  • 71
  • 72
  • 73
  • 74
  • 75
  • 76
  • 77
  • 78
  • 79
  • 80
  • 81
  • 82
  • 83
  • 84
  • 85
  • 86
  • 87
  • 88
  • 89
  • 90
  • 91
  • 92
  • 93
  • 94
  • 95
  • 96
  • 97
  • 98
  • 99
  • 100
  • 101
  • 102
  • 103
  • 104
  • 105
  • 106
  • 107
  • 108
  • 109
  • 110
  • 111
  • 112
  • 113
  • 114
  • 115
  • 116
  • 117
  • 118
  • 119
  • 120
  • 121
  • 122

结果是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);
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 19
  • 20
  • 21
  • 22
  • 23
  • 24
  • 25
  • 26
  • 27
  • 28
  • 29
  • 30
  • 31
  • 32
  • 33
  • 34
  • 35
  • 36
  • 37
  • 38
  • 39
  • 40
  • 41
  • 42
  • 43
  • 44
  • 45
  • 46
  • 47
  • 48
  • 49
  • 50
  • 51
  • 52
  • 53
  • 54
  • 55
  • 56
  • 57
  • 58
  • 59
  • 60
  • 61
  • 62
  • 63
  • 64
  • 65
  • 66
  • 67
  • 68
  • 69
  • 70
  • 71
  • 72
  • 73
  • 74
  • 75
  • 76
  • 77
  • 78
  • 79
  • 80
  • 81
  • 82
  • 83
  • 84
  • 85
  • 86
  • 87
  • 88
  • 89
  • 90
  • 91
  • 92
  • 93
  • 94
  • 95
  • 96
  • 97
  • 98
  • 99
  • 100
  • 101
  • 102
  • 103

结果是GPU的速度远优于CPU的速度;
在这里插入图片描述

6 CUDA编程个人心得

1.CUDA编程调用GPU运算,会增加CPU与GPU传输数据的开销,也就是说使用CUDA编程GPU加速,本身就会出现一部分额外开销;若CPU与GPU交互的数据量一定,则在GPU上执行的计算量越大,则使用GPU加速的效果越明显。因此不可盲目地使用CUDA的GPU加速。
2.

声明:本文内容由网友自发贡献,不代表【wpsshop博客】立场,版权归原作者所有,本站不承担相应法律责任。如您发现有侵权的内容,请联系我们。转载请注明出处:https://www.wpsshop.cn/w/天景科技苑/article/detail/811554
推荐阅读
相关标签
  

闽ICP备14008679号