NVPTX架构中的通用内存空间位置

Question

在用于CUDA程序的NVPTX(LLVM IR)中,存储器地址空间的标识符从0到5(参见下表).

我在同一个LLVM IR程序中看到,内存地址被标识为"Generic"或其他类型,如图片所示.

对于'Generic'(默认情况下,没有标识符):

用以分享':

我的问题是,对于通用存储器地址空间,数据实际位于硬件,片外,片上存储器或本地寄存器中的哪个位置？有人能解释一下最终管理泛型类型的地址空间吗？

Answer 1

答案很简单:通用地址空间没有硬件表示.

您可以将通用地址空间(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;

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