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

大黑山修道

已于 2024-04-20 14:21:51 修改

阅读量3.1k

点赞数 4

分类专栏：【15】CUDA 文章标签：学习笔记 php

于 2020-08-20 17:59:16 首次发布

本文链接：https://blog.csdn.net/sazass/article/details/108128820

版权

【15】CUDA 专栏收录该内容

5 篇文章 1 订阅

订阅专栏

在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维组织；

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.

大黑山修道

关注

4
点赞
踩
64

收藏

觉得还不错? 一键收藏
1
评论
【CUDA编程系列】CUDA编程基本入门学习笔记

参考博客：https://www.cnblogs.com/skyfsm/p/9673960.html1.GPU特点核心很多；规则数据结构，高度统一；相互无依赖的大规模数据；不需要被打断的纯净的计算环境；可预测存储模式；在并行化计算表现极佳；非常适合处理图像领域的计算，图像在计算机中的表现形式是矩阵，而很多矩阵运算可以实现并行计算；一般，串行部分跑CPU，并行部分跑GPU，异构计算；2.CUDA线程模型从小到大：**1.Thread：**线程，并行的基本单位**2.Thread B
复制链接

扫一扫

专栏目录