标签：上下文 其它 结构 必须 set 方式 共享 port api
- 需要有支持CUDA的Nvidia显卡
linux查看显卡信息：lspci | grep -i vga
使用nvidia显卡可以这样查看：lspci | grep -i nvidia
上一个命令可以得到类似"03.00.0"的显卡代号，查看详细信息：lspic -v -s 03.00.0
查看显卡使用情况（nvidia专用）：nvidia-smi
持续周期性输出使用情况（1秒1次）：watch -n 1 nvidia-smi- 需要安装pycuda（linux安装：apt install python3-pycuda）
- 准备工作: sudo apt install python3-pycuda
- 在cpu使用的内存中创建矩阵
- 将矩阵从cpu使用的内存移动到gpu的显存中
- 编辑c语言代码，让gpu计算
- 将矩阵从gpu显存移动到cpu使用的内存中
内核与线程层级：
CUDA程序的一个最重要元素就是内核（kernel），它代表可以并行执行的代码
每个内核的执行均有叫做线程（thread）的计算单元完成，与cpu的线程不同，gpu线程更加轻量，上下文切换不会影响性能
为了确定运行一个内核所需的线程数机器逻辑组织形式，CUDA定义了一个二层结构。在最高一层，定义了所谓的区块网格（grid of blocks），这个网格代表了线程区块所在的二维结构，而这些线程区块则是三维的（简单来说，一个cuda结构包含多个blocks，每个blocks包含多个thread）（在下面还会对每个线程区块细分操作）
一个线程区块会被指派给一个流式多处理器(SM)，然后这些线程被进一步划分为被称为warp的线程组，其大小有GPU的架构决定
为了充分发挥SM本身的并发性，同一组内的线程必须执行相同的指令，否则会出现线程分歧（divergence of thread）
示例：
用gpu将矩阵每个元素x2
import pycuda.driver as cuda
import pycuda.autoinit # init GPU
from pycuda.compiler import SourceModule
import numpy as np
# 1.cpu create matrix
a = np.random.randn(5, 5) # matrix:m*n
a = a.astype(np.float32) # nvidia only support float calculate
# 2.move to gpu from cpu
a_gpu = cuda.mem_alloc(a.nbytes) # alloc memory of gpu, this is 1 dim
cuda.memcpy_htod(a_gpu, a) # copy cpu memory to gpu memory
# 3.gpu calculate
# create module of gpu calculate by c
mod = SourceModule('''
__global__ void doubleMatrix(float *a)
{
int idx = threadIdx.x + threadIdx.y * 5; // (x,y,z), gpu -> sm -> warp
a[idx] *= 2;
}
''')
func = mod.get_function('doubleMatrix') # get function from module
func(a_gpu, block = (5, 5, 1)) # set var and thread number from (x,y,z) orient to function
# 4.move to cpu from gpu
a_doubled = np.empty_like(a) # create memory of cpu
cuda.memcpy_dtoh(a_doubled, a_gpu) # copy gpu memory to cpu memory
print('original matrix')
print(a)
print('double matrix')
print(a_doubled)注1：import pycuda.autoinit 语句自动根据GPU可用性和数量选择要使用的GPU，这将创建一个接下来的代码运行中所需的GPU上下文（只需导入即可完成）
注: astype(numpy.float32)：将矩阵中的项转换为单精度模式，因为许多Nvidia显卡只支持单精度
注3：在调用gpu的c函数时，通过block参数设定，分配线程的方式，（5, 5, 1）是对应这gpu的(x, y, x)的分配
在c函数中，threadIdx是一个结构体，它有三个字段x、y、z，每个线程中的这个变量都不同（结合线程层级理解），故用此索引数组，由于动态分配的gpu内存是一维数组，所以需要在c函数内，使用threadIdx.y乘以矩阵每行的元素个数转换
更多详情：自行搜索cuda的线程区块划分
注4：gpu执行的c函数中 __global__ 关键字表示该函数是一个内核函数，必须从主机上调用才能在gpu设备上生成线程层级
涉及到pycuda内存，为了最大限度地利用可用资源，在支持CUDA的GPU显卡中，有4类内存：
寄存器(registers)：每个线程将被分配一个寄存器，每个线程只能访问自身的寄存器，即使同属于一个线程区块
共享存储器(shared memory)：在共享存储器中，每个线程区块都有一个其内部线程共享的内存，这部分内存速度极快
常数存储器(constant memory)：一个网格中的所有线程一直都可以访问这部分内存，但只能在读取时访问。常数存储器中的数据在应用持续期间一直存在
全局存储器(global memory)，：所有网格中的线程（也包含所有内核）都可访问全局存储器
更多详情：自行搜索PyCUDA内存模型
用gpuarray调用内核，它可以直接将数据保存在计算设备(gpu)中，并在该设备中进行计算
示例：
用gpu将矩阵每个元素x2
import pycuda.autoinit
import pycuda.gpuarray as gpuarray
import numpy as np
x = np.random.randn(5, 5)
x_gpu = gpuarray.to_gpu(x)
x_doubled = (2 * x_gpu).get()
print('x:')
print(x)
print('x_doubled:')
print(x_doubled)还有其它的库可以方便对cuda编程，如NumbaPro是一个Python编译器，提供了基于CUDA的API编程接口，可以编写CUDA持续，它专门设计用来执行与数组相关的计算任务，和广泛使用的numpy库类似
NumbaPro: 对GPU编程的库，提供许多的数值计算库，GPU加速库
1.参考书籍：《Python并行编程手册》
标签：上下文 其它 结构 必须 set 方式 共享 port api
原文地址：https://www.cnblogs.com/maplesnow/p/12044372.html
