NVPTX 通用内存 space 在体系结构中的位置

Question

在用于 CUDA 程序的 NVPTX(LLVM IR) 中，内存地址 space 有从 0 到 5 的标识符（参见下面的 Table）。

我在同一个LLVM IR程序中看到，内存地址被识别为'Generic'或其他类型，如图所示。

对于 'Generic'（默认情况下，没有标识符）：

对于'Shared'：

我的问题是，对于通用内存地址space，数据实际位于硬件、片外、片上内存或本地寄存器中的什么位置？有人可以解释通用类型的地址 space 最终是如何管理的吗？

Answer 1

答案很简单：通用地址 space 没有硬件表示。

您可以将通用地址 space (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 中，您必须显式地为每个指针添加一个修饰符，并且必须提供一个采用特定 AS 指针的函数 bar。

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;

您可以指定地址 space 或让硬件选择。如果您想在没有指针减法开销的情况下生成正确的原子指令，后端负责将指针强制转换回正确的地址 space.

NVPTX 通用内存 space 在体系结构中的位置

NVPTX generic memory space location in architecture

architecture

cuda

llvm

llvm-ir

ptx