NVPTX 后端用户指南¶
简介¶
为了支持 GPU 编程，NVPTX 后端支持 LLVM IR 的一个子集，以及用于表示 GPU 编程概念的一组定义的约定。本文档概述了后端的常规用法，包括对所用约定的描述以及一组可接受的 LLVM IR。
注意
本文档假设您已基本熟悉 CUDA 和 PTX 汇编语言。有关 CUDA 驱动程序 API 和 PTX 汇编语言的信息，请参阅 CUDA 文档。
约定¶
将函数标记为内核¶
在 PTX 中，有两种类型的函数：设备函数，只能由设备代码调用，以及内核函数，可以由主机代码调用。默认情况下，后端将发出设备函数。元数据用于将函数声明为内核函数。此元数据附加到名为 nvvm.annotations 的元数据对象，并具有以下格式
!0 = !{<function-ref>, metadata !"kernel", i32 1}
第一个参数是对内核函数的引用。以下示例显示了在 LLVM IR 中调用设备函数的内核函数。函数 @my_kernel 可从主机代码调用，但 @my_fmad 不可调用。
define float @my_fmad(float %x, float %y, float %z) {
%mul = fmul float %x, %y
%add = fadd float %mul, %z
ret float %add
}
define void @my_kernel(ptr %ptr) {
%val = load float, ptr %ptr
%ret = call float @my_fmad(float %val, float %val, float %val)
store float %ret, ptr %ptr
ret void
}
!nvvm.annotations = !{!1}
!1 = !{ptr @my_kernel, !"kernel", i32 1}
编译后，PTX 内核函数可由主机端代码调用。
地址空间¶
NVPTX 后端使用以下地址空间映射
地址空间
内存空间
0
通用
1
全局
2
内部使用
3
共享
4
常量
5
局部
每个全局变量和指针类型都分配给这些地址空间之一，其中 0 是默认地址空间。提供了内联函数，可用于在通用和非通用地址空间之间转换指针。
例如，以下 IR 将定义一个位于全局设备内存中的数组 @g。
@g = internal addrspace(1) global [4 x i32] [ i32 0, i32 1, i32 2, i32 3 ]
LLVM IR 函数可以读取和写入此数组，主机端代码可以通过名称使用 CUDA 驱动程序 API 将数据复制到其中。
请注意，由于地址空间 0 是通用空间，因此在地址空间 0 中拥有全局变量是非法的。地址空间 0 是 LLVM 中的默认地址空间，因此对于全局变量，需要使用 addrspace(N) 注解。
三元组¶
NVPTX 目标使用模块三元组在 32/64 位代码生成之间进行选择，并使用驱动程序编译器接口。三元组架构可以是 nvptx（32 位 PTX）或 nvptx64（64 位 PTX）。操作系统应该是 cuda 或 nvcl 之一，它确定生成的代码用于与驱动程序通信的接口。大多数用户希望使用 cuda 作为操作系统，这使得生成的 PTX 与 CUDA 驱动程序 API 兼容。
示例：用于 CUDA 驱动程序 API 的 32 位 PTX：nvptx-nvidia-cuda
示例：用于 CUDA 驱动程序 API 的 64 位 PTX：nvptx64-nvidia-cuda
NVPTX 内联函数¶
地址空间转换¶
‘llvm.nvvm.ptr.*.to.gen’ 内联函数¶
语法：¶
这些是重载的内联函数。您可以在任何指针类型上使用它们。
declare ptr @llvm.nvvm.ptr.global.to.gen.p0.p1(ptr addrspace(1))
declare ptr @llvm.nvvm.ptr.shared.to.gen.p0.p3(ptr addrspace(3))
declare ptr @llvm.nvvm.ptr.constant.to.gen.p0.p4(ptr addrspace(4))
declare ptr @llvm.nvvm.ptr.local.to.gen.p0.p5(ptr addrspace(5))
概述：¶
‘llvm.nvvm.ptr.*.to.gen’ 内联函数将非通用地址空间中的指针转换为通用地址空间指针。
语义：¶
这些内联函数修改指针值以成为有效的通用地址空间指针。
‘llvm.nvvm.ptr.gen.to.*’ 内联函数¶
语法：¶
这些是重载的内联函数。您可以在任何指针类型上使用它们。
declare ptr addrspace(1) @llvm.nvvm.ptr.gen.to.global.p1.p0(ptr)
declare ptr addrspace(3) @llvm.nvvm.ptr.gen.to.shared.p3.p0(ptr)
declare ptr addrspace(4) @llvm.nvvm.ptr.gen.to.constant.p4.p0(ptr)
declare ptr addrspace(5) @llvm.nvvm.ptr.gen.to.local.p5.p0(ptr)
概述：¶
‘llvm.nvvm.ptr.gen.to.*’ 内联函数将通用地址空间中的指针转换为目标地址空间中的指针。请注意，只有在已知指针的目标地址空间的地址空间时，这些内联函数才有用。使用地址空间转换内联函数将指针从一个非通用地址空间转换为另一个非通用地址空间是非法的。
语义：¶
这些内联函数修改指针值以成为目标非通用地址空间中的有效指针。
读取 PTX 特殊寄存器¶
‘llvm.nvvm.read.ptx.sreg.*’¶
语法：¶
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.warpsize()
概述：¶
‘@llvm.nvvm.read.ptx.sreg.*’ 内联函数提供对 PTX 特殊寄存器的访问，特别是内核启动边界。这些寄存器以以下方式映射到 CUDA 内建函数
CUDA 内建函数
PTX 特殊寄存器内联函数
threadId
@llvm.nvvm.read.ptx.sreg.tid.*
blockIdx
@llvm.nvvm.read.ptx.sreg.ctaid.*
blockDim
@llvm.nvvm.read.ptx.sreg.ntid.*
gridDim
@llvm.nvvm.read.ptx.sreg.nctaid.*
屏障¶
‘llvm.nvvm.barrier0’¶
语法：¶
declare void @llvm.nvvm.barrier0()
概述：¶
‘@llvm.nvvm.barrier0()’ 内联函数发出 PTX bar.sync 0 指令，相当于 CUDA 中的 __syncthreads() 调用。
选举线程¶
‘llvm.nvvm.elect.sync’¶
语法：¶
declare {i32, i1} @llvm.nvvm.elect.sync(i32 %membermask)
概述：¶
‘@llvm.nvvm.elect.sync’ 内联函数生成 elect.sync PTX 指令，该指令从由 membermask 指定的一组线程中选举一个预测的活动领导者线程。如果正在执行的线程不在 membermask 中，则行为未定义。当选线程的 laneid 捕获在 i32 返回值中。对于领导者线程，i1 返回值设置为 True，对于所有其他线程，i1 返回值设置为 False。领导者线程的选举以确定性方式发生，即每次对于相同的 membermask 都会选举相同的领导者线程。有关更多信息，请参阅 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-elect-sync。
内存屏障/栅栏¶
‘llvm.nvvm.fence.proxy.tensormap_generic.*’¶
语法：¶
declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.cta()
declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.cluster()
declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.gpu()
declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.sys()
declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cta(ptr %addr, i32 %size)
declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cluster(ptr %addr, i32 %size)
declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.gpu(ptr %addr, i32 %size)
declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.sys(ptr %addr, i32 %size)
概述：¶
@llvm.nvvm.fence.proxy.tensormap_generic.* 是一个单向栅栏，用于在通过通用 proxy<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#proxies>_ 执行的先前内存访问与通过 tensormap 代理执行的后续内存访问之间建立顺序。 nvvm.fence.proxy.tensormap_generic.release 可以形成一个释放序列，该序列与包含 nvvm.fence.proxy.tensormap_generic.acquire 代理栅栏的获取序列同步。下表描述了 LLVM 内在函数与 PTX 指令之间的映射关系。
NVVM 内在函数
PTX 指令
@llvm.nvvm.fence.proxy.tensormap_generic.release.*
fence.proxy.tensormap::generic.release.*
@llvm.nvvm.fence.proxy.tensormap_generic.acquire.*
fence.proxy.tensormap::generic.acquire.* [addr], size
地址操作数 addr 和操作数 size 共同指定内存范围 [addr, addr+size)，代理之间内存访问的顺序保证将在此范围内提供。 size 操作数唯一支持的值是 128，并且必须是立即数。无条件使用通用寻址，并且操作数 addr 指定的地址必须位于 .global 状态空间内。否则，行为未定义。有关更多信息，请参阅 PTX ISA。
算术内在函数¶
‘llvm.nvvm.idp2a.[us].[us]’ 内在函数¶
语法:¶
declare i32 @llvm.nvvm.idp2a.s.s(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)
declare i32 @llvm.nvvm.idp2a.s.u(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)
declare i32 @llvm.nvvm.idp2a.u.s(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)
declare i32 @llvm.nvvm.idp2a.u.u(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)
概述:¶
‘llvm.nvvm.idp2a.[us].[us]’ 内在函数执行 2 元素向量点积，然后进行加法。它们直接对应于 dp2a PTX 指令。
语义:¶
%a 中的 32 位值被分成 2 个 16 位值，然后扩展到 32 位。对于 ‘llvm.nvvm.idp2a.u.[us]’ 变体，使用零扩展，而对于 ‘llvm.nvvm.idp2a.s.[us]’ 变体，使用符号扩展。从 %b 中选择两个字节，如果 %is.hi 为真，则选择最高有效字节，否则选择最低有效字节。然后将这些字节扩展到 32 位。对于 ‘llvm.nvvm.idp2a.[us].u’ 变体，使用零扩展，而对于 ‘llvm.nvvm.idp2a.[us].s’ 变体，使用符号扩展。将这两个 2 元素向量的点积加到 %c 以生成返回值。
‘llvm.nvvm.idp4a.[us].[us]’ 内在函数¶
语法:¶
declare i32 @llvm.nvvm.idp4a.s.s(i32 %a, i32 %b, i32 %c)
declare i32 @llvm.nvvm.idp4a.s.u(i32 %a, i32 %b, i32 %c)
declare i32 @llvm.nvvm.idp4a.u.s(i32 %a, i32 %b, i32 %c)
declare i32 @llvm.nvvm.idp4a.u.u(i32 %a, i32 %b, i32 %c)
概述:¶
‘llvm.nvvm.idp4a.[us].[us]’ 内在函数执行 4 元素向量点积，然后进行加法。它们直接对应于 dp4a PTX 指令。
语义:¶
%a 和 %b 中的每个 4 个字节都扩展为 32 位整数，形成 2 个 <4 x i32>。对于 %a，在 ‘llvm.nvvm.idp4a.u.[us]’ 变体中使用零扩展，而在 ‘llvm.nvvm.idp4a.s.[us]’ 变体中使用符号扩展。类似地，对于 %b，在 ‘llvm.nvvm.idp4a.[us].u’ 变体中使用零扩展，而在 ‘llvm.nvvm.idp4a.[us].s’ 变体中使用符号扩展。将这两个 4 元素向量的点积加到 %c 以生成返回值。
其他内在函数¶
有关 NVPTX 内在函数的完整列表，请参阅 LLVM 源代码树中的 include/llvm/IR/IntrinsicsNVVM.td 文件。
与 Libdevice 链接¶
CUDA Toolkit 带有一个名为 libdevice 的 LLVM 位代码库，该库实现了许多常见的数学函数。此库可用作任何使用 LLVM NVPTX 目标的编译器的 高性能数学库。该库可以在 CUDA Toolkit 中的 nvvm/libdevice/ 下找到，并且每个计算架构都有一个单独的版本。
有关 libdevice 中实现的所有数学函数的列表，请参阅 libdevice 用户指南。
为了适应可能影响 libdevice 代码生成的各种与数学相关的编译器标志，库代码依赖于一个特殊的 LLVM IR 传递 (NVVMReflect) 来处理 LLVM IR 中的条件编译。此传递查找对 @__nvvm_reflect 函数的调用，并根据定义的反射参数用常量替换它们。此类条件代码通常遵循以下模式
float my_function(float a) {
if (__nvvm_reflect("FASTMATH"))
return my_function_fast(a);
else
return my_function_precise(a);
}
所有未指定反射参数的默认值为零。
NVVMReflect 传递应在优化管道的早期执行，在链接阶段之后立即执行。还建议使用 internalize 传递从生成的 PTX 中删除未使用的数学函数。对于输入 IR 模块 module.bc，建议使用以下编译流程
NVVMReflect 传递将尝试删除死代码，即使没有优化也是如此。这允许通过使用 __CUDA_ARCH 参数在所有优化级别上避免潜在的不兼容指令。
保存
module.bc中的外部函数列表将
module.bc与libdevice.compute_XX.YY.bc链接内部化列表 (1) 中不存在的所有函数
消除所有未使用的内部函数
运行
NVVMReflect传递运行标准优化管道
注意
linkonce 和 linkonce_odr 链接类型不适用于 libdevice 函数。可以使用不同的反射变量链接已链接到 libdevice 的两个 IR 模块。
由于 NVVMReflect 传递用常量替换条件语句，因此它通常会留下以下形式的死代码
entry:
..
br i1 true, label %foo, label %bar
foo:
..
bar:
; Dead code
..
因此，建议在优化管道的早期，在死代码消除之前执行 NVVMReflect。
NVPTX TargetMachine 知道如何在传递管理器开头安排 NVVMReflect；在设置传递管理器时只需使用以下代码，PassBuilder 将使用 registerPassBuilderCallbacks 让 NVPTXTargetMachine::registerPassBuilderCallbacks 将传递添加到传递管理器中
std::unique_ptr<TargetMachine> TM = ...;
PassBuilder PB(TM);
ModulePassManager MPM;
PB.parsePassPipeline(MPM, ...);
反射参数¶
libdevice 库当前使用以下反射参数来控制代码生成
标志 |
描述 |
|---|---|
|
使用将次正规数刷新为零的优化代码路径 |
此标志的值由“nvvm-reflect-ftz”模块标志确定。以下将 ftz 标志设置为 1。
!llvm.module.flags = !{!0}
!0 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
(i32 4 表示此处设置的值覆盖我们链接的其他模块中的值。有关详细信息，请参阅 LangRef <LangRef.html#module-flags-metadata>。)
执行 PTX¶
在 GPU 设备上执行 PTX 汇编的最常见方法是使用 CUDA 驱动程序 API。此 API 是 GPU 驱动程序的低级接口，允许将 PTX 代码 JIT 编译为本机 GPU 机器代码。
初始化驱动程序 API
CUdevice device;
CUcontext context;
// Initialize the driver API
cuInit(0);
// Get a handle to the first compute device
cuDeviceGet(&device, 0);
// Create a compute device context
cuCtxCreate(&context, 0, device);
将 PTX 字符串 JIT 编译为设备二进制文件
CUmodule module;
CUfunction function;
// JIT compile a null-terminated PTX string
cuModuleLoadData(&module, (void*)PTXString);
// Get a handle to the "myfunction" kernel function
cuModuleGetFunction(&function, module, "myfunction");
有关执行 PTX 汇编的完整示例，请参阅 CUDA 示例 分发版。
常见问题¶
ptxas 抱怨未定义函数：__nvvm_reflect¶
当与 libdevice 链接时，必须使用 NVVMReflect 传递。有关更多信息，请参阅 与 Libdevice 链接。
教程：简单的计算内核¶
首先，让我们看一下直接用 LLVM IR 编写的简单计算内核。该内核实现了向量加法，其中每个线程从输入向量 A 和 B 计算输出向量 C 的一个元素。为了简化起见，我们还假设只启动了一个 CTA（线程块），并且它将是一维的。
内核¶
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-nvidia-cuda"
; Intrinsic to read X component of thread ID
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind
define void @kernel(ptr addrspace(1) %A,
ptr addrspace(1) %B,
ptr addrspace(1) %C) {
entry:
; What is my ID?
%id = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind
; Compute pointers into A, B, and C
%ptrA = getelementptr float, ptr addrspace(1) %A, i32 %id
%ptrB = getelementptr float, ptr addrspace(1) %B, i32 %id
%ptrC = getelementptr float, ptr addrspace(1) %C, i32 %id
; Read A, B
%valA = load float, ptr addrspace(1) %ptrA, align 4
%valB = load float, ptr addrspace(1) %ptrB, align 4
; Compute C = A + B
%valC = fadd float %valA, %valB
; Store back to C
store float %valC, ptr addrspace(1) %ptrC, align 4
ret void
}
!nvvm.annotations = !{!0}
!0 = !{ptr @kernel, !"kernel", i32 1}
我们可以使用 LLVM llc 工具直接运行 NVPTX 代码生成器
# llc -mcpu=sm_20 kernel.ll -o kernel.ptx
注意
如果要生成 32 位代码，请将模块数据布局字符串中的 p:64:64:64 更改为 p:32:32:32，并将 nvptx-nvidia-cuda 用作目标三元组。
我们从 llc（截至 LLVM 3.4）获得的输出
//
// Generated by LLVM NVPTX Back-End
//
.version 3.1
.target sm_20
.address_size 64
// .globl kernel
// @kernel
.visible .entry kernel(
.param .u64 kernel_param_0,
.param .u64 kernel_param_1,
.param .u64 kernel_param_2
)
{
.reg .f32 %f<4>;
.reg .s32 %r<2>;
.reg .s64 %rl<8>;
// %bb.0: // %entry
ld.param.u64 %rl1, [kernel_param_0];
mov.u32 %r1, %tid.x;
mul.wide.s32 %rl2, %r1, 4;
add.s64 %rl3, %rl1, %rl2;
ld.param.u64 %rl4, [kernel_param_1];
add.s64 %rl5, %rl4, %rl2;
ld.param.u64 %rl6, [kernel_param_2];
add.s64 %rl7, %rl6, %rl2;
ld.global.f32 %f1, [%rl3];
ld.global.f32 %f2, [%rl5];
add.f32 %f3, %f1, %f2;
st.global.f32 [%rl7], %f3;
ret;
}
内核剖析¶
现在让我们剖析构成这个内核的 LLVM IR。
数据布局¶
数据布局字符串确定常用数据类型的位大小、其 ABI 对齐方式以及存储大小。对于 NVPTX，您应该使用以下之一
32 位 PTX
target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
64 位 PTX
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
目标内联函数¶
在此示例中，我们使用 @llvm.nvvm.read.ptx.sreg.tid.x 内联函数读取当前线程 ID 的 X 分量，这对应于在 PTX 中读取寄存器 %tid.x。NVPTX 后端支持大量内联函数。下面列出了一个简短的列表；有关完整列表，请参见 include/llvm/IR/IntrinsicsNVVM.td。
内联函数 |
CUDA 等效项 |
|---|---|
|
threadIdx.{x,y,z} |
|
blockIdx.{x,y,z} |
|
blockDim.{x,y,z} |
|
gridDim.{x,y,z} |
|
__syncthreads() |
地址空间¶
您可能已经注意到，LLVM IR 示例中的所有指针类型都具有显式的地址空间说明符。地址空间 1 是什么？NVIDIA GPU 设备（通常）具有四种类型的内存
全局：大型的片外内存
共享：小型片上内存，CTA 中的所有线程共享
局部：每个线程的私有内存
常量：所有线程共享的只读内存
这些不同类型的内存在 LLVM IR 中表示为地址空间。NVPTX 代码生成器还使用第五个地址空间，对应于“通用”地址空间。此地址空间可以表示任何其他地址空间中的地址（有一些例外）。这允许用户编写可以使用相同指令加载/存储内存的 IR 函数。提供内联函数来在通用地址空间和非通用地址空间之间转换指针。
有关更多信息，请参阅 地址空间 和 NVPTX 内联函数。
内核元数据¶
在 PTX 中，函数可以是 内核 函数（可从主机程序调用），也可以是 设备 函数（只能从 GPU 代码调用）。您可以将 内核 函数视为 GPU 程序中的入口点。为了将 LLVM IR 函数标记为 内核 函数，我们使用了特殊的 LLVM 元数据。NVPTX 后端将查找名为 nvvm.annotations 的命名元数据节点。此命名元数据必须包含描述 IR 的元数据列表。出于我们的目的，我们需要声明一个元数据节点，将“内核”属性分配给应作为 PTX 内核 函数发出的 LLVM IR 函数。这些元数据节点采用以下形式
!{<function ref>, metadata !"kernel", i32 1}
对于前面的示例，我们有
!nvvm.annotations = !{!0}
!0 = !{ptr @kernel, !"kernel", i32 1}
在这里，我们在 nvvm.annotations 中有一个元数据声明。此元数据使用 kernel 属性注释我们的 @kernel 函数。
运行内核¶
从 LLVM IR 生成 PTX 很好，但是我们如何在真实的 GPU 设备上执行它？CUDA 驱动程序 API 提供了一种方便的机制，用于将 PTX 加载并 JIT 编译为本机 GPU 设备，以及启动内核。该 API 类似于 OpenCL。下面显示了一个简单的示例，说明如何加载和执行我们的向量加法代码。请注意，为了简洁起见，此代码没有执行太多错误检查！
注意
您还可以使用 CUDA 工具包提供的 ptxas 工具将 PTX 离线编译为特定 GPU 架构的机器代码 (SASS)。CUDA 驱动程序 API 可以像 PTX 一样加载此类二进制文件。这对于通过预编译 PTX 内核来减少启动时间很有用。
#include <iostream>
#include <fstream>
#include <cassert>
#include "cuda.h"
void checkCudaErrors(CUresult err) {
assert(err == CUDA_SUCCESS);
}
/// main - Program entry point
int main(int argc, char **argv) {
CUdevice device;
CUmodule cudaModule;
CUcontext context;
CUfunction function;
CUlinkState linker;
int devCount;
// CUDA initialization
checkCudaErrors(cuInit(0));
checkCudaErrors(cuDeviceGetCount(&devCount));
checkCudaErrors(cuDeviceGet(&device, 0));
char name[128];
checkCudaErrors(cuDeviceGetName(name, 128, device));
std::cout << "Using CUDA Device [0]: " << name << "\n";
int devMajor, devMinor;
checkCudaErrors(cuDeviceComputeCapability(&devMajor, &devMinor, device));
std::cout << "Device Compute Capability: "
<< devMajor << "." << devMinor << "\n";
if (devMajor < 2) {
std::cerr << "ERROR: Device 0 is not SM 2.0 or greater\n";
return 1;
}
std::ifstream t("kernel.ptx");
if (!t.is_open()) {
std::cerr << "kernel.ptx not found\n";
return 1;
}
std::string str((std::istreambuf_iterator<char>(t)),
std::istreambuf_iterator<char>());
// Create driver context
checkCudaErrors(cuCtxCreate(&context, 0, device));
// Create module for object
checkCudaErrors(cuModuleLoadDataEx(&cudaModule, str.c_str(), 0, 0, 0));
// Get kernel function
checkCudaErrors(cuModuleGetFunction(&function, cudaModule, "kernel"));
// Device data
CUdeviceptr devBufferA;
CUdeviceptr devBufferB;
CUdeviceptr devBufferC;
checkCudaErrors(cuMemAlloc(&devBufferA, sizeof(float)*16));
checkCudaErrors(cuMemAlloc(&devBufferB, sizeof(float)*16));
checkCudaErrors(cuMemAlloc(&devBufferC, sizeof(float)*16));
float* hostA = new float[16];
float* hostB = new float[16];
float* hostC = new float[16];
// Populate input
for (unsigned i = 0; i != 16; ++i) {
hostA[i] = (float)i;
hostB[i] = (float)(2*i);
hostC[i] = 0.0f;
}
checkCudaErrors(cuMemcpyHtoD(devBufferA, &hostA[0], sizeof(float)*16));
checkCudaErrors(cuMemcpyHtoD(devBufferB, &hostB[0], sizeof(float)*16));
unsigned blockSizeX = 16;
unsigned blockSizeY = 1;
unsigned blockSizeZ = 1;
unsigned gridSizeX = 1;
unsigned gridSizeY = 1;
unsigned gridSizeZ = 1;
// Kernel parameters
void *KernelParams[] = { &devBufferA, &devBufferB, &devBufferC };
std::cout << "Launching kernel\n";
// Kernel launch
checkCudaErrors(cuLaunchKernel(function, gridSizeX, gridSizeY, gridSizeZ,
blockSizeX, blockSizeY, blockSizeZ,
0, NULL, KernelParams, NULL));
// Retrieve device data
checkCudaErrors(cuMemcpyDtoH(&hostC[0], devBufferC, sizeof(float)*16));
std::cout << "Results:\n";
for (unsigned i = 0; i != 16; ++i) {
std::cout << hostA[i] << " + " << hostB[i] << " = " << hostC[i] << "\n";
}
// Clean up after ourselves
delete [] hostA;
delete [] hostB;
delete [] hostC;
// Clean-up
checkCudaErrors(cuMemFree(devBufferA));
checkCudaErrors(cuMemFree(devBufferB));
checkCudaErrors(cuMemFree(devBufferC));
checkCudaErrors(cuModuleUnload(cudaModule));
checkCudaErrors(cuCtxDestroy(context));
return 0;
}
您需要链接 CUDA 驱动程序并指定 cuda.h 的路径。
# clang++ sample.cpp -o sample -O2 -g -I/usr/local/cuda-5.5/include -lcuda
我们不需要指定 libcuda.so 的路径，因为驱动程序（而不是 CUDA 工具包）将其安装在系统位置。
如果一切按计划进行，则运行编译后的程序时，您应该会看到以下输出
Using CUDA Device [0]: GeForce GTX 680
Device Compute Capability: 3.0
Launching kernel
Results:
0 + 0 = 0
1 + 2 = 3
2 + 4 = 6
3 + 6 = 9
4 + 8 = 12
5 + 10 = 15
6 + 12 = 18
7 + 14 = 21
8 + 16 = 24
9 + 18 = 27
10 + 20 = 30
11 + 22 = 33
12 + 24 = 36
13 + 26 = 39
14 + 28 = 42
15 + 30 = 45
注意
您可能会根据您的硬件看到不同的设备标识符
教程：与 Libdevice 链接¶
在本教程中，我们展示了一个与 libdevice 库链接 LLVM IR 的简单示例。我们将使用与上一个教程相同的内核，但我们将计算 C = pow(A, B) 而不是 C = A + B。Libdevice 提供了一个我们将要使用的 __nv_powf 函数。
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-nvidia-cuda"
; Intrinsic to read X component of thread ID
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind
; libdevice function
declare float @__nv_powf(float, float)
define void @kernel(ptr addrspace(1) %A,
ptr addrspace(1) %B,
ptr addrspace(1) %C) {
entry:
; What is my ID?
%id = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind
; Compute pointers into A, B, and C
%ptrA = getelementptr float, ptr addrspace(1) %A, i32 %id
%ptrB = getelementptr float, ptr addrspace(1) %B, i32 %id
%ptrC = getelementptr float, ptr addrspace(1) %C, i32 %id
; Read A, B
%valA = load float, ptr addrspace(1) %ptrA, align 4
%valB = load float, ptr addrspace(1) %ptrB, align 4
; Compute C = pow(A, B)
%valC = call float @__nv_powf(float %valA, float %valB)
; Store back to C
store float %valC, ptr addrspace(1) %ptrC, align 4
ret void
}
!nvvm.annotations = !{!0}
!0 = !{ptr @kernel, !"kernel", i32 1}
要编译此内核，我们执行以下步骤
与 libdevice 链接
内部化除公共内核函数之外的所有函数
运行
NVVMReflect并将__CUDA_FTZ设置为 0优化链接的模块
生成模块代码
这些步骤可以通过 LLVM llvm-link、opt 和 llc 工具执行。在完整的编译器中，这些步骤也可以完全以编程方式执行，方法是设置适当的传递配置（请参阅 与 Libdevice 链接）。
# llvm-link t2.bc libdevice.compute_20.10.bc -o t2.linked.bc
# opt -internalize -internalize-public-api-list=kernel -nvvm-reflect-list=__CUDA_FTZ=0 -nvvm-reflect -O3 t2.linked.bc -o t2.opt.bc
# llc -mcpu=sm_20 t2.opt.bc -o t2.ptx
注意
-nvvm-reflect-list=_CUDA_FTZ=0 不是严格必需的，因为任何未定义的变量都将默认为零。此处显示它用于评估目的。
这为我们提供了以下 PTX（摘录）
//
// Generated by LLVM NVPTX Back-End
//
.version 3.1
.target sm_20
.address_size 64
// .globl kernel
// @kernel
.visible .entry kernel(
.param .u64 kernel_param_0,
.param .u64 kernel_param_1,
.param .u64 kernel_param_2
)
{
.reg .pred %p<30>;
.reg .f32 %f<111>;
.reg .s32 %r<21>;
.reg .s64 %rl<8>;
// %bb.0: // %entry
ld.param.u64 %rl2, [kernel_param_0];
mov.u32 %r3, %tid.x;
ld.param.u64 %rl3, [kernel_param_1];
mul.wide.s32 %rl4, %r3, 4;
add.s64 %rl5, %rl2, %rl4;
ld.param.u64 %rl6, [kernel_param_2];
add.s64 %rl7, %rl3, %rl4;
add.s64 %rl1, %rl6, %rl4;
ld.global.f32 %f1, [%rl5];
ld.global.f32 %f2, [%rl7];
setp.eq.f32 %p1, %f1, 0f3F800000;
setp.eq.f32 %p2, %f2, 0f00000000;
or.pred %p3, %p1, %p2;
@%p3 bra BB0_1;
bra.uni BB0_2;
BB0_1:
mov.f32 %f110, 0f3F800000;
st.global.f32 [%rl1], %f110;
ret;
BB0_2: // %__nv_isnanf.exit.i
abs.f32 %f4, %f1;
setp.gtu.f32 %p4, %f4, 0f7F800000;
@%p4 bra BB0_4;
// %bb.3: // %__nv_isnanf.exit5.i
abs.f32 %f5, %f2;
setp.le.f32 %p5, %f5, 0f7F800000;
@%p5 bra BB0_5;
BB0_4: // %.critedge1.i
add.f32 %f110, %f1, %f2;
st.global.f32 [%rl1], %f110;
ret;
BB0_5: // %__nv_isinff.exit.i
...
BB0_26: // %__nv_truncf.exit.i.i.i.i.i
mul.f32 %f90, %f107, 0f3FB8AA3B;
cvt.rzi.f32.f32 %f91, %f90;
mov.f32 %f92, 0fBF317200;
fma.rn.f32 %f93, %f91, %f92, %f107;
mov.f32 %f94, 0fB5BFBE8E;
fma.rn.f32 %f95, %f91, %f94, %f93;
mul.f32 %f89, %f95, 0f3FB8AA3B;
// inline asm
ex2.approx.ftz.f32 %f88,%f89;
// inline asm
add.f32 %f96, %f91, 0f00000000;
ex2.approx.f32 %f97, %f96;
mul.f32 %f98, %f88, %f97;
setp.lt.f32 %p15, %f107, 0fC2D20000;
selp.f32 %f99, 0f00000000, %f98, %p15;
setp.gt.f32 %p16, %f107, 0f42D20000;
selp.f32 %f110, 0f7F800000, %f99, %p16;
setp.eq.f32 %p17, %f110, 0f7F800000;
@%p17 bra BB0_28;
// %bb.27:
fma.rn.f32 %f110, %f110, %f108, %f110;
BB0_28: // %__internal_accurate_powf.exit.i
setp.lt.f32 %p18, %f1, 0f00000000;
setp.eq.f32 %p19, %f3, 0f3F800000;
and.pred %p20, %p18, %p19;
@!%p20 bra BB0_30;
bra.uni BB0_29;
BB0_29:
mov.b32 %r9, %f110;
xor.b32 %r10, %r9, -2147483648;
mov.b32 %f110, %r10;
BB0_30: // %__nv_powf.exit
st.global.f32 [%rl1], %f110;
ret;
}