源程序设计--Metaprogramming

在“传统”编程中，一个人写一个程序来完成一项任务。在元编程中，一个人写一个程序来写一个程序来完成一个任务。

这听起来相当复杂——所以首先，我们将看看为什么它可能是一个好主意。

什么是元编程

自动编译

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

• 每个块的最佳线程数是多少?
• 我应该一次处理多少数据?
• 应该将哪些数据加载到共享内存中，相应的块应该有多大?

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

这是PyCUDA相对于CUDA运行时API的一个重要优势:它允许您在代码运行时做出这些决策。许多著名的计算程序包都使用了。

数据类型

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

为给定的问题专门化代码

如果您正在编写一个库，那么您的用户将要求您的库执行一些任务。想象一下，如果您能够针对被要求解决的问题有目的地生成代码，而不必保持代码不必要的泛型，从而降低速度，这将是多么自由的一件事。PyCUDA让这成为现实。

常数比变量快

如果您的问题大小因运行而异，但是您对相同大小的数据执行了大量内核调用，那么您可能需要考虑将数据大小作为常量编译到代码中。这可以带来显著的性能优势，主要是由于减少了获取时间和更少的寄存器压力。特别是，与一般的变变量乘法相比，常数乘法的效率要高得多。

循环展开

CUDA编程指南介绍了nvcc的优点，以及它将如何为您展开循环。在2.1版本中，这是不正确的，#pragma unroll是一个简单的no-op，至少根据我的经验是这样。使用元编程，您可以用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中找到。您还可以在demo_meta_matrixmul_cheetah.py和demo_meta_matrixmul_cheetah.template.cu中找到一个使用Cheetah进行模板元编程的矩阵乘法优化示例。 

元编程使用codepy

对于更复杂的元程序，对源代码集的编程控制可能比模板引擎所能提供的更多。codepy包提供了一种从Python数据结构生成CUDA源代码的方法。

下面的示例演示如何使用代码页进行元编程。其实现与上述程序完全相同

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中找到。