热门标签
热门文章
- 1洞察数据未来：治理与提取技术的前沿探索_数据治理前沿技术有哪些
- 2Oracle多个表的联接（JOIN）、子查询（Subquery）、聚合函数（Aggregate Functions）、分组（GROUP BY）以及HAVING子句的使用
- 3【Web】CTFSHOW元旦水友赛部分wp_ctfshow元旦水友赛wp
- 4Redis - ziplist (压缩列表) 图文详解
- 5算法系列--动态规划--背包问题(3)--完全背包介绍_动态规划完全背包问题
- 6nlp mrc的损失是什么_自然语言处理是什么，自然语言处理能做什么，主要有哪些难点？...
- 7Column count doesn't match value count at row 1问题解决
- 8ChatGPT最新功能“Text To Speech (TTS，文本转语音)”详细解读！_chatgpttts
- 9Spring整合Hibernate实现JPA持久化_hibernate怎么硬编码
- 10Logisim 运算器设计(HUST)_运算器设计logisim
当前位置: article > 正文
【CUDA编程系列】CUDA编程基本入门学习笔记
作者：天景科技苑 | 2024-07-11 20:41:27
赞
踩
cuda编程
参考博客：
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
推荐阅读
相关标签
