python gpu加速函数_【GPU加速系列】PyCUDA（一）：上手简单操作

最新推荐文章于 2024-04-12 14:12:47 发布

weixin_39803207

最新推荐文章于 2024-04-12 14:12:47 发布

阅读量870

点赞数

文章标签： python gpu加速函数

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

本文涵盖的内容有：

通过 PyCUDA 查询 GPU 信息。

NumPy array 和 gpuarray 之间的相互转换。

使用 gpuarray 进行基本的运算。

使用 ElementwiseKernel 进行按元素的运算。

使用 InclusiveScanKernel 和 ReductionKernel 的 reduce 操作。

本文示例在 GPU 环境下，使用 Jupyter Notebook 导入了以下包：

1 importsys2 from time importtime3 from functools importreduce4

5 importnumpy as np6 importpandas as pd7 importmatplotlib8 from matplotlib importpyplot as plt9 from IPython.core.interactiveshell importInteractiveShell10

11 importpycuda12 importpycuda.autoinit13 importpycuda.driver as drv14 from pycuda importgpuarray15 from pycuda.elementwise importElementwiseKernel16 from pycuda.scan importInclusiveScanKernel17 from pycuda.reduction importReductionKernel18

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 defquery_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 inrange(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.

18 device_attributes_tuples =gpu_device.get_attributes().items()19 device_attributes ={}20

21 for k, v indevice_attributes_tuples:22 device_attributes[str(k)] =v23

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

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 indevice_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_data4 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_host9 (x_device +y_device).get()10

11 x_host **z_host12 (x_device **z_device).get()13

14 x_host /x_host15 (x_device /x_device).get()16

17 z_host -x_host18 (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([1., 4., 9.], dtype=float32)

array([1., 1., 1.], dtype=float32)

array([ 1., 0., -1.], dtype=float32)

array([1., 1., 1.], dtype=float32)

array([0., 1., 2.], dtype=float32)

性能比较

1 defsimple_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 defelementwise_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 的参数：

classpycuda.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])

输出：

与 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 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

对于 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

首先对两向量的每个元素进行 map_expr 的计算，其结果再进行 reduce_expr 的计算(neutral 表示初始值)，最终得到两向量的内积。

好了，到此为止，就是初识 PyCUDA 的一些操作。

参考

《Hands-On GPU Programming with Python and CUDA》by Dr. Brian Tuomanen

weixin_39803207

关注

0
点赞
踩
0

收藏

觉得还不错? 一键收藏
0
评论
python gpu加速函数_【GPU加速系列】PyCUDA（一）：上手简单操作

PyCUDA 可以通过 Python 访问 NVIDIA 的 CUDA 并行计算 API。本文涵盖的内容有：通过 PyCUDA 查询 GPU 信息。NumPy array 和 gpuarray 之间的相互转换。使用 gpuarray 进行基本的运算。使用 ElementwiseKernel 进行按元素的运算。使用 InclusiveScanKernel 和 ReductionKernel 的 re...
复制链接

扫一扫