共享内存与Thread的同步

给出3072*3072大小的数组, 每一个元素都是整数, 现在要做的就是, 将每个元素的立方相加, 并求出最终的结果. 首先,我们先用PyCuda基础知识写出来一个可以运行的程序.

import time
import numpy as npimport pycuda.autoinit
import pycuda.driver as cuda
from pycuda.compiler import SourceModulemod = SourceModule("""__global__ void sumOfSquares(int* num, int *result, size_t N){int index = threadIdx.x + blockIdx.x * blockDim.x;int stride = blockDim.x * gridDim.x;int sum = 0;for (int i = index; i < N; i += stride) {sum += num[i]*num[i]*num[i];}result[index] = sum;}""")def test(N, np_seed):np.random.seed(np_seed)a = np.random.randint(1, 10, N)N = np.int32(N)thread_size = 256# block_size = int((N + thread_size - 1) / thread_size)# power = len(str(block_size)) - 2# block_size = int(block_size / (10**power))block_size = 32b = np.empty(thread_size*block_size, dtype=np.int32)sumOfSquares = mod.get_function("sumOfSquares")sumOfSquares(cuda.In(a), cuda.Out(b), N,block=(thread_size, 1, 1), grid=(block_size, 1))t_sum = 0for item in b:t_sum += itemprint('sum: %d'%(t_sum))if __name__ == "__main__":N = 3072*3072time_sum = 0cnt = 10for i in range(cnt):time_start = time.time()test(N, i)time_run = time.time()-time_starttime_sum += time_runprint('time: %f'%(time_sum/cnt))

从上面的代码看, 我们写并行的形式是将3072*3072(943718)的数据分成了256*32(8192)份让它们并行计算, 然后再将计算出来的8192份结果再串行向加得出最终的结果.这样就保留出了一个问题就是如果数据非常多的情况下, 我们cpu计算的压力会非常大. 为了解决这个问题, 我们就可以使用共享内存与线程同步.

ShareMemory 和Thread同步
shareMemory: 而接下来我们要使用的shared memory，是一个 block 中每个 thread 都共享的内存。它会使用在 GPU 上的内存，所以存取的速度相当快，不需要担心 latency 的问题。

声明一块ShareMemory也是十分简单的：

__shared__ int sharedata[128];

Thread同步: 因为我们要想让每个block把自己Thread的结果加起来，需要等到所有的Thread都将自己的结果结算出来。不过同步问题也没什么好说的，因为这是无论使用哪种语言在使用多线程时都需要考虑的一个问题。在Cuda编程中只需要调用一个函数:

__syncthreads()

我们现在要做的是将每个block中的每个Thread的结果再相加到一起, 也就是现在CPU只需要将每个block中的结果相加就行了, 只需要进行32次的加运算而不是8192次了. 所以我们对上面的程序要进行两方面的修改, 一方面是减小b数组的大小, 一方面是在CUDA程序中申请共享内存并且实现Thread的同步.

import time
import sys
import numpy as npimport pycuda.autoinit
import pycuda.driver as cuda
from pycuda.compiler import SourceModulemod = SourceModule("""__global__ void sumOfSquares(int* num, int *result, size_t N){extern __shared__ int shared[256];int index = threadIdx.x + blockIdx.x * blockDim.x;int stride = blockDim.x * gridDim.x;int tid = threadIdx.x;int bid = blockIdx.x;shared[tid] = 0;for (int i = index; i < N; i += stride) {shared[tid] += num[i]*num[i]*num[i];}//同步 保证每个 thread 都已经把结果写到 shared[tid] 里面__syncthreads();if (tid == 0) {for (int i = 1; i < blockDim.x; i++) {shared[0] += shared[i];}result[bid] = shared[0];}}""")def test(N, np_seed):np.random.seed(np_seed)a = np.random.randint(1, 10, N)N = np.int32(N)thread_size = 256block_size = int((N + thread_size - 1) / thread_size)power = len(str(block_size)) - 2block_size = int(block_size / (10**power))# block_size = 64b = np.empty(block_size, dtype=np.int32)sumOfSquares = mod.get_function("sumOfSquares")# b = np.int32(b)sumOfSquares(cuda.In(a), cuda.Out(b), N,block=(thread_size, 1, 1), grid=(block_size, 1))# t_sum = np.sum(b)t_sum = 0for item in b:t_sum += itemprint('sum: %d'%(t_sum))if __name__ == "__main__":N = 3072*3072time_sum = 0cnt = 10for i in range(cnt):time_start = time.time()test(N, i)time_run = time.time()-time_starttime_sum += time_runprint('time: %f'%(time_sum/cnt))

