taifyang的博客

以像素强度值作为索引的直方图内存位置的值随着由线程索引递增的线程计算而递增，线程总数等于展平数组的大小。这里定义一个小的 INDEX函数来计算256×256大小的二维图像的特定索引值，以三通道彩色图像的展平图像数组作为内核函数的输入，其输出是相同大小的灰度图像。代码几乎与上一节中的代码相同，唯一的区别在于内核函数的调用。调用内核时需定义共享内存的大小，可以使用内核调用函数中的共享参数来指定，这里指定为256*4，因为共享内存的大小为256个整数元素，每个元素需要4个字节的存储空间。

从代码中可以看出，我们正在启动一个有10个并行线程块的内核，每个块都只有一个线程，在内核代码中打印内核执行的块 ID，可以把它视为是并行执行10个myfirst_kernel的副本，每个副本都有个唯一的线程块ID(可直接访问的blockIdx.x）和唯一的线程ID(可直接访问的threadldx.x)，这些ID将告诉我们哪个块和线程正在执行内核。多次运行该程序时，将发现每次都会以不同顺序执行。

使用CPU测量的时间也将取决于CPU高精度计时器的可用性，很多时候主机在GPU内核运行时执行异步计算，因此Python 的CPU计时器可能无法为内核执行提供正确的时间。使用record 函数记录调用内核函数前后的时间，并计算时间差得出内核函数的时间，打印在控制台上。我们可以记录两个事件:一个在代码的开头，一个在代码的结尾。创建start和 end两个事件来测量GPU代码的计时，driver类中的Event()函数用于定义事件对象，然后使用get_function创建指向内核函数的指针引用。

此矩阵通过使用gpuarray类的to_gpu()方法上载到设备显存,将上载的矩阵作为此方法的参数，然后将矩阵转换成单精度数字，上载到矩阵上的所有操作都在设备上执行。接着5×5的矩阵在主机上用随机数初始化，通过numpy.random模块的randint函数完成，这里需要3个参数:前两个参数定义用于生成随机数的数字范围，第一个参数用于生成数字的最小值，第二个参数是用于生成数字的最大值;内核使用多维线程调用，x和y方向的值为5，因此启动的线程总数是25，每个线程计算矩阵中单个元素的平方。

首先，确定系统中存在多少支持CUDA的设备是很重要的，因为一个系统可能包含多个GPU设备。系统上所有设备都会被迭代以确定每个设备的属性，使用drv.Device函数创建每个设备的指针对象，此指针用于确定特定设备的所有属性。name函数为每个特定设备提供对应名称，total_memory将给出该设备上可用的GPU全局内存的大小。其他属性存储为Python字典，可通过get_attributes().items()函数获取,通过Python使用列表解译将其转换为元组列表。

两个向量各有100个元素，用随机整数初始化以计算点乘，用Python 的time模块来计算点乘所需的时间。*运算符用于计算两个向量的元素相乘，其结果相加以计算整体点乘。请注意，numpy.dot方法只用于矩阵乘法，而不能用于点乘。使用to_gpu方法将两个向量上传给GPU执行点乘计算。gpuarray类提供了一个点乘方法可以直接计算点乘，需要两个GPU数组作为参数，使用get()方法将计算结果下载回主机。计算结果和使用CUDA事件测量的时间显示在控制台。

两个3×3的矩阵初始化为1到5的随机整数，这些矩阵使用gpuarray类的 to_gpu方法上载到设备显存。创建空的GPU数组以将结果存储在设备上，这三个变量作为参数传递给内核函数。内核函数是以矩阵大小作为x和y方向的维度来调用的。使用get()方法将结果下载回主机。控制台上打印两个输入矩阵和GPU计算的结果，矩阵乘法也在CPU上使用numpy库的dot方法计算。并与GPU结果进行了比较，验证了内核计算的结果。

通过使用可选的线程块和网格参数，还可以将每个块的线程数和内核启动的每个网格的块数作为参数提供。然后可以包含用逗号分隔的内核函数的参数，block参数表示要启动的线程数,Grid参数表示网格中的线程块数。block 和 Grid参数使用1×3的Python元组指定，该元组指示三维中的线程块和线程，内核启动的线程总数是这两个数字的乘积。的方式从Python代码启动。开发PyCUDA代码的第一步是导入代码所需的所有库，import 指令用于在文件中导入库、模块、类或函数，类似于在C或C++中的导入指令。

内核函数与前面看到的代码完全相同，在主机上创建两个由10个随机数组成的数组，这次不再创建内存并将其传输到设备，而是直接调用内核，通过使用drv.out或drv.In指定数据的方向来修改内核函数，这简化了PyCUDA代码。调用内核时，启动一个具有N个线程的线程块，这N个线程并行添加数组的N个元素，这将加速加法操作。PyCUDA为内核调用提供一个更简单的API，不需要内存分配和内存复制，这是由API隐式完成的，可以通过使用PyCUDA中 driver类的In和Out函数来实现。

PyCUDA提供了pycuda.scan.InclusiveScanKernel类来定义包含scan 操作的内核函数,要求以输出与扫描操作的数据类型为参数，为这个累加和的函数指定加法运算，以随机整数的数组作为此内核函数的输人，内核输出将与输入具有相同的大小。计算点乘的归约内核函数需要两个向量之间的元素相乘，然后求和。scan 操作是另一个非常重要的并行计算范例，指派特定函数去计算输人序列的第一项,再将这个计算结果与输人序列的第二项作为输入提供给这个特定函数计算，所有的中间计算结果形成输出的序列。