码迷,mamicode.com
首页 > 编程语言 > 详细

[spring 并行5]GPU

时间:2019-12-15 16:49:08      阅读:92      评论:0      收藏:0      [点我收藏+]

标签:上下文   其它   结构   必须   set   方式   共享   port   api   

GPU篇 1

准备

  1. 需要有支持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
  2. 需要安装pycuda(linux安装:apt install python3-pycuda)

基本使用

1.底层操作

  1. 准备工作: sudo apt install python3-pycuda
  2. 在cpu使用的内存中创建矩阵
  3. 将矩阵从cpu使用的内存移动到gpu的显存中
  4. 编辑c语言代码,让gpu计算
  5. 将矩阵从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)

注1import pycuda.autoinit 语句自动根据GPU可用性和数量选择要使用的GPU,这将创建一个接下来的代码运行中所需的GPU上下文(只需导入即可完成)

: astype(numpy.float32):将矩阵中的项转换为单精度模式,因为许多Nvidia显卡只支持单精度

注3:在调用gpu的c函数时,通过block参数设定,分配线程的方式,(5, 5, 1)是对应这gpu的(x, y, x)的分配
在c函数中,threadIdx是一个结构体,它有三个字段xyz,每个线程中的这个变量都不同(结合线程层级理解),故用此索引数组,由于动态分配的gpu内存是一维数组,所以需要在c函数内,使用threadIdx.y乘以矩阵每行的元素个数转换
更多详情:自行搜索cuda的线程区块划分

注4:gpu执行的c函数中 __global__ 关键字表示该函数是一个内核函数,必须从主机上调用才能在gpu设备上生成线程层级
涉及到pycuda内存,为了最大限度地利用可用资源,在支持CUDA的GPU显卡中,有4类内存:
寄存器(registers):每个线程将被分配一个寄存器,每个线程只能访问自身的寄存器,即使同属于一个线程区块
共享存储器(shared memory):在共享存储器中,每个线程区块都有一个其内部线程共享的内存,这部分内存速度极快
常数存储器(constant memory):一个网格中的所有线程一直都可以访问这部分内存,但只能在读取时访问。常数存储器中的数据在应用持续期间一直存在
全局存储器(global memory),:所有网格中的线程(也包含所有内核)都可访问全局存储器

更多详情:自行搜索PyCUDA内存模型

2.python封装控制

用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并行编程手册》

[spring 并行5]GPU

标签:上下文   其它   结构   必须   set   方式   共享   port   api   

原文地址:https://www.cnblogs.com/maplesnow/p/12044372.html

(0)
(0)
   
举报
评论 一句话评论(0
登录后才能评论!
© 2014 mamicode.com 版权所有  联系我们:gaon5@hotmail.com
迷上了代码!