注: 在实现上述代码时我还发现了一个定义block大小的一个规律

如果是计算的形式是将矩阵中的所有元素经过计算从而取得一个值, 这样的选取的block需要自己定义一个数, 一般是32或者64, 这种形式要比第二种形式快, 而且快的很明显, 我使用3072*3072的数据用量进行对比, 发现要比第二种快了10倍左右.
如果计算的形式类似与两个矩阵的加减这种形式, 可以使用block_size = int(N + thread_size - 1) / thread_size) (N: 数据量的大小, thread_size是定义的thread的大小)这种形式会比使用上一种形式的要快. 但是快的可能不太明显.
附两种比较的例子

PyCuda学习三之--共享内存与Thread的同步相关推荐

CUDA学习(九)：共享内存
博主CUDA学习系列汇总传送门(持续更新):编程语言|CUDA入门转自 CUDA学习笔记(6) 共享内存与全局内存共享内存(Shared memory)是位于每个流处理器组(SM)中的高速内存 ...
进程间通信学习小结（共享内存）
要使用共享内存,应该有如下步骤:1.开辟一块共享内存 shmget() 2.允许本进程使用共某块共享内存 shmat() 3.写入/读出 4.禁止本进程使用这块共享内存 shmdt() 5.删除这块共 ...
IPC--进程间通信三（共享内存）
1.什么是共享内存? 共享内存就是允许两个不相关的进程访问同一个逻辑内存.共享内存是在两个正在运行的进程之间共享和传递数据的一种非常有效的方式.不同进程之间共享的内存通常安排为同一段物理内存.进程可以 ...
Linux-进程间通信（三）: 共享内存
1. 共享内存: 共享内存方式可以在多个进程直接共享数据,因为其直接使用内存,不要多余的拷贝,是速度最快的IPC方式: 共享内存有两种实现方式,使用mmap和shm方式,如下图: (1) mmap方式 ...
linux进程间通信：system V 信号量和共享内存实现进程间同步
关于信号量和共享内存的相关描述已经在前几篇提到过: 信号量:即内核维护的一个正整数,可以使用内核提供的p/v接口进行该正整数的+/-操作,它主要用来表示系统中可用资源的个数,协调各个进程有序访问资源, ...
进程间共享内存（信号量实现同步）
原创作品,允许转载,转载时请务必以超链接形式标明文章原始出处 .作者信息和本声明.否则将追究法律责任.http://yongjiuzhizhen.blog.51cto.com/7980250/132 ...
Android系统匿名共享内存Ashmem（Anonymous Shared Memory）驱动程序源代码分析
文章转载至CSDN社区罗升阳的安卓之旅,原文地址:http://blog.csdn.net/luoshengyang/article/details/6664554 在上一文章Android系统匿名共 ...
210307共享内存的读写
目录一.学习的知识点 systemV POSIX(可移植操作系统接口)是一种标准共享内存的读写二.上课没有听懂或者没有理解的地方三.当天学习的收获一.学习的知识点 systemV msgge ...
210305设计共享内存
目录一.学习的知识点共享内存的设计共享内存的头部信息共享内存的属性共享内存类共享内存类的构造函数二.上课没有听懂或者没有理解的地方三.当天学习的收获一.学习的知识点共享内存的作用 ...

PyCuda学习三之--共享内存与Thread的同步

共享内存与Thread的同步

注: 在实现上述代码时我还发现了一个定义block大小的一个规律

PyCuda学习三之--共享内存与Thread的同步相关推荐

最新文章

热门文章