【GPU加速系列】PyCUDA(一):上手简单操作
PyCUDA 可以通过 Python 访问 Navidia 的 CUDA 并行计算 API。
具体介绍和安装可以参考 PyCUDA 官网文档和 pycuda PyPI。
本文涵盖的内容有:
- 通过 PyCUDA 查询 GPU 信息。
- NumPy array 和 gpuarray 之间的相互转换。
- 使用 gpuarray 进行基本的运算。
- 使用 ElementwiseKernel 进行按元素的运算。
- 使用 InclusiveScanKernel 和 ReductionKernel 的 reduce 操作。
本文示例在 GPU 环境下,使用 Jupyter Notebook 导入了以下包:
1 import sys 2 from time import time 3 from functools import reduce 4 5 import numpy as np 6 import pandas as pd 7 import matplotlib 8 from matplotlib import pyplot as plt 9 from IPython.core.interactiveshell import InteractiveShell 10 11 import pycuda 12 import pycuda.autoinit 13 import pycuda.driver as drv 14 from pycuda import gpuarray 15 from pycuda.elementwise import ElementwiseKernel 16 from pycuda.scan import InclusiveScanKernel 17 from pycuda.reduction import ReductionKernel 18 19 InteractiveShell.ast_node_interactivity = "all" 20 print(f'The version of PyCUDA: {pycuda.VERSION}') 21 print(f'The version of Python: {sys.version}')
输出:
The version of PyCUDA: (2019, 1, 2) The version of Python: 3.6.6 |Anaconda, Inc.| (default, Oct 9 2018, 12:34:16) [GCC 7.3.0]
查询 GPU 信息
GPU 查询是一个非常基本的操作,比较常用的重要信息有 GPU 设备名、GPU 显存、核心数量等。
定义函数:
1 def query_device(): 2 drv.init() 3 print('CUDA device query (PyCUDA version) \n') 4 print(f'Detected {drv.Device.count()} CUDA Capable device(s) \n') 5 for i in range(drv.Device.count()): 6 7 gpu_device = drv.Device(i) 8 print(f'Device {i}: {gpu_device.name()}') 9 compute_capability = float( '%d.%d' % gpu_device.compute_capability() ) 10 print(f'\t Compute Capability: {compute_capability}') 11 print(f'\t Total Memory: {gpu_device.total_memory()//(1024**2)} megabytes') 12 13 # The following will give us all remaining device attributes as seen 14 # in the original deviceQuery. 15 # We set up a dictionary as such so that we can easily index 16 # the values using a string descriptor. 17 18 device_attributes_tuples = gpu_device.get_attributes().items() 19 device_attributes = {} 20 21 for k, v in device_attributes_tuples: 22 device_attributes[str(k)] = v 23 24 num_mp = device_attributes['MULTIPROCESSOR_COUNT'] 25 26 # Cores per multiprocessor is not reported by the GPU! 27 # We must use a lookup table based on compute capability. 28 # See the following: 29 # http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities 30 31 cuda_cores_per_mp = { 5.0 : 128, 5.1 : 128, 5.2 : 128, 6.0 : 64, 6.1 : 128, 6.2 : 128}[compute_capability] 32 33 print(f'\t ({num_mp}) Multiprocessors, ({cuda_cores_per_mp}) CUDA Cores / Multiprocessor: {num_mp*cuda_cores_per_mp} CUDA Cores') 34 35 device_attributes.pop('MULTIPROCESSOR_COUNT') 36 37 for k in device_attributes.keys(): 38 print(f'\t {k}: {device_attributes[k]}')
执行 GPU 查询操作:
CUDA device query (PyCUDA version) Detected 1 CUDA Capable device(s) Device 0: Tesla P100-PCIE-16GBCompute Capability: 6.0Total Memory: 16280 megabytes(56) Multiprocessors, (64) CUDA Cores / Multiprocessor: 3584 CUDA CoresASYNC_ENGINE_COUNT: 2CAN_MAP_HOST_MEMORY: 1CLOCK_RATE: 1328500COMPUTE_CAPABILITY_MAJOR: 6COMPUTE_CAPABILITY_MINOR: 0COMPUTE_MODE: DEFAULTCONCURRENT_KERNELS: 1ECC_ENABLED: 1GLOBAL_L1_CACHE_SUPPORTED: 1GLOBAL_MEMORY_BUS_WIDTH: 4096GPU_OVERLAP: 1INTEGRATED: 0KERNEL_EXEC_TIMEOUT: 0L2_CACHE_SIZE: 4194304LOCAL_L1_CACHE_SUPPORTED: 1MANAGED_MEMORY: 1MAXIMUM_SURFACE1D_LAYERED_LAYERS: 2048MAXIMUM_SURFACE1D_LAYERED_WIDTH: 32768MAXIMUM_SURFACE1D_WIDTH: 32768MAXIMUM_SURFACE2D_HEIGHT: 65536MAXIMUM_SURFACE2D_LAYERED_HEIGHT: 32768MAXIMUM_SURFACE2D_LAYERED_LAYERS: 2048MAXIMUM_SURFACE2D_LAYERED_WIDTH: 32768MAXIMUM_SURFACE2D_WIDTH: 131072MAXIMUM_SURFACE3D_DEPTH: 16384MAXIMUM_SURFACE3D_HEIGHT: 16384MAXIMUM_SURFACE3D_WIDTH: 16384MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS: 2046MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH: 32768MAXIMUM_SURFACECUBEMAP_WIDTH: 32768MAXIMUM_TEXTURE1D_LAYERED_LAYERS: 2048MAXIMUM_TEXTURE1D_LAYERED_WIDTH: 32768MAXIMUM_TEXTURE1D_LINEAR_WIDTH: 134217728MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH: 16384MAXIMUM_TEXTURE1D_WIDTH: 131072MAXIMUM_TEXTURE2D_ARRAY_HEIGHT: 32768MAXIMUM_TEXTURE2D_ARRAY_NUMSLICES: 2048MAXIMUM_TEXTURE2D_ARRAY_WIDTH: 32768MAXIMUM_TEXTURE2D_GATHER_HEIGHT: 32768MAXIMUM_TEXTURE2D_GATHER_WIDTH: 32768MAXIMUM_TEXTURE2D_HEIGHT: 65536MAXIMUM_TEXTURE2D_LINEAR_HEIGHT: 65000MAXIMUM_TEXTURE2D_LINEAR_PITCH: 2097120MAXIMUM_TEXTURE2D_LINEAR_WIDTH: 131072MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT: 32768MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH: 32768MAXIMUM_TEXTURE2D_WIDTH: 131072MAXIMUM_TEXTURE3D_DEPTH: 16384MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE: 32768MAXIMUM_TEXTURE3D_HEIGHT: 16384MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE: 8192MAXIMUM_TEXTURE3D_WIDTH: 16384MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE: 8192MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS: 2046MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH: 32768MAXIMUM_TEXTURECUBEMAP_WIDTH: 32768MAX_BLOCK_DIM_X: 1024MAX_BLOCK_DIM_Y: 1024MAX_BLOCK_DIM_Z: 64MAX_GRID_DIM_X: 2147483647MAX_GRID_DIM_Y: 65535MAX_GRID_DIM_Z: 65535MAX_PITCH: 2147483647MAX_REGISTERS_PER_BLOCK: 65536MAX_REGISTERS_PER_MULTIPROCESSOR: 65536MAX_SHARED_MEMORY_PER_BLOCK: 49152MAX_SHARED_MEMORY_PER_MULTIPROCESSOR: 65536MAX_THREADS_PER_BLOCK: 1024MAX_THREADS_PER_MULTIPROCESSOR: 2048MEMORY_CLOCK_RATE: 715000MULTI_GPU_BOARD: 0MULTI_GPU_BOARD_GROUP_ID: 0PCI_BUS_ID: 0PCI_DEVICE_ID: 4PCI_DOMAIN_ID: 0STREAM_PRIORITIES_SUPPORTED: 1SURFACE_ALIGNMENT: 512TCC_DRIVER: 0TEXTURE_ALIGNMENT: 512TEXTURE_PITCH_ALIGNMENT: 32TOTAL_CONSTANT_MEMORY: 65536UNIFIED_ADDRESSING: 1WARP_SIZE: 32
在这里,我们发现了有一个 GPU 设备 Tesla P100-PCIE-16GB,其显存为 16G,核心数目为 3584 个。
NumPy array 和 gpuarray 之间的相互转换
GPU 有自己的显存,这区别于主机上的内存,这又称为设备内存(device memory)。
NumPy array 运行在 CPU 环境(主机端),而 gpuarray 运行在 GPU 环境(设备端),两者常常需要相互转换,即 CPU 数据和 GPU 数据之间的传输转换。
1 host_data = np.array([1, 2, 3, 4, 5], dtype=np.float32) 2 device_data = gpuarray.to_gpu(host_data) 3 device_data_x2 = 2 * device_data 4 host_data_x2 = device_data_x2.get() 5 print(host_data_x2)
其输出:
[ 2. 4. 6. 8. 10.]
进行转换的时候应该尽可能通过 dtype 指定类型,以避免不必要的性能损失。
gpuarray 的基本运算
按元素运算是天生的可并行计算的操作类型,在进行这种运算时 gpuarray 会自动利用多核进行并行计算。
1 x_host = np.array([1, 2, 3], dtype=np.float32) 2 y_host = np.array([1, 1, 1], dtype=np.float32) 3 z_host = np.array([2, 2, 2], dtype=np.float32) 4 x_device = gpuarray.to_gpu(x_host) 5 y_device = gpuarray.to_gpu(y_host) 6 z_device = gpuarray.to_gpu(z_host) 7 8 x_host + y_host 9 (x_device + y_device).get() 10 11 x_host ** z_host 12 (x_device ** z_device).get() 13 14 x_host / x_host 15 (x_device / x_device).get() 16 17 z_host - x_host 18 (z_device - x_device).get() 19 20 z_host / 2 21 (z_device / 2).get() 22 23 x_host - 1 24 (x_device - 1).get()
输出:
array([2., 3., 4.], dtype=float32) array([2., 3., 4.], dtype=float32) array([1., 4., 9.], dtype=float32) array([1., 4., 9.], dtype=float32) array([1., 1., 1.], dtype=float32) array([1., 1., 1.], dtype=float32) array([ 1., 0., -1.], dtype=float32) array([ 1., 0., -1.], dtype=float32) array([1., 1., 1.], dtype=float32) array([1., 1., 1.], dtype=float32) array([0., 1., 2.], dtype=float32) array([0., 1., 2.], dtype=float32)
性能比较
1 def simple_speed_test(): 2 host_data = np.float32(np.random.random(50000000)) 3 4 t1 = time() 5 host_data_2x = host_data * np.float32(2) 6 t2 = time() 7 8 print(f'total time to compute on CPU: {t2 - t1}') 9 10 device_data = gpuarray.to_gpu(host_data) 11 12 t1 = time() 13 device_data_2x = device_data * np.float32(2) 14 t2 = time() 15 16 from_device = device_data_2x.get() 17 18 print(f'total time to compute on GPU: {t2 - t1}') 19 print(f'Is the host computation the same as the GPU computation? : {np.allclose(from_device, host_data_2x)}') 20 21 simple_speed_test()
如果是第一次执行会输出类似:
total time to compute on CPU: 0.14141535758972168 total time to compute on GPU: 2.010883092880249 Is the host computation the same as the GPU computation? : True
而后面再继续执行几次,会有类似的输出:
total time to compute on CPU: 0.1373155117034912 total time to compute on GPU: 0.0006959438323974609 Is the host computation the same as the GPU computation? : True
这是因为在 PyCUDA 中,通常会在程序第一次运行过程中,nvcc 编译器会对 GPU 代码进行编译,然后由 PyCUDA 进行调用。这个编译时间就是额外的性能损耗。
ElementwiseKernel:按元素运算
我们先看一下 Python 的内置函数 map。
第一个参数 function 以参数序列中的每一个元素调用 function 函数,返回包含每次 function 函数返回值的迭代器(Python2 中 map 输出的是列表),我们用 list() 把迭代器转换为列表观察结果。
list(map(lambda x: x + 10, [1, 2, 3, 4, 5]))
输出:
[11, 12, 13, 14, 15]
ElementWiseKernel 非常类似于 map 函数。
ElementwiseKernel 函数可以自定义按元素运算的内核。使用时需要嵌入 CUDA C 的代码。
内核(kernel)在这里可以简单理解为 CUDA 直接运行在 GPU 的函数。
看代码:
1 gpu_2x_ker = ElementwiseKernel( 2 "float *in, float *out", 3 "out[i] = 2 * in[i];", 4 "gpu_2x_ker" 5 ) 6 7 def elementwise_kernel_example(): 8 host_data = np.float32(np.random.random(50000000)) 9 t1 = time() 10 host_data_2x = host_data * np.float32(2) 11 t2 = time() 12 print(f'total time to compute on CPU: {t2 - t1}') 13 14 device_data = gpuarray.to_gpu(host_data) 15 # allocate memory for output 16 device_data_2x = gpuarray.empty_like(device_data) 17 18 t1 = time() 19 gpu_2x_ker(device_data, device_data_2x) 20 t2 = time() 21 from_device = device_data_2x.get() 22 print(f'total time to compute on GPU: {t2 - t1}') 23 print(f'Is the host computation the same as the GPU computation? : {np.allclose(from_device, host_data_2x)}') 24 25 elementwise_kernel_example() 26 elementwise_kernel_example() 27 elementwise_kernel_example() 28 elementwise_kernel_example() 29 elementwise_kernel_example()
输出:
total time to compute on CPU: 0.13545799255371094 total time to compute on GPU: 0.4059629440307617 Is the host computation the same as the GPU computation? : True total time to compute on CPU: 0.13948774337768555 total time to compute on GPU: 0.0001266002655029297 Is the host computation the same as the GPU computation? : True total time to compute on CPU: 0.1357274055480957 total time to compute on GPU: 0.0001552104949951172 Is the host computation the same as the GPU computation? : True total time to compute on CPU: 0.13451647758483887 total time to compute on GPU: 0.0001761913299560547 Is the host computation the same as the GPU computation? : True total time to compute on CPU: 0.1362597942352295 total time to compute on GPU: 0.00011849403381347656 Is the host computation the same as the GPU computation? : True
同样我们发现在第一次运行时,出现了 nvcc 编译产生的性能损耗。
ElementwiseKernel 的参数:
class
pycuda.elementwise.
ElementwiseKernel
(arguments, operation, name="kernel", keep=False, options=[], preamble="")
- arguments:该内核定义的传参。
- operation:该内核定义的内嵌 CUDA C 代码。
- name:定义的内核名称。
gpuarray.empty_like 用于分配与 device_data 相同形状和类型的内存空间。
InclusiveScanKernel 和 ReductionKernel 的 reduce 操作
我们先看一下 Python 标准包 functools 中的 reduce 函数。
reduce(lambda x, y : x + y, [1, 2, 3, 4])
输出:
10
与 map 函数不同,reduce 执行迭代的二元运算,只输出一个单值。
我们将使用 InclusiveScan 和 ReductionKernel 来实现类似于 reduce 的操作。
InclusiveScanKernel
InclusiveScanKernel 类似于 reduce,因为它并非输出单值,输出与输入形状相同。
计算求和的操作,输出是一个累加的序列:
1 seq = np.array([1, 2, 3, 4], dtype=np.int32) 2 seq_gpu = gpuarray.to_gpu(seq) 3 sum_gpu = InclusiveScanKernel(np.int32, "a+b") 4 print(sum_gpu(seq_gpu).get()) 5 print(np.cumsum(seq))
输出:
[ 1 3 6 10] [ 1 3 6 10]
查找最大值(最大值向后冒泡):
1 seq = np.array([1,100,-3,-10000, 4, 10000, 66, 14, 21], dtype=np.int32) 2 seq_gpu = gpuarray.to_gpu(seq) 3 max_gpu = InclusiveScanKernel(np.int32, "a > b ? a : b") 4 seq_max_bubble = max_gpu(seq_gpu) 5 print(seq_max_bubble) 6 print(seq_max_bubble.get()[-1]) 7 print(np.max(seq))
输出:
[ 1 100 100 100 100 10000 10000 10000 10000] 10000 10000
对于 a > b ? a : b ,我们可以想象是做从前往后做一个遍历(实际是并行的),而对于每个当前元素 cur,都和前一个元素做比较,把最大值赋值给 cur。
这样,最大值就好像“冒泡”一样往后移动,最终取最后一个元素即可。
ReductionKernel
实际上,ReductionKernel 就像是执行 ElementWiseKernel 后再执行一个并行扫描内核。
一个计算两向量内积的例子:
1 a_host = np.array([1, 2, 3], dtype=np.float32) 2 b_host = np.array([4, 5, 6], dtype=np.float32) 3 print(a_host.dot(b_host)) 4 5 dot_prod = ReductionKernel(np.float32, neutral="0", reduce_expr="a+b", 6 map_expr="x[i]*y[i]", arguments="float *x, float *y") 7 a_device = gpuarray.to_gpu(a_host) 8 b_device = gpuarray.to_gpu(b_host) 9 print(dot_prod(a_device, b_device).get())
32.0 32.0
首先对两向量的每个元素进行 map_expr 的计算,其结果再进行 reduce_expr 的计算(neutral 表示初始值),最终得到两向量的内积。
好了,到此为止,就是初识 PyCUDA 的一些操作。
原文作者:雨先生
原文链接:https://www.cnblogs.com/noluye/p/11465389.html
许可协议:知识共享署名-非商业性使用 4.0 国际许可协议
参考
- PyCUDA 官网文档
- pycuda PyPI
- 《Hands-On GPU Programming with Python and CUDA》by Dr. Brian Tuomanen
- CUDA系列学习(五)GPU基础算法: Reduce, Scan, Histogram
转载于:https://www.cnblogs.com/noluye/p/11517489.html
【GPU加速系列】PyCUDA(一):上手简单操作相关推荐
- 【Python - GPU】基于Python的GPU加速并行计算 -- pyCUDA
Python实现的CUDA – pyCUDA Nvidia的CUDA 架构为我们提供了一种便捷的方式来直接操纵GPU 并进行编程,但是基于 C语言的CUDA实现较为复杂,开发周期较长.而python ...
- python gpu并行计算_【Python - GPU】基于Python的GPU加速并行计算 -- pyCUDA
Python实现的CUDA – pyCUDA Nvidia的CUDA 架构为我们提供了一种便捷的方式来直接操纵GPU 并进行编程,但是基于 C语言的CUDA实现较为复杂,开发周期较长.而python ...
- 从入门到入土:Python爬虫学习|Selenium自动化模块学习|简单入门|轻松上手|自动操作浏览器进行处理|chrome|PART01
此博客仅用于记录个人学习进度,学识浅薄,若有错误观点欢迎评论区指出.欢迎各位前来交流.(部分材料来源网络,若有侵权,立即删除) 本人博客所有文章纯属学习之用,不涉及商业利益.不合适引用,自当删除! 若 ...
- Mac Apple Silicon M1/M2 homebrew miniforge conda pytorch yolov5深度学习环境搭建并简单测试MPS GPU加速
目录 开始安装 零,获取代理 一,配置代理 配置zsh走代理 配置git走代理 二,安装homebrew 三,安装miniforge 四,创建conda环境 五,安装pytorch 六,运行yolov ...
- python gpu编程_Python笔记_第四篇_高阶编程_进程、线程、协程_5.GPU加速
Numba:高性能计算的高生产率 在这篇文章中,笔者将向你介绍一个来自Anaconda的Python编译器Numba,它可以在CUDA-capable GPU或多核cpu上编译Python代码.Pyt ...
- 利用gpu加速神经网络算法,外接gpu 训练神经网络
神经网络做图像分类一定要用到gpu吗? GPU最大的价值一直是"accelerating"(加速),GPU不是取代CPU,而是利用GPU的并行计算架构,来将并行计算的负载放到GPU ...
- 56 Marvin: 一个支持GPU加速、且不依赖其他库(除cuda和cudnn)的轻量化多维深度学习(deep learning)框架介绍...
0 引言 Marvin是普林斯顿视觉实验室(PrincetonVision)于2015年提出的轻量化GPU加速的多维深度学习网络框架.该框架采用纯c/c++编写,除了cuda和cudnn以外,不依赖其 ...
- GPU加速原理浅析及代码实现
GPU加速原理浅析及代码实现 一.CUDA简介 二.GPU架构特点 三.CUDA线程模型 四.CUDA内存模型 五.CUDA编程规范 **第一个要掌握的编程要点**:**我们怎么写一个能在GPU跑的程 ...
- Python GPU加速
Numba:高性能计算的高生产率 在这篇文章中,笔者将向你介绍一个来自Anaconda的Python编译器Numba,它可以在CUDA-capable GPU或多核cpu上编译Python代码.Pyt ...
最新文章
- 标准访问控制列表和扩展的访问控制列表有什么区别??
- http://hi.baidu.com/grayworm/blog/item/c76c4046307d380a6a63e56a.html
- 机器学习、深度学习所需掌握的数学知识全都在这里了!
- C/C++中各种类型int、long、double、char表示范围(最大最小值)
- 链接器怎样使用静态库来解决符号引用
- TensorFlow RNN tutorial解读
- MyBatis中使用#{}和${}的区别
- 女朋友当众甩了我一巴掌,我扑上去......
- php上传文件表单,php中关于普通表单多文件上传的处理方法
- 学会 配置文件+反射,走遍全球都不怕.
- Windows11配置Java开发环境
- NB-IoT在无线烟感监控系统中的优势
- LeCun、Bengio、Hinton三巨头曾合体,Nature发文综述深度学习(论文精华)
- ajax jsonp跨域访问,jquery ajax怎么使用jsonp跨域访问
- Codeforces 1013 A. Piles With Stones
- nginx “504 Gateway Time-out”错误
- Android逆向-实战so分析-某洲_v3.5.8_unidbg学习
- 什么是下一代防火墙NGFW(Next Generation Firewall)?
- 格式化后如何修复损坏的JPEG文件?
- 【备忘】AAD Intune维护
热门文章
- Codeforces Zepto Code Rush 2014 -C - Dungeons and Candies
- android使用webview加载flash文件
- 难度炸裂!DeepChange:一个新的超大规模的换衣行人再识别数据集
- CVPR2021 双图层实例分割,大幅提升遮挡处理性能
- 胶囊网络与计算机视觉教程 @CVPR 2019
- ICCV 2019 | 旷视提出MetaPruning:基于元学习和AutoML的模型压缩新方法
- 推荐!京东开源姿态跟踪新框架LightTrack!
- debian10树莓派4安装mysql_树莓派 4 发布!新 OS 基于 Debian 10 Buster
- 想转行,是要入坑Python还是Java?这问题还用问?
- 定了!2022考研时间公布!