在用于 CUDA 程序的 NVPTX(LLVM IR) 中,有从 0 到 5 的内存地址空间标识符(见下表)。
我在同一个 LLVM IR 程序中看到,内存地址被标识为“通用”或其他类型,如图所示。
对于“通用”(默认情况下,没有标识符):
用以分享':
我的问题是,对于通用内存地址空间,数据实际上位于硬件、片外、片上内存还是本地寄存器中?有人能解释一下最终是如何管理通用类型的地址空间的吗?
最佳答案
答案很简单:通用地址空间没有硬件表示。
您可以将通用地址空间 (AS) 视为一个逻辑 AS,其中组合了其他每个 AS。
例如:以下内核调用和一个接受指针的设备函数。
__device__ void bar(int* x){
*x = *x + 1;
}
__global__ void foo(int* x){
__shared__ int y[1];
bar(x);
bar(y);
}
您可以传递任何指向该函数的指针。从语言的角度来看,指针是在 AS 1(全局)还是 AS 3(共享)中并不重要。
在 C++(和 CUDA C/C++)中,您不必明确指定 AS。例如,在 OpenCL < 2.0 中,您必须为每个指针显式添加一个修饰符,并且必须提供一个函数
bar
来获取特定的 AS 指针。在 LLVM IR 中发生的事情是,传递给函数的指针通过
addresspacecast
指令被强制转换为通用 AS。在 PTX 中
addresspacecast
由 cvta
指令表示:// convert const, global, local, or shared address to generic address
cvta.space.size p, a; // source address in register a
cvta.space.size p, var; // get generic address of var
cvta.space.size p, var+imm; // generic address of var+offset
// convert generic address to const, global, local, or shared address
cvta.to.space.size p, a;
.space = { .const, .global, .local, .shared };
.size = { .u32, .u64 };
通用指针映射到全局内存,除非它位于为另一个 AS 保留的地址区域内。硬件从通用指针中减去 AS 的起始地址以确定正确的内存区域。
原子是一个很好的例子:
atom{.space}.op.type d, [a], b;
atom{.space}.op.type d, [a], b, c;
您可以指定地址空间或让硬件选择。如果你想在没有指针减法开销的情况下生成正确的原子指令,后端负责将指针转换回正确的地址空间。
关于architecture - NVPTX 通用内存空间在架构中的位置,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/32468913/