元程序设计#

在“传统”编程中,一个人编写一个完成任务的程序。在 元程序设计 ,一个人写程序 编写程序的 完成了一项任务。

听起来很复杂,所以首先,我们来看看为什么这可能是个好主意。

为什么是元编程?#

自动调谐#

CUDA程序员的很大一部分时间通常花在优化代码上。此调整可回答以下问题:

  • 每个块的最佳线程数是多少?

  • 我应该一次处理多少数据?

  • 哪些数据应该加载到共享内存中,相应的块应该有多大?

如果幸运的话,您将能够在代码的执行时间内找到一个模式,并提出一个启发式方法,使您能够可靠地选择最快的版本。不幸的是,这种启发式可能变得不可靠,甚至在新一代硬件中完全失败。Pycuda试图推广的这个问题的解决方案是:

忘了启发式吧。在运行时进行基准测试,并使用最快的工具。

这是pycuda相对于cuda运行时api的一个重要优势:它允许您做出这些决策 当你的代码运行时 . 一些著名的计算软件包使用了类似的技术,其中包括atlas和fftw。虽然这些都需要相当复杂的优化驱动程序,但是可以从Python的舒适性驱动PycUDA。

数据类型#

在运行时,您的代码可能必须处理不同的数据类型。例如,它可能必须同时处理单精度浮点数和双精度浮点数。你可以对两个版本都进行预编译,但是为什么?只要生成所需的代码就行了 when 这是必要的。

为给定问题专门化代码#

如果您正在编写库,那么您的用户将要求您的库执行许多任务。想象一下,如果你能有目的地为你被要求解决的问题生成代码,而不是让代码保持不必要的泛型,从而变得缓慢,那将是多么的自由。皮库达让这成为现实。

常数比变量快#

如果问题大小因运行而异,但对相同大小的数据执行了大量内核调用,则可能需要考虑将数据大小作为常量编译到代码中。这可以有显著的性能效益,主要是由于减少提取时间和较少的寄存器压力。尤其是常数乘法比一般变量乘法更有效。

循环展开#

cuda编程指南讲述了关于 nvcc 以及它将如何为你展开循环。从2.1版开始,这根本不是真的,而且 #pragma unroll 至少根据我的经验,这只是一次不行动。使用元编程,可以动态地将Python展开到Python中所需的大小。

使用模板引擎的元编程#

如果您的元编程需求相当简单,那么在运行时生成代码的最简单方法可能是通过模板引擎。有许多python模板引擎,其中两个最突出的是 Jinja 2Cheetah .

下面是一个简单的元程序,它对可配置的块大小执行向量加法。它说明了基于模板的元编程技术:

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.pydemo_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 .