【GPU加速系列】PyCUDA(一):上手简单操作

PyCUDA 可以通过 Python 访问 Navidia 的 CUDA 并行计算 API。

具体介绍和安装可以参考 PyCUDA 官网文档和 pycuda PyPI

本文涵盖的内容有:

  1. 通过 PyCUDA 查询 GPU 信息。
  2. NumPy array 和 gpuarray 之间的相互转换。
  3. 使用 gpuarray 进行基本的运算。
  4. 使用 ElementwiseKernel 进行按元素的运算。
  5. 使用 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-16GB
	 Compute Capability: 6.0
	 Total Memory: 16280 megabytes
	 (56) Multiprocessors, (64) CUDA Cores / Multiprocessor: 3584 CUDA Cores
	 ASYNC_ENGINE_COUNT: 2
	 CAN_MAP_HOST_MEMORY: 1
	 CLOCK_RATE: 1328500
	 COMPUTE_CAPABILITY_MAJOR: 6
	 COMPUTE_CAPABILITY_MINOR: 0
	 COMPUTE_MODE: DEFAULT
	 CONCURRENT_KERNELS: 1
	 ECC_ENABLED: 1
	 GLOBAL_L1_CACHE_SUPPORTED: 1
	 GLOBAL_MEMORY_BUS_WIDTH: 4096
	 GPU_OVERLAP: 1
	 INTEGRATED: 0
	 KERNEL_EXEC_TIMEOUT: 0
	 L2_CACHE_SIZE: 4194304
	 LOCAL_L1_CACHE_SUPPORTED: 1
	 MANAGED_MEMORY: 1
	 MAXIMUM_SURFACE1D_LAYERED_LAYERS: 2048
	 MAXIMUM_SURFACE1D_LAYERED_WIDTH: 32768
	 MAXIMUM_SURFACE1D_WIDTH: 32768
	 MAXIMUM_SURFACE2D_HEIGHT: 65536
	 MAXIMUM_SURFACE2D_LAYERED_HEIGHT: 32768
	 MAXIMUM_SURFACE2D_LAYERED_LAYERS: 2048
	 MAXIMUM_SURFACE2D_LAYERED_WIDTH: 32768
	 MAXIMUM_SURFACE2D_WIDTH: 131072
	 MAXIMUM_SURFACE3D_DEPTH: 16384
	 MAXIMUM_SURFACE3D_HEIGHT: 16384
	 MAXIMUM_SURFACE3D_WIDTH: 16384
	 MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS: 2046
	 MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH: 32768
	 MAXIMUM_SURFACECUBEMAP_WIDTH: 32768
	 MAXIMUM_TEXTURE1D_LAYERED_LAYERS: 2048
	 MAXIMUM_TEXTURE1D_LAYERED_WIDTH: 32768
	 MAXIMUM_TEXTURE1D_LINEAR_WIDTH: 134217728
	 MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH: 16384
	 MAXIMUM_TEXTURE1D_WIDTH: 131072
	 MAXIMUM_TEXTURE2D_ARRAY_HEIGHT: 32768
	 MAXIMUM_TEXTURE2D_ARRAY_NUMSLICES: 2048
	 MAXIMUM_TEXTURE2D_ARRAY_WIDTH: 32768
	 MAXIMUM_TEXTURE2D_GATHER_HEIGHT: 32768
	 MAXIMUM_TEXTURE2D_GATHER_WIDTH: 32768
	 MAXIMUM_TEXTURE2D_HEIGHT: 65536
	 MAXIMUM_TEXTURE2D_LINEAR_HEIGHT: 65000
	 MAXIMUM_TEXTURE2D_LINEAR_PITCH: 2097120
	 MAXIMUM_TEXTURE2D_LINEAR_WIDTH: 131072
	 MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT: 32768
	 MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH: 32768
	 MAXIMUM_TEXTURE2D_WIDTH: 131072
	 MAXIMUM_TEXTURE3D_DEPTH: 16384
	 MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE: 32768
	 MAXIMUM_TEXTURE3D_HEIGHT: 16384
	 MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE: 8192
	 MAXIMUM_TEXTURE3D_WIDTH: 16384
	 MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE: 8192
	 MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS: 2046
	 MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH: 32768
	 MAXIMUM_TEXTURECUBEMAP_WIDTH: 32768
	 MAX_BLOCK_DIM_X: 1024
	 MAX_BLOCK_DIM_Y: 1024
	 MAX_BLOCK_DIM_Z: 64
	 MAX_GRID_DIM_X: 2147483647
	 MAX_GRID_DIM_Y: 65535
	 MAX_GRID_DIM_Z: 65535
	 MAX_PITCH: 2147483647
	 MAX_REGISTERS_PER_BLOCK: 65536
	 MAX_REGISTERS_PER_MULTIPROCESSOR: 65536
	 MAX_SHARED_MEMORY_PER_BLOCK: 49152
	 MAX_SHARED_MEMORY_PER_MULTIPROCESSOR: 65536
	 MAX_THREADS_PER_BLOCK: 1024
	 MAX_THREADS_PER_MULTIPROCESSOR: 2048
	 MEMORY_CLOCK_RATE: 715000
	 MULTI_GPU_BOARD: 0
	 MULTI_GPU_BOARD_GROUP_ID: 0
	 PCI_BUS_ID: 0
	 PCI_DEVICE_ID: 4
	 PCI_DOMAIN_ID: 0
	 STREAM_PRIORITIES_SUPPORTED: 1
	 SURFACE_ALIGNMENT: 512
	 TCC_DRIVER: 0
	 TEXTURE_ALIGNMENT: 512
	 TEXTURE_PITCH_ALIGNMENT: 32
	 TOTAL_CONSTANT_MEMORY: 65536
	 UNIFIED_ADDRESSING: 1
	 WARP_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(argumentsoperationname="kernel"keep=Falseoptions=[]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 执行迭代的二元运算,只输出一个单值

我们将使用 InclusiveScanReductionKernel 来实现类似于 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 国际许可协议

参考

原文地址:https://www.cnblogs.com/noluye/p/11517489.html

时间: 2024-10-11 04:24:12

【GPU加速系列】PyCUDA(一):上手简单操作的相关文章

Python基于pyCUDA实现GPU加速并行计算功能入门教程

https://www.jb51.net/article/142212.htm 这篇文章主要介绍了Python基于pyCUDA实现GPU加速并行计算功能,结合实例形式分析了Python使用pyCUDA进行GPU加速并行计算的原理与相关实现操作技巧,需要的朋友可以参考下 目录 pyCUDA特点 pyCUDA的工作流程 调用基本例子 具体内容 本文实例讲述了Python基于pyCUDA实现GPU加速并行计算功能.分享给大家供大家参考,具体如下: Nvidia的CUDA 架构为我们提供了一种便捷的方式

ZooKeeper系列3:ZooKeeper命令、命令行工具及简单操作

问题导读1.ZooKeeper包含哪些常用命令?2.通过什么命令可以列出服务器 watch 的详细信息?3.ZooKeeper包含哪些操作?4.ZooKeeper如何创建zookeeper? 常用命令 ZooKeeper 支持某些特定的四字命令字母与其的交互.它们大多是查询命令,用来获取 ZooKeeper 服务的当前状态及相关信息.用户在客户端可以通过 telnet 或 nc 向 ZooKeeper 提交相应的命令. ZooKeeper 常用四字命令见下表 1 所示: 表 1 : ZooKee

Python GPU加速

https://blog.csdn.net/weixin_41923961/article/details/83687809 Numba:高性能计算的高生产率 在这篇文章中,笔者将向你介绍一个来自Anaconda的Python编译器Numba,它可以在CUDA-capable GPU或多核cpu上编译Python代码.Python通常不是一种编译语言,你可能想知道为什么要使用Python编译器.答案当然是:运行本地编译的代码要比运行动态的.解译的代码快很多倍.Numba允许你为Python函数指

windows 10 64bit下安装Tensorflow+Keras+VS2015+CUDA8.0 GPU加速

原文地址:http://www.jianshu.com/p/c245d46d43f0 写在前面的话 2016年11月29日,Google Brain 工程师团队宣布在 TensorFlow 0.12 中加入初步的 Windows 支持.但是目前只支持64位,而且Python版本为3.5版本,需要CUDA 8.0 .之前Tensorflow对windows的支持并不好,导致如果需要使用它,需要转移到Linux平台,或者使用Cygwin什么的,总之挺麻烦,现在好了.麻烦事google帮我们解决了.感

Javascript如何实现GPU加速?

一.什么是Javascript实现GPU加速? CPU与GPU设计目标不同,导致它们之间内部结构差异很大.CPU需要应对通用场景,内部结构非常复杂.而GPU往往面向数据类型统一,且相互无依赖的计算.所以,我们在Web上实现3D场景时,通常使用WebGL利用GPU运算(大量顶点).但是,如果只是通用的计算场景呢?比如处理图片中大量像素信息,我们有办法使用GPU资源吗?这正是本文要讲的,GPU通用计算,简称GPGPU. 二.实例演示:色块识别. 如下图所示,我们识别图片中彩虹糖色块,给糖果添加表情.

编译GDAL支持OpenCL使用GPU加速

前言 GDAL库中提供的gdalwarp支持各种高性能的图像重采样算法,图像重采样算法广泛应用于图像校正,重投影,裁切,镶嵌等算法中,而且对于这些算法来说,计算坐标变换的运算量是相当少的,绝大部分运算量都在图像的重采样算法中,尤其是三次卷积采样以及更高级的重采样算法来说,运算量会成倍的增加,所以提升这些算法的处理效率优先是提高重采样的效率.由于GPU的多核心使得目前对于GPU的并行处理非常热,同时也能大幅度的提升处理速度.基于上述原因,GDALWARP也提供了基于OPENCL的GPU加速,之前在

GPU 加速NLP任务(Theano+CUDA)

之前学习了CNN的相关知识,提到Yoon Kim(2014)的论文,利用CNN进行文本分类,虽然该CNN网络结构简单效果可观,但论文没有给出具体训练时间,这便值得进一步探讨. Yoon Kim代码:https://github.com/yoonkim/CNN_sentence 利用作者提供的源码进行学习,在本人机子上训练时,做一次CV的平均训练时间如下,纵坐标为min/CV(供参考): 机子配置:Intel(R) Core(TM) i3-4150 CPU @ 3.50GHz, 32G,x64 显

小程序开发工具教程汇总,亲测无需代码,简单操作,可二次开发!

制作小程序要简单点,开发的方式简单点.所以!不得不扒一扒这款开发工具某即速应用,无需代码,简单操作,还可套用模板快速生成小程序代码包,支持二次开发.当然市面上还有很多款小程序编辑器,但大部分还是那种开发接单的,水太深了,交不起学费,还是自己来琢磨吧!大家有什么好的小程序工具类可以分享给我! 以下将我网上参与学习培训总结的一些资料文档分享出来,供大家参考学习.特别是不懂代码的朋友,这款开发工具还是可以试试的!容易上手! 当然首先,我们还是要先下载:最新版的微信开发者工具(方便进入代码二次编辑) 接

Apache Derby-02通过IJ简单操作DERBY

上回说到了Derby的历史以及需要准备的环境,这章将为大家介绍Apache Derby的简单操作 1.配置Derby环境 1.1去官网下载Derby_BIN并解压在文件夹中 http://mirrors.cnnic.cn/apache//db/derby/db-derby-10.12.1.1/db-derby-10.12.1.1-bin.zip 1.2如果需要在本地使用DERBY需要在系统变量中设置变量,derby是基于JAVA环境的所以在之前需要JAVA环境变量的配置 1.3 在系统变量创建D