PTX编译器API
PTX编译器API用户指南。
1. 简介
PTX编译器API是一组用于将PTX程序编译成GPU汇编代码的应用程序接口。
这些API接口接受字符串形式的PTX程序，并创建编译器句柄，可用于获取GPU汇编代码。由API生成的GPU汇编代码字符串可以通过cuModuleLoadData和cuModuleLoadDataEx加载，并通过CUDA驱动API的cuLinkAddData或nvjitlink中的nvJitLinkAddData API与其他模块链接。
这些PTX编译器API的主要使用场景包括：
2. 快速入门
2.1. 系统要求
PTX编译器库需要以下系统配置：
非Windows平台支持POSIX线程。
GPU: 任何支持CUDA计算能力5.0或更高版本的GPU。
CUDA工具包和驱动程序。
2.2. 安装
PTX编译器库是CUDA工具包发布的一部分，其组件在CUDA工具包安装目录中的组织结构如下：
-
在Windows上：
include\nvPTXCompiler.hlib\x64\nvptxcompiler_static.libdoc\pdf\PTX_Compiler_API_用户指南.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版本来实现。
Functions
- 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工具包的版本控制。PTX编译器API版本所支持的PTX ISA版本列在此处。
- Parameters
-
major – [out] PTX编译器API的主版本号
minor – [out] PTX编译器API的次要版本号
- Returns
4.4. 编译API
Functions
- nvPTXCompileResult nvPTXCompilerCompile(nvPTXCompilerHandle compiler, int numCompileOptions, const char *const *compileOptions)
-
使用给定的编译器选项编译PTX程序。
- nvPTXCompileResult nvPTXCompilerCreate(nvPTXCompilerHandle *compiler, size_t ptxCodeLen, const char *ptxCode)
-
获取已使用给定PTX程序
ptxCode初始化的PTX编译器实例的句柄。 - 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) 是一个必选参数。
- Parameters
-
compiler – [inout] 一个指向PTX编译器的句柄，该编译器已使用待编译的PTX程序初始化。可以通过该句柄访问编译后的程序。
numCompileOptions – [输入] 数组
compileOptions的长度compileOptions – [输入] 用于指定编译过程的编译器选项。该编译器选项字符串是一个以空字符结尾的字符数组。有效的编译器选项列表请参见link。
- Returns
-
nvPTXCompileResult nvPTXCompilerCreate(nvPTXCompilerHandle *compiler, size_t ptxCodeLen, const char *ptxCode)
-
获取已使用给定PTX程序
ptxCode初始化的PTX编译器实例的句柄。- Parameters
-
compiler – [out] 返回一个已用PTX程序
ptxCode初始化的PTX编译器句柄ptxCodeLen – [输入] 作为字符串传递的PTX程序
ptxCode的大小ptxCode – [输入] 要编译的PTX程序，以字符串形式传递。
- Returns
-
nvPTXCompileResult nvPTXCompilerDestroy(nvPTXCompilerHandle *compiler)
-
销毁并清理已创建的PTX编译器。
- Parameters
-
compiler – [输入] 指向待销毁的PTX编译器句柄
- Returns
-
nvPTXCompileResult nvPTXCompilerGetCompiledProgram(nvPTXCompilerHandle compiler, void *binaryImage)
-
获取已编译程序的镜像。
注意
nvPTXCompilerCompile() API应在调用此API前先对句柄进行调用。否则将返回NVPTXCOMPILE_ERROR_COMPILER_INVOCATION_INCOMPLETE错误。
- Parameters
-
compiler – [in] 一个已执行nvPTXCompilerCompile()操作的PTX编译器句柄。
binaryImage – [输出] 编译后程序的映像。客户端需要为
binaryImage分配内存。
- Returns
-
nvPTXCompileResult nvPTXCompilerGetCompiledProgramSize(nvPTXCompilerHandle compiler, size_t *binaryImageSize)
-
获取已编译程序映像的大小。
注意
nvPTXCompilerCompile() API应在调用此API之前为句柄调用。否则将返回NVPTXCOMPILE_ERROR_COMPILER_INVOCATION_INCOMPLETE。
- Parameters
-
compiler – [in] 一个指向PTX编译器的句柄，该编译器上已执行过nvPTXCompilerCompile()操作。
binaryImageSize – [输出] 编译后程序映像的大小
- Returns
-
nvPTXCompileResult nvPTXCompilerGetErrorLog(nvPTXCompilerHandle compiler, char *errorLog)
-
查询该句柄之前看到的错误信息。
- Parameters
-
compiler – [in] 一个已执行过nvPTXCompilerCompile()的PTX编译器句柄。
errorLog – [输出] 在之前调用nvPTXCompilerCompiler()时生成的错误日志。客户端需要为
errorLog分配内存。
- Returns
-
nvPTXCompileResult nvPTXCompilerGetErrorLogSize(nvPTXCompilerHandle compiler, size_t *errorLogSize)
-
查询之前为该句柄看到的错误消息的大小。
- Parameters
-
compiler – [输入] 一个指向PTX编译器的句柄，该句柄已执行过nvPTXCompilerCompile()操作。
errorLogSize – [out] 错误日志的大小（以字节为单位），该日志是在之前调用nvPTXCompilerCompiler()时生成的。
- Returns
-
nvPTXCompileResult nvPTXCompilerGetInfoLog(nvPTXCompilerHandle compiler, char *infoLog)
-
查询之前为该句柄看到的信息消息。
- Parameters
-
compiler – [in] 一个已执行过nvPTXCompilerCompile()操作的PTX编译器句柄。
infoLog – [out] 信息日志，由之前调用nvPTXCompilerCompiler()时生成。客户端需要为
infoLog分配内存。
- Returns
-
nvPTXCompileResult nvPTXCompilerGetInfoLogSize(nvPTXCompilerHandle compiler, size_t *infoLogSize)
-
查询之前为该句柄看到的信息消息的大小。
- Parameters
-
compiler – [in] 一个已执行过nvPTXCompilerCompile()的PTX编译器句柄。
infoLogSize – [输出] 表示前次调用nvPTXCompilerCompiler()时生成的信息日志大小（以字节为单位）。
- Returns
-
nvPTXCompileResult nvPTXCompilerSetFlowCallback(nvPTXCompilerHandle compiler, int (*callback)(void*, void*), void *payload)
-
注册一个回调函数，编译器在调用nvPTXCompilerCompile()期间会在PTX编译的不同阶段调用该函数。
回调函数通过返回特定值来决定取消编译。
回调函数必须满足以下约束条件 (1) 其签名应为
When invoking the callback, the compiler will always passint callback(void* param1, void* param2);
payloadto param1 so that the callback may make decisions based onpayload. It’ll always pass NULL to param2 for now which is reserved for future extensions.(2) 它必须返回1以取消编译，或返回0以继续。其他返回值保留供未来使用。
(3) 它必须返回一致的值。一旦在某一点返回1，在当前进行中的nvPTXCompilerCompile调用期间，后续所有调用都必须返回1。
(4) 它必须是线程安全的。
(5) 它不得调用任何nvrtc/libnvvm/ptx API。
- Parameters
-
compiler – [in] 一个指向已初始化PTX编译器的句柄，用于引入回调函数。
callback – [in] 指向回调函数的函数指针。
payload – [in] 调用回调函数时作为参数传递的有效载荷。
- Returns
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)
指定必须为其生成代码的入口函数。
此选项的入口函数名称必须以修饰名形式指定。
--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)
警告：如果在指令中使用了双精度浮点数。
--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)
启用或禁用设备代码中栈保护机制的生成。
栈保护机制（Stack canaries）使得利用涉及栈局部变量的某些内存安全漏洞变得更加困难。编译器会使用启发式方法来评估每个函数中此类漏洞的风险。只有那些被判定为高风险的函数才会使用栈保护机制。
--g-tensor-memory-access-check (-g-tmem-access-check)
为tcgen05操作启用张量内存访问检查。
--split-compile (-split-compile)
指定运行编译器优化时可使用的最大并发线程数。
如果指定的值为
1，该选项将被忽略。 如果指定的值为0，则线程数将等于底层机器的CPU数量。
6. 基础用法
本文档的这一部分使用一个简单示例向量加法（如图1所示）来说明如何使用PTX编译器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驱动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
-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公司（“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对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。
7.2.2. OpenCL
OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。
7.2.3. 商标
NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。