GPU共享内存：pycuda使用教程

最近在学习pycuda，因为资料比较少，而且杂乱，所以我就拷贝到了自己的博客

使用Python写CUDA程序有两种方式：

* Numba
* PyCUDA

numbapro现在已经不推荐使用了，功能被拆分并分别被集成到accelerate和Numba了。

例子

numba

Numba通过及时编译机制（JIT）优化Python代码，Numba可以针对本机的硬件环境进行优化，同时支持CPU和GPU的优化，并且可以和Numpy集成，使Python代码可以在GPU上运行，只需在函数上方加上相关的指令标记，

如下所示：

import numpy as np

from timeit import default_timer as timer

from numba import vectorize

@vectorize(["float32(float32, float32)"], target='cuda')

def vectorAdd(a, b):

return a + b

def main():

N = 320000000

A = np.ones(N, dtype=np.float32 )

B = np.ones(N, dtype=np.float32 )

C = np.zeros(N, dtype=np.float32 )

start = timer()

C = vectorAdd(A, B)

vectorAdd_time = timer() - start

print("c[:5] = " + str(C[:5]))

print("c[-5:] = " + str(C[-5:]))

print("vectorAdd took %f seconds " % vectorAdd_time)

if __name__ == '__main__':

main()

PyCUDA

PyCUDA的内核函数（kernel）其实就是使用C/C++编写的，通过动态编译为GPU微码，Python代码与GPU代码进行交互，如下所示：

import pycuda.autoinit

import pycuda.driver as drv

import numpy as np

from timeit import default_timer as timer

from pycuda.compiler import SourceModule

mod = SourceModule("""

__global__ void func(float *a, float *b, size_t N)

{

const int i = blockIdx.x * blockDim.x + threadIdx.x;

if (i >= N)

{

return;

}

float temp_a = a[i];

float temp_b = b[i];

a[i] = (temp_a * 10 + 2 ) * ((temp_b + 2) * 10 - 5 ) * 5;

// a[i] = a[i] + b[i];

}

""")

func = mod.get_function("func")

def test(N):

# N = 1024 * 1024 * 90 # float: 4M = 1024 * 1024

print("N = %d" % N)

N = np.int32(N)

a = np.random.randn(N).astype(np.float32)

b = np.random.randn(N).astype(np.float32)

# copy a to aa

aa = np.empty_like(a)

aa[:] = a

# GPU run

nTheads = 256

nBlocks = int( ( N + nTheads - 1 ) / nTheads )

start = timer()

func(

drv.InOut(a), drv.In(b), N,

block=( nTheads, 1, 1 ), grid=( nBlocks, 1 ) )

run_time = timer() - start

print("gpu run time %f seconds " % run_time)

# cpu run

start = timer()

aa = (aa * 10 + 2 ) * ((b + 2) * 10 - 5 ) * 5

run_time = timer() - start

print("cpu run time %f seconds " % run_time)

# check result

r = a - aa

print( min(r), max(r) )

def main():

for n in range(1, 10):

N = 1024 * 1024 * (n * 10)

print("------------%d---------------" % n)

test(N)

if __name__ == '__main__':

main()

对比

numba使用一些指令标记某些函数进行加速（也可以使用Python编写内核函数），这一点类似于OpenACC，而PyCUDA需要自己写kernel，在运行时进行编译，底层是基于C/C++实现的。通过测试，这两种方式的加速比基本差不多。但是，numba更像是一个黑盒，不知道内部到底做了什么，而PyCUDA就显得很直观。因此，这两种方式具有不同的应用：

* 如果只是为了加速自己的算法而不关心CUDA编程，那么直接使用numba会更好。

* 如果为了学习、研究CUDA编程或者实验某一个算法在CUDA下的可行性，那么使用PyCUDA。

* 如果写的程序将来要移植到C/C++，那么就一定要使用PyCUDA了，因为使用PyCUDA写的kernel本身就是用CUDA C/C++写的。

GPU共享内存：pycuda使用教程相关推荐

GPU 共享内存bank冲突(shared memory bank conflicts)
GPU 共享内存bank冲突(shared memory bank conflicts) 时间 2016-11-05 21:47:58 FindSpace 原文 http://www.findspac ...
cuda合并访问的要求_在 CUDA C / C ++ 中使用共享内存
在上一篇文章中,我研究了如何将一组线程访问的全局内存合并到一个事务中,以及对齐和跨步如何影响 CUDA 各代硬件的合并.对于最新版本的 CUDA 硬件,未对齐的数据访问不是一个大问题.然而,不管 ...
PyCuda学习三之--共享内存与Thread的同步
共享内存与Thread的同步给出3072*3072大小的数组, 每一个元素都是整数, 现在要做的就是, 将每个元素的立方相加, 并求出最终的结果. 首先,我们先用PyCuda基础知识写出来一个可以运 ...
实验9Linux共享内存通信,操作系统原理与Linux实践教程(卓越工程师培养计划系列教材)...
导语由申丰山和王黎明共同编著的<操作系统原理与Linux实践教程(卓越工程师培养计划系列教材)>一书理论与实践并重,全面.系统地阐述了操作系统的重要概念和原理,深入.细致地剖析了操作系统 ...
【GPU结构与CUDA系列4】GPU存储资源：寄存器，本地内存，共享内存，缓存，显存等存储器细节
0 软件抽象和硬件结构对应关系的例子把GPU跟一个学校对应起来,学校里有教学楼.操场.食堂,还有老师和学生们:很快有领导(CPU)来检查卫生(需要执行的任务Host程序),因此这个学校的学生们要完成 ...
盘点来自工业界的GPU共享方案
作者丨阎姝含@知乎(已授权) 来源丨https://zhuanlan.zhihu.com/p/398369404 编辑丨极市平台导读本文总结了目前有公开PR的.来自工业界的部分GPU容器计算共享方 ...
windows共享内存
在windows编程中避免不了使用共享内存,因为他是进程间通信.文件读取最简单的方式,有书上还说其他进程间通讯机制如管.油槽.WM_COPYDATA底层也是用的共享内存机制.关于使用方法还是请参考MS ...
共享内存简介及docker容器的shm设置与修改
共享内存简介及docker容器的shm设置与修改共享内存简介共享内存指 (shared memory)在多处理器的计算机系统中,可以被不同中央处理器(CPU)访问的大容量内存.由于多个CPU需要快 ...
GPU事务性内存技术研究
点击上方蓝字关注我们 GPU事务性内存技术研究林玉哲1,2, 张为华1,2 1 复旦大学软件学院,上海 201203 2 上海市数据科学重点实验室,上海 201203 论文引用格式: 林玉哲,张为华 ...
docker 不包含依赖打包_Docker打包深度学习项目（解决：Opencv依赖库、共享内存）...
最近比赛要提交打包好的CNN模型给主办方去测试,使用Docker镜像.这可难倒了我这个小白,搞了好几天才打包好.在这里记录一下Docker使用过程,希望能给广大同学们提供点帮助. 文章内容包含以下3个 ...

GPU共享内存：pycuda使用教程

GPU共享内存：pycuda使用教程相关推荐

最新文章

热门文章