在“传统”编程中,一个人写一个程序来完成一项任务。在元编程中,一个人写一个程序来写一个程序来完成一个任务。
这听起来相当复杂——所以首先,我们将看看为什么它可能是一个好主意。
什么是元编程
自动编译
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中找到。