PTX Compiler APIs
PTX 编译器 API 用户指南。
1. 简介
PTX 编译器 API 是一组 API，可用于将 PTX 程序编译为 GPU 汇编代码。
这些 API 接受字符字符串形式的 PTX 程序，并创建编译器句柄，这些句柄可用于获取 GPU 汇编代码。API 生成的 GPU 汇编代码字符串可以被 cuModuleLoadData 和 cuModuleLoadDataEx 加载，并通过 CUDA Driver API 的 cuLinkAddData 或来自 nvjitlink 的 nvJitLinkAddData API 与其他模块链接。
这些 PTX 编译器 API 的主要用例是
使用 CUDA 驱动程序 API，编译和加载是紧密结合的。PTX 编译器 API 将这两个操作解耦。这允许应用程序执行早期编译和 GPU 汇编代码缓存。
PTX 编译器 API 允许用户对 CUDA Toolkit 版本中支持的最新 PTX 版本使用运行时编译。如果应用程序在系统上安装了较旧的驱动程序，则 CUDA 驱动程序中存在的 PTX JIT 编译器可能无法提供此支持。有关更多详细信息，请参阅 CUDA 兼容性。
使用 PTX 编译器 API，客户端可以使用编译后的 GPU 汇编实现自定义缓存机制。使用 CUDA 驱动程序，无法控制 JIT 编译结果的缓存。
客户端可以精细地控制，并可以在编译期间指定编译器选项。
2. 入门指南
2.1. 系统要求
PTX 编译器库需要以下系统配置
非 Windows 平台的 POSIX 线程支持。
GPU：任何 CUDA Compute Capability 5.0 或更高的 GPU。
CUDA Toolkit 和驱动程序。
2.2. 安装
PTX 编译器库是 CUDA Toolkit 版本的一部分，组件按以下方式在 CUDA toolkit 安装目录中组织
-
在 Windows 上
include\nvPTXCompiler.hlib\x64\nvptxcompiler_static.libdoc\pdf\PTX_Compiler_API_User_Guide.pdf
-
在 Linux 上
include/nvPTXCompiler.hlib64/libnvptxcompiler_static.adoc/pdf/PTX_Compiler_API_User_Guide.pdf
3. 线程安全
所有 PTX 编译器 API 函数都是线程安全的，可以由多个线程并发调用。
4. 用户界面
本章介绍 PTX 编译器 API。API 的基本用法在 基本用法 中进行了解释。
4.1. PTX-编译器句柄
类型定义
- nvPTXCompilerHandle
-
nvPTXCompilerHandle 表示 PTX 编译器的句柄。
4.1.1. 类型定义
-
typedef struct nvPTXCompiler *nvPTXCompilerHandle
-
nvPTXCompilerHandle 表示 PTX 编译器的句柄。
要编译 PTX 程序字符串，必须创建 nvPTXCompiler 的实例，并且必须使用 API nvPTXCompilerCreate() 获取其句柄。然后可以使用 API nvPTXCompilerCompile() 完成编译。
4.2. 错误代码
枚举
- nvPTXCompileResult
-
nvPTXCompiler API 返回 nvPTXCompileResult 代码以指示调用结果。
4.2.1. 枚举
-
enum nvPTXCompileResult
-
nvPTXCompiler API 返回 nvPTXCompileResult 代码以指示调用结果。
值
-
enumerator NVPTXCOMPILE_SUCCESS
-
enumerator NVPTXCOMPILE_ERROR_INVALID_COMPILER_HANDLE
-
enumerator NVPTXCOMPILE_ERROR_INVALID_INPUT
-
enumerator NVPTXCOMPILE_ERROR_COMPILATION_FAILURE
-
enumerator NVPTXCOMPILE_ERROR_INTERNAL
-
enumerator NVPTXCOMPILE_ERROR_OUT_OF_MEMORY
-
enumerator NVPTXCOMPILE_ERROR_COMPILER_INVOCATION_INCOMPLETE
-
enumerator NVPTXCOMPILE_ERROR_UNSUPPORTED_PTX_VERSION
-
enumerator NVPTXCOMPILE_ERROR_UNSUPPORTED_DEVSIDE_SYNC
-
enumerator NVPTXCOMPILE_ERROR_CANCELLED
-
enumerator NVPTXCOMPILE_SUCCESS
4.3. API 版本控制
PTX 编译器 API 具有版本控制，因此任何新功能或 API 更改都可以通过增加 API 版本来完成。
函数
- nvPTXCompileResult nvPTXCompilerGetVersion(unsigned int *major, unsigned int *minor)
-
查询当前正在使用的 PTX 编译器 API 的
major和minor版本。
4.3.1. 函数
-
nvPTXCompileResult nvPTXCompilerGetVersion(unsigned int *major, unsigned int *minor)
-
查询当前正在使用的 PTX 编译器 API 的
major和minor版本。注意
PTX 编译器 API 的版本遵循 CUDA Toolkit 版本控制。PTX 编译器 API 版本支持的 PTX ISA 版本在此处列出 here。
- 参数
-
major – [out] PTX 编译器 API 的主版本
minor – [out] PTX 编译器 API 的次版本
- 返回值
4.4. 编译 API
函数
- nvPTXCompileResult nvPTXCompilerCompile(nvPTXCompilerHandle compiler, int numCompileOptions, const char *const *compileOptions)
-
使用给定的编译器选项编译 PTX 程序。
- nvPTXCompileResult nvPTXCompilerCreate(nvPTXCompilerHandle *compiler, size_t ptxCodeLen, const char *ptxCode)
-
获取 PTX 编译器实例的句柄，该实例使用给定的 PTX 程序
ptxCode初始化。 - nvPTXCompileResult nvPTXCompilerDestroy(nvPTXCompilerHandle *compiler)
-
销毁并清理已创建的 PTX 编译器。
- nvPTXCompileResult nvPTXCompilerGetCompiledProgram(nvPTXCompilerHandle compiler, void *binaryImage)
-
获取已编译程序的映像。
- nvPTXCompileResult nvPTXCompilerGetCompiledProgramSize(nvPTXCompilerHandle compiler, size_t *binaryImageSize)
-
获取已编译程序的映像大小。
- nvPTXCompileResult nvPTXCompilerGetErrorLog(nvPTXCompilerHandle compiler, char *errorLog)
-
查询先前为句柄看到的错误消息。
- nvPTXCompileResult nvPTXCompilerGetErrorLogSize(nvPTXCompilerHandle compiler, size_t *errorLogSize)
-
查询先前为句柄看到的错误消息的大小。
- nvPTXCompileResult nvPTXCompilerGetInfoLog(nvPTXCompilerHandle compiler, char *infoLog)
-
查询先前为句柄看到的信息消息。
- nvPTXCompileResult nvPTXCompilerGetInfoLogSize(nvPTXCompilerHandle compiler, size_t *infoLogSize)
-
查询先前为句柄看到的信息消息的大小。
- nvPTXCompileResult nvPTXCompilerSetFlowCallback(nvPTXCompilerHandle compiler, int(*callback)(void *, void *), void *payload)
-
注册一个回调函数，编译器将在调用 nvPTXCompilerCompile() 期间在 PTX 编译的不同阶段调用该函数。
4.4.1. 函数
-
nvPTXCompileResult nvPTXCompilerCompile(nvPTXCompilerHandle compiler, int numCompileOptions, const char *const *compileOptions)
-
使用给定的编译器选项编译 PTX 程序。
注意
—gpu-name (-arch) 是一个强制选项。
- 参数
-
compiler – [inout] PTX 编译器的句柄，已使用要编译的 PTX 程序初始化。可以使用句柄访问编译后的程序
numCompileOptions – [in] 数组
compileOptions的长度compileOptions – [in] 应该用于编译的编译器选项。编译器选项字符串是以 null 结尾的字符数组。有效的编译器选项列表位于 link。
- 返回值
-
nvPTXCompileResult nvPTXCompilerCreate(nvPTXCompilerHandle *compiler, size_t ptxCodeLen, const char *ptxCode)
-
获取 PTX 编译器实例的句柄，该实例使用给定的 PTX 程序
ptxCode初始化。- 参数
-
compiler – [out] 返回 PTX 编译器的句柄，该句柄使用 PTX 程序
ptxCode初始化ptxCodeLen – [in] 作为字符串传递的 PTX 程序
ptxCode的大小ptxCode – [in] 要编译的 PTX 程序，作为字符串传递。
- 返回值
-
nvPTXCompileResult nvPTXCompilerDestroy(nvPTXCompilerHandle *compiler)
-
销毁并清理已创建的 PTX 编译器。
- 参数
-
compiler – [in] 要销毁的 PTX 编译器的句柄
- 返回值
-
nvPTXCompileResult nvPTXCompilerGetCompiledProgram(nvPTXCompilerHandle compiler, void *binaryImage)
-
获取已编译程序的映像。
注意
在调用此 API 之前，应为句柄调用 nvPTXCompilerCompile() API。否则，将返回 NVPTXCOMPILE_ERROR_COMPILER_INVOCATION_INCOMPLETE。
- 参数
-
compiler – [in] 已在其上执行 nvPTXCompilerCompile() 的 PTX 编译器的句柄。
binaryImage – [out] 已编译程序的映像。客户端应为
binaryImage分配内存
- 返回值
-
nvPTXCompileResult nvPTXCompilerGetCompiledProgramSize(nvPTXCompilerHandle compiler, size_t *binaryImageSize)
-
获取已编译程序的映像大小。
注意
在调用此 API 之前，应为句柄调用 nvPTXCompilerCompile() API。否则，将返回 NVPTXCOMPILE_ERROR_COMPILER_INVOCATION_INCOMPLETE。
- 参数
-
compiler – [in] 已在其上执行 nvPTXCompilerCompile() 的 PTX 编译器的句柄。
binaryImageSize – [out] 已编译程序的映像大小
- 返回值
-
nvPTXCompileResult nvPTXCompilerGetErrorLog(nvPTXCompilerHandle compiler, char *errorLog)
-
查询先前为句柄看到的错误消息。
- 参数
-
compiler – [in] 已在其上执行 nvPTXCompilerCompile() 的 PTX 编译器的句柄。
errorLog – [out] 在先前调用 nvPTXCompilerCompiler() 中生成的错误日志。客户端应为
errorLog分配内存
- 返回值
-
nvPTXCompileResult nvPTXCompilerGetErrorLogSize(nvPTXCompilerHandle compiler, size_t *errorLogSize)
-
查询先前为句柄看到的错误消息的大小。
- 参数
-
compiler – [in] 已在其上执行 nvPTXCompilerCompile() 的 PTX 编译器的句柄。
errorLogSize – [out] 在先前调用 nvPTXCompilerCompiler() 中生成的错误日志的大小（以字节为单位）。
- 返回值
-
nvPTXCompileResult nvPTXCompilerGetInfoLog(nvPTXCompilerHandle compiler, char *infoLog)
-
查询先前为句柄看到的信息消息。
- 参数
-
compiler – [in] 已在其上执行 nvPTXCompilerCompile() 的 PTX 编译器的句柄。
infoLog – [out] 在先前调用 nvPTXCompilerCompiler() 中生成的信息日志。客户端应为
infoLog分配内存
- 返回值
-
nvPTXCompileResult nvPTXCompilerGetInfoLogSize(nvPTXCompilerHandle compiler, size_t *infoLogSize)
-
查询先前为句柄看到的信息消息的大小。
- 参数
-
compiler – [in] 已在其上执行 nvPTXCompilerCompile() 的 PTX 编译器的句柄。
infoLogSize – [out] 在先前调用 nvPTXCompilerCompiler() 中生成的信息日志的大小（以字节为单位）。
- 返回值
-
nvPTXCompileResult nvPTXCompilerSetFlowCallback(nvPTXCompilerHandle compiler, int (*callback)(void*, void*), void *payload)
-
注册一个回调函数，编译器将在调用 nvPTXCompilerCompile() 期间在 PTX 编译的不同阶段调用该函数。
回调函数决定通过返回特定值来取消编译。
回调函数必须满足以下约束条件 (1) 其签名应为
在调用回调时，编译器将始终将int callback(void* param1, void* param2);
payload传递给 param1，以便回调可以根据payload做出决策。目前它始终将 NULL 传递给 param2，param2 保留供将来扩展使用。(2) 它必须返回 1 以取消编译，或返回 0 以继续。其他返回值保留供将来使用。
(3) 它必须返回一致的值。一旦它在某一点返回 1，它必须在当前 nvPTXCompilerCompile 调用进行期间的所有后续调用中返回 1。
(4) 它必须是线程安全的。
(5) 它不得调用任何 nvrtc/libnvvm/ptx API。
- 参数
-
compiler – [in] 要在其中引入回调的已初始化 PTX 编译器的句柄。
callback – [in] 指向回调函数的函数指针。
payload – [in] 调用回调时要作为参数传递的 payload。
- 返回值
5. 编译选项
本章介绍 nvPTXCompilerCompile() API 支持的选项。
带有两个前导破折号 (--) 的选项名称是长选项名称，带有一个前导破折号 (-) 的选项名称是短选项名称。短选项名称可以代替长选项名称使用。当编译选项接受参数时，赋值运算符 (=) 用于分隔编译选项参数和编译选项名称，例如，"--gpu-name=sm_70"。或者，编译选项名称和参数可以在没有赋值运算符的情况下在单独的字符串中指定，例如，"--gpu-name""sm_70"。
--allow-expensive-optimizations (-allow-expensive-optimizations)
启用（禁用）以允许编译器使用最大可用资源（内存和编译时间）执行昂贵的优化。
如果未指定，则默认行为是为优化级别 >=
O2启用此功能。
--compile-as-tools-patch (-astoolspatch)
为 CUDA 工具编译补丁代码。
不得与
-c或-ewp结合使用。某些 PTX ISA 功能可能在此编译模式下不可用。
--compile-only (-c)
生成可重定位对象。
--def-load-cache (-dlcm)
全局/通用加载的默认缓存修饰符。
--def-store-cache (-dscm)
全局/通用存储的默认缓存修饰符。
--device-debug (-g)
为设备代码生成调试信息。
--device-function-maxrregcount N (-func-maxrregcount)
使用 -c 选项编译时，指定设备函数可以使用的最大寄存器数量。
此选项在整个程序编译中被忽略，并且不影响入口函数使用的寄存器。对于设备函数，此选项会覆盖
--maxrregcount选项指定的值。如果既未指定--device-function-maxrregcount也未指定--maxrregcount，则不假定最大值。注意
在某些情况下，
static设备函数可以安全地从调用方入口函数继承更高的寄存器计数。在这种情况下，ptx 编译器可能会应用更高的计数来编译静态函数。小于 ABI 所需的最小寄存器数的值将被编译器提升到 ABI 最小限制。
--disable-optimizer-constants (-disable-optimizer-consts)
禁用优化器常量库的使用。
--disable-warnings (-w)
禁止所有警告消息。
--dont-merge-basicblocks (-no-bb-merge)
阻止基本块合并，但会略微降低性能。
通常，ptx 编译器会尝试合并连续的基本块作为其优化过程的一部分。但是，对于可调试代码，这非常令人困惑。此选项阻止合并连续的基本块。
--entry entry,... (-e)
指定必须为其生成代码的入口函数。
此选项的入口函数名称必须以 mangled 名称指定。
--extensible-whole-program (-ewp)
生成可扩展的整个程序设备代码，这允许某些调用在与 libcudadevrt 链接之前不被解析。
--fmad (-fmad)
启用（禁用）将浮点乘法和加法/减法收缩为浮点乘加运算（FMAD、FFMA 或 DFMA）
默认值：
true
--force-load-cache (-flcm)
强制全局/通用加载使用指定的缓存修饰符。
--force-store-cache (-fscm)
强制全局/通用存储使用指定的缓存修饰符。
--generate-line-info (-lineinfo)
为设备代码生成行号信息。
--gpu-name gpuname (-arch)
指定要为其生成代码的 NVIDIA GPU 的名称。
此选项也接受虚拟计算架构，在这种情况下，代码生成将被抑制。这可以仅用于解析。
此选项的允许值：
compute_50,compute_52,compute_53,compute_60,compute_61,compute_62,compute_70,compute_72,compute_73,compute_75,compute_80,compute_86,compute_87,compute_89,compute_90,compute_90a,compute_100,compute_100a,compute_101,compute_101a,compute_120,compute_120a,sm_50,sm_52,sm_53,sm_60,sm_61,sm_62,sm_70,sm_72,sm_73,sm_75,sm_80,sm_86,sm_87,sm_89,sm_90,sm_90a,sm_100,sm_100a,sm_101,sm_101a,sm_120,sm_120a默认值：
sm_52。
--maxrregcount N (-maxrregcount)
指定 GPU 函数可以使用的最大寄存器数量。
在函数特定限制之前，更高的值通常会提高执行此函数的单个 GPU 线程的性能。但是，由于线程寄存器是从每个 GPU 上的全局寄存器池分配的，因此此选项的更高值也会减小最大线程块大小，从而减少线程并行度。因此，好的 maxrregcount 值是权衡的结果。
如果未指定此选项，则不假定最大值。小于 ABI 所需的最小寄存器数的值将被编译器提升到 ABI 最小限制。用户程序可能无法使用所有寄存器，因为某些寄存器由编译器保留。
--opt-level N (-O)
指定优化级别。
默认值：
3。
--position-independent-code (-pic)
生成位置无关代码。
默认值
对于整个程序编译：
true。否则：
false。
--preserve-relocs (-preserve-relocs)
此选项将使 ptx 编译器为变量生成可重定位引用，并在链接的可执行文件中保留为它们生成的重定位。
--return-at-end (-ret-end)
阻止优化程序末尾的返回指令
通常，ptx 编译器会优化程序末尾的返回。但是，对于可调试代码，这会导致在末尾设置断点时出现问题。此选项阻止 ptxas 优化最后一条返回指令。
--suppress-async-bulk-multicast-advisory-warning (-suppress-async-bulk-multicast-advisory-warning)
在使用 sm_90 的 cp.async.bulk{.tensor} 指令上使用 .multicast::cluster 修饰符时，禁止显示警告。
--suppress-stack-size-warning (-suppress-stack-size-warning)
禁止在无法确定堆栈大小时打印的警告。
--verbose (-v)
启用详细模式，该模式打印代码生成统计信息。
--warn-on-double-precision-use (-warn-double-usage)
如果在指令中使用 double，则发出警告。
--warn-on-local-memory-usage (-warn-lmem-usage)
如果使用本地内存，则发出警告。
--warn-on-spills (-warn-spills)
如果寄存器溢出到本地内存，则发出警告。
--warning-as-error (-Werror)
将所有警告转换为错误。
--maxntid (-maxntid)
指定线程块可以拥有的最大线程数。
如果与
-maxrregcount选项一起使用，则将忽略此选项。对于指定了.maxntid指令的入口函数，此选项也会被忽略。
--minnctapersm (-minnctapersm)
指定要映射到 SM 的最小 CTA 数量。
如果与
-maxrregcount选项一起使用，则将忽略此选项。对于指定了.minnctapersm指令的入口函数，此选项也会被忽略。
--override-directive-values (-override-directive-values)
通过相应的选项值覆盖 PTX 指令值。
此选项仅对
-minnctapersm、-maxntid和-maxregcount选项有效。
--make-errors-visible-at-exit (-make-errors-visible-at-exit)
在退出点生成所需的指令，以使内存错误和错误在退出时可见。
--oFast-compile (-Ofc)
指定优先考虑设备代码编译速度的级别。
默认值：
0。
--device-stack-protector (-device-stack-protector)
启用或禁用在设备代码中生成堆栈金丝雀。
堆栈金丝雀使利用涉及堆栈局部变量的某些类型的内存安全漏洞变得更加困难。编译器使用启发式方法来评估每个函数中此类漏洞的风险。只有那些被认为是高风险的函数才会使用堆栈金丝雀。
--g-tensor-memory-access-check (-g-tmem-access-check)
为 tcgen05 操作启用张量内存访问检查。
--split-compile (-split-compile)
指定运行编译器优化时要利用的最大并发线程数。
如果指定的值为
1，则将忽略该选项。如果指定的值为0，则线程数将为底层计算机上的 CPU 数量。
6. 基本用法
本文档的此部分使用一个简单的示例向量加法，如图 1所示，解释如何使用 PTX Compiler API 编译此 PTX 程序。为了简洁和可读性，未显示 API 返回值的错误检查。
图 1. 简单向量加法的 PTX 源代码字符串
const char *ptxCode = " \n \
.version 7.0 \n \
.target sm_50 \n \
.address_size 64 \n \
.visible .entry simpleVectorAdd( \n \
.param .u64 simpleVectorAdd_param_0, \n \
.param .u64 simpleVectorAdd_param_1, \n \
.param .u64 simpleVectorAdd_param_2 \n \
) { \n \
.reg .f32 %f<4>; \n \
.reg .b32 %r<5>; \n \
.reg .b64 %rd<11>; \n \
ld.param.u64 %rd1, [simpleVectorAdd_param_0]; \n \
ld.param.u64 %rd2, [simpleVectorAdd_param_1]; \n \
ld.param.u64 %rd3, [simpleVectorAdd_param_2]; \n \
cvta.to.global.u64 %rd4, %rd3; \n \
cvta.to.global.u64 %rd5, %rd2; \n \
cvta.to.global.u64 %rd6, %rd1; \n \
mov.u32 %r1, %ctaid.x; \n \
mov.u32 %r2, %ntid.x; \n \
mov.u32 %r3, %tid.x; \n \
mad.lo.s32 %r4, %r2, %r1, %r3; \n \
mul.wide.u32 %rd7, %r4, 4; \n \
add.s64 %rd8, %rd6, %rd7; \n \
ld.global.f32 %f1, [%rd8]; \n \
add.s64 %rd9, %rd5, %rd7; \n \
ld.global.f32 %f2, [%rd9]; \n \
add.f32 %f3, %f1, %f2; \n \
add.s64 %rd10, %rd4, %rd7; \n \
st.global.f32 [%rd10], %f3; \n \
ret; \n \
} ";
与此 PTX 程序对应的 CUDA 代码如下所示
图 2. 简单向量加法的等效 CUDA 源代码
extern "C"
__global__ void simpleVectorAdd(float *x, float *y, float *out)
{
size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
out[tid] = x[tid] + y[tid];
}
使用此 PTX 程序作为字符串，我们可以创建编译器并获取其句柄，如图 3所示。
图 3. 编译器创建和程序的初始化
nvPTXCompilerHandle compiler;
nvPTXCompilerCreate(&compiler, (size_t)strlen(ptxCode), ptxCode);
现在可以通过指定编译选项来完成编译，如图 4所示。
图 4. PTX 程序的编译
const char* compile_options[] = { "--gpu-name=sm_70",
"--verbose"
};
nvPTXCompilerCompile(compiler, 2, compile_options);
现在可以获得编译后的 GPU 汇编代码。为了获得此代码，我们首先为其分配内存。为了分配内存，我们需要查询编译后的 GPU 汇编代码映像的大小，这如图 5所示。
图 5. 查询编译后的汇编映像的大小
nvPTXCompilerGetCompiledProgramSize(compiler, &elfSize);
现在可以查询编译后的 GPU 汇编代码的映像，如图 6所示。然后可以通过将此映像传递给 CUDA Driver API 在 GPU 上执行此映像。
图 6. 查询编译后的汇编映像
elf = (char*) malloc(elfSize);
nvPTXCompilerGetCompiledProgram(compiler, (void*)elf);
当不再需要编译器时，可以销毁它，如图 7所示。
图 7. 销毁编译器
nvPTXCompilerDestroy(&compiler);
7. 示例：简单向量加法
代码 (simpleVectorAddition.c)
#include <stdio.h>
#include <string.h>
#include "cuda.h"
#include "nvPTXCompiler.h"
#define NUM_THREADS 128
#define NUM_BLOCKS 32
#define SIZE NUM_THREADS * NUM_BLOCKS
#define CUDA_SAFE_CALL(x) \
do { \
CUresult result = x; \
if (result != CUDA_SUCCESS) { \
const char *msg; \
cuGetErrorName(result, &msg); \
printf("error: %s failed with error %s\n", #x, msg); \
exit(1); \
} \
} while(0)
#define NVPTXCOMPILER_SAFE_CALL(x) \
do { \
nvPTXCompileResult result = x; \
if (result != NVPTXCOMPILE_SUCCESS) { \
printf("error: %s failed with error code %d\n", #x, result); \
exit(1); \
} \
} while(0)
const char *ptxCode = " \
.version 7.0 \n \
.target sm_50 \n \
.address_size 64 \n \
.visible .entry simpleVectorAdd( \n \
.param .u64 simpleVectorAdd_param_0, \n \
.param .u64 simpleVectorAdd_param_1, \n \
.param .u64 simpleVectorAdd_param_2 \n \
) { \n \
.reg .f32 %f<4>; \n \
.reg .b32 %r<5>; \n \
.reg .b64 %rd<11>; \n \
ld.param.u64 %rd1, [simpleVectorAdd_param_0]; \n \
ld.param.u64 %rd2, [simpleVectorAdd_param_1]; \n \
ld.param.u64 %rd3, [simpleVectorAdd_param_2]; \n \
cvta.to.global.u64 %rd4, %rd3; \n \
cvta.to.global.u64 %rd5, %rd2; \n \
cvta.to.global.u64 %rd6, %rd1; \n \
mov.u32 %r1, %ctaid.x; \n \
mov.u32 %r2, %ntid.x; \n \
mov.u32 %r3, %tid.x; \n \
mad.lo.s32 %r4, %r2, %r1, %r3; \n \
mul.wide.u32 %rd7, %r4, 4; \n \
add.s64 %rd8, %rd6, %rd7; \n \
ld.global.f32 %f1, [%rd8]; \n \
add.s64 %rd9, %rd5, %rd7; \n \
ld.global.f32 %f2, [%rd9]; \n \
add.f32 %f3, %f1, %f2; \n \
add.s64 %rd10, %rd4, %rd7; \n \
st.global.f32 [%rd10], %f3; \n \
ret; \n \
} ";
int elfLoadAndKernelLaunch(void* elf, size_t elfSize)
{
CUdevice cuDevice;
CUcontext context;
CUmodule module;
CUfunction kernel;
CUdeviceptr dX, dY, dOut;
size_t i;
size_t bufferSize = SIZE * sizeof(float);
float a;
float hX[SIZE], hY[SIZE], hOut[SIZE];
void* args[3];
CUDA_SAFE_CALL(cuInit(0));
CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, elf, 0, 0, 0));
CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, "simpleVectorAdd"));
// Generate input for execution, and create output buffers.
for (i = 0; i < SIZE; ++i) {
hX[i] = (float)i;
hY[i] = (float)i * 2;
}
CUDA_SAFE_CALL(cuMemAlloc(&dX, bufferSize));
CUDA_SAFE_CALL(cuMemAlloc(&dY, bufferSize));
CUDA_SAFE_CALL(cuMemAlloc(&dOut, bufferSize));
CUDA_SAFE_CALL(cuMemcpyHtoD(dX, hX, bufferSize));
CUDA_SAFE_CALL(cuMemcpyHtoD(dY, hY, bufferSize));
args[0] = &dX;
args[1] = &dY;
args[2] = &dOut;
CUDA_SAFE_CALL( cuLaunchKernel(kernel,
NUM_BLOCKS, 1, 1, // grid dim
NUM_THREADS, 1, 1, // block dim
0, NULL, // shared mem and stream
args, 0)); // arguments
CUDA_SAFE_CALL(cuCtxSynchronize()); // Retrieve and print output.
CUDA_SAFE_CALL(cuMemcpyDtoH(hOut, dOut, bufferSize));
for (i = 0; i < SIZE; ++i) {
printf("Result:[%ld]:%f\n", i, hOut[i]);
}
// Release resources.
CUDA_SAFE_CALL(cuMemFree(dX));
CUDA_SAFE_CALL(cuMemFree(dY));
CUDA_SAFE_CALL(cuMemFree(dOut));
CUDA_SAFE_CALL(cuModuleUnload(module));
CUDA_SAFE_CALL(cuCtxDestroy(context));
return 0;
}
int main(int _argc, char *_argv[])
{
nvPTXCompilerHandle compiler = NULL;
nvPTXCompileResult status;
size_t elfSize, infoSize, errorSize;
char *elf, *infoLog, *errorLog;
unsigned int minorVer, majorVer;
const char* compile_options[] = { "--gpu-name=sm_70",
"--verbose"
};
NVPTXCOMPILER_SAFE_CALL(nvPTXCompilerGetVersion(&majorVer, &minorVer));
printf("Current PTX Compiler API Version : %d.%d\n", majorVer, minorVer);
NVPTXCOMPILER_SAFE_CALL(nvPTXCompilerCreate(&compiler,
(size_t)strlen(ptxCode), /* ptxCodeLen */
ptxCode) /* ptxCode */
);
status = nvPTXCompilerCompile(compiler,
2, /* numCompileOptions */
compile_options); /* compileOptions */
if (status != NVPTXCOMPILE_SUCCESS) {
NVPTXCOMPILER_SAFE_CALL(nvPTXCompilerGetErrorLogSize(compiler, &errorSize));
if (errorSize != 0) {
errorLog = (char*)malloc(errorSize+1);
NVPTXCOMPILER_SAFE_CALL(nvPTXCompilerGetErrorLog(compiler, errorLog));
printf("Error log: %s\n", errorLog);
free(errorLog);
}
exit(1);
}
NVPTXCOMPILER_SAFE_CALL(nvPTXCompilerGetCompiledProgramSize(compiler, &elfSize));
elf = (char*) malloc(elfSize);
NVPTXCOMPILER_SAFE_CALL(nvPTXCompilerGetCompiledProgram(compiler, (void*)elf));
NVPTXCOMPILER_SAFE_CALL(nvPTXCompilerGetInfoLogSize(compiler, &infoSize));
if (infoSize != 0) {
infoLog = (char*)malloc(infoSize+1);
NVPTXCOMPILER_SAFE_CALL(nvPTXCompilerGetInfoLog(compiler, infoLog));
printf("Info log: %s\n", infoLog);
free(infoLog);
}
NVPTXCOMPILER_SAFE_CALL(nvPTXCompilerDestroy(&compiler));
// Load the compiled GPU assembly code 'elf'
elfLoadAndKernelLaunch(elf, elfSize);
free(elf);
return 0;
}
7.1. 构建说明
假设环境变量 CUDA_PATH 指向 CUDA 工具包安装目录，请按如下方式构建此示例
-
Windows
cl.exe simpleVectorAddition.c /FesimpleVectorAddition ^ /I "%CUDA_PATH%"\include ^ "%CUDA_PATH%"\lib\x64\nvptxcompiler_static.lib "%CUDA_PATH%"\lib\x64\cuda.lib或
nvcc simpleVectorAddition.c -ccbin <PATH_TO_cl.exe> -I $CUDA_PATH/include -L $CUDA_PATH/lib/x64/ -lcuda nvptxcompiler_static.lib -
Linux
gcc simpleVectorAddition.c -o simpleVectorAddition \ -I $CUDA_PATH/include \ -L $CUDA_PATH/lib64 \ libnvptxcompiler_static.a -lcuda -lm -lpthread \ -Wl,-rpath,$CUDA_PATH/lib64
7.2. 声明
7.2.1. 声明
本文档仅供参考，不得视为对产品的特定功能、状况或质量的保证。NVIDIA Corporation (“NVIDIA”) 对本文档中包含的信息的准确性或完整性不作任何明示或暗示的陈述或保证，并且对本文档中包含的任何错误不承担任何责任。NVIDIA 对因使用此类信息或因使用此类信息而可能导致的侵犯第三方专利或其他权利的行为的后果或使用不承担任何责任。本文档不构成对开发、发布或交付任何材料（如下定义）、代码或功能的承诺。
NVIDIA 保留随时更正、修改、增强、改进和对本文档进行任何其他更改的权利，恕不另行通知。
客户在下订单前应获取最新的相关信息，并应验证此类信息是否为最新且完整。
NVIDIA 产品的销售受订单确认时提供的 NVIDIA 标准销售条款和条件的约束，除非 NVIDIA 和客户的授权代表签署的个别销售协议（“销售条款”）另有约定。NVIDIA 在此明确反对将任何客户通用条款和条件应用于购买本文档中引用的 NVIDIA 产品。本文档未直接或间接形成任何合同义务。
NVIDIA 产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于 NVIDIA 产品的故障或失灵可能合理预期会导致人身伤害、死亡或财产或环境损害的应用。NVIDIA 对在上述设备或应用中包含和/或使用 NVIDIA 产品不承担任何责任，因此，此类包含和/或使用由客户自行承担风险。
NVIDIA 不保证或声明基于本文档的产品将适用于任何特定用途。NVIDIA 不一定会执行每个产品的所有参数的测试。客户全权负责评估和确定本文档中包含的任何信息的适用性，确保产品适合客户计划的应用并符合要求，并为应用执行必要的测试，以避免应用或产品出现缺陷。客户产品设计的缺陷可能会影响 NVIDIA 产品的质量和可靠性，并可能导致超出本文档中包含的附加或不同条件和/或要求。NVIDIA 对可能基于或归因于以下原因的任何缺陷、损坏、成本或问题不承担任何责任：(i) 以任何违反本文档的方式使用 NVIDIA 产品或 (ii) 客户产品设计。
本文档未授予 NVIDIA 专利权、版权或其他 NVIDIA 知识产权下的任何明示或暗示的许可。NVIDIA 发布的有关第三方产品或服务的信息不构成 NVIDIA 授予使用此类产品或服务的许可，也不构成对此类产品或服务的保证或认可。使用此类信息可能需要获得第三方的专利或其他知识产权下的第三方许可，或者获得 NVIDIA 的专利或其他 NVIDIA 知识产权下的 NVIDIA 许可。
只有在事先获得 NVIDIA 书面批准的情况下，才允许复制本文档中的信息，并且复制时不得进行更改，必须完全遵守所有适用的出口法律和法规，并附带所有相关的条件、限制和声明。
本文档和所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“材料”，单独称为“材料”）均按“原样”提供。NVIDIA 对材料不作任何明示、暗示、法定或其他形式的保证，并明确否认所有关于不侵权、适销性和针对特定用途的适用性的暗示保证。在法律未禁止的范围内，在任何情况下，NVIDIA 均不对因使用本文档而引起的任何损害（包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害，无论如何引起，也无论责任理论如何）承担责任，即使 NVIDIA 已被告知可能发生此类损害。尽管客户可能因任何原因遭受任何损害，但 NVIDIA 对本文所述产品的客户承担的总体和累积责任应根据产品的销售条款进行限制。
7.2.2. OpenCL
OpenCL 是 Apple Inc. 的商标，已授权 Khronos Group Inc. 使用。
7.2.3. 商标
NVIDIA 和 NVIDIA 徽标是 NVIDIA Corporation 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。