元程序设计#
在“传统”编程中,一个人编写一个完成任务的程序。在 元程序设计 ,一个人写程序 编写程序的 完成了一项任务。
听起来很复杂,所以首先,我们来看看为什么这可能是个好主意。
为什么是元编程?#
自动调谐#
CUDA程序员的很大一部分时间通常花在优化代码上。此调整可回答以下问题:
每个块的最佳线程数是多少?
我应该一次处理多少数据?
哪些数据应该加载到共享内存中,相应的块应该有多大?
如果幸运的话,您将能够在代码的执行时间内找到一个模式,并提出一个启发式方法,使您能够可靠地选择最快的版本。不幸的是,这种启发式可能变得不可靠,甚至在新一代硬件中完全失败。Pycuda试图推广的这个问题的解决方案是:
忘了启发式吧。在运行时进行基准测试,并使用最快的工具。
这是pycuda相对于cuda运行时api的一个重要优势:它允许您做出这些决策 当你的代码运行时 . 一些著名的计算软件包使用了类似的技术,其中包括atlas和fftw。虽然这些都需要相当复杂的优化驱动程序,但是可以从Python的舒适性驱动PycUDA。
数据类型#
在运行时,您的代码可能必须处理不同的数据类型。例如,它可能必须同时处理单精度浮点数和双精度浮点数。你可以对两个版本都进行预编译,但是为什么?只要生成所需的代码就行了 when 这是必要的。
为给定问题专门化代码#
如果您正在编写库,那么您的用户将要求您的库执行许多任务。想象一下,如果你能有目的地为你被要求解决的问题生成代码,而不是让代码保持不必要的泛型,从而变得缓慢,那将是多么的自由。皮库达让这成为现实。
常数比变量快#
如果问题大小因运行而异,但对相同大小的数据执行了大量内核调用,则可能需要考虑将数据大小作为常量编译到代码中。这可以有显著的性能效益,主要是由于减少提取时间和较少的寄存器压力。尤其是常数乘法比一般变量乘法更有效。
循环展开#
cuda编程指南讲述了关于 nvcc 以及它将如何为你展开循环。从2.1版开始,这根本不是真的,而且 #pragma unroll
至少根据我的经验,这只是一次不行动。使用元编程,可以动态地将Python展开到Python中所需的大小。
使用模板引擎的元编程#
如果您的元编程需求相当简单,那么在运行时生成代码的最简单方法可能是通过模板引擎。有许多python模板引擎,其中两个最突出的是 Jinja 2 和 Cheetah .
下面是一个简单的元程序,它对可配置的块大小执行向量加法。它说明了基于模板的元编程技术:
from jinja2 import Template
tpl = Template("""
__global__ void add(
{{ type_name }} *tgt,
{{ type_name }} *op1,
{{ type_name }} *op2)
{
int idx = threadIdx.x +
{{ thread_block_size }} * {{block_size}}
* blockIdx.x;
{% for i in range(block_size) %}
{% set offset = i*thread_block_size %}
tgt[idx + {{ offset }}] =
op1[idx + {{ offset }}]
+ op2[idx + {{ offset }}];
{% endfor %}
}""")
rendered_tpl = tpl.render(
type_name="float", block_size=block_size,
thread_block_size=thread_block_size)
mod = SourceModule(rendered_tpl)
在工作上下文中的这个片段可以在 examples/demo_meta_template.py
.
你还可以找到一个使用CytAin模板模板编程的矩阵乘法优化的例子。 demo_meta_matrixmul_cheetah.py
和 demo_meta_matrixmul_cheetah.template.cu
.
元编程使用 codepy
#
对于更复杂的元程序,可能需要比模板引擎能够提供更多的编程控制源代码的汇编。这个 codepy
包提供了从python数据结构生成cuda源代码的方法。
下面的示例演示了 codepy
用于元编程。它的实现与上述程序完全相同:
from codepy.cgen import FunctionBody, \
FunctionDeclaration, Typedef, POD, Value, \
Pointer, Module, Block, Initializer, Assign
from codepy.cgen.cuda import CudaGlobal
mod = Module([
FunctionBody(
CudaGlobal(FunctionDeclaration(
Value("void", "add"),
arg_decls=[Pointer(POD(dtype, name))
for name in ["tgt", "op1", "op2"]])),
Block([
Initializer(
POD(numpy.int32, "idx"),
"threadIdx.x + %d*blockIdx.x"
% (thread_block_size*block_size)),
]+[
Assign(
"tgt[idx+%d]" % (o*thread_block_size),
"op1[idx+%d] + op2[idx+%d]" % (
o*thread_block_size,
o*thread_block_size))
for o in range(block_size)]))])
mod = SourceModule(mod)
在工作上下文中的这个片段可以在 examples/demo_meta_codepy.py
.