Местоположение общей памяти NVPTX в архитектуре

Question

В NVPTX (LLVM IR) для программ CUDA есть идентификаторы для адресного пространства памяти от 0 до 5 (см. Таблицу ниже).

я видел в одной и той же программы LLVM IR, адреса памяти обозначаются как «Generic» или других типов, как показано на картинках.

Для 'Generic' (по умолчанию, нет идентификатора):

For 'Общий':

Мой вопрос заключается в том, что для общего адресного пространства памяти, где данные на самом деле находится в аппаратном обеспечении, вне чипа, встроенной памяти или локальных регистрах? Может ли кто-нибудь объяснить, как окончательно управляется общий тип адресного пространства?

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 из общего указателя для определения правильной области памяти.

Atomics являются хорошим примером:

atom{.space}.op.type d, [a], b; 
atom{.space}.op.type d, [a], b, c; 

Вы можете указать адресное пространство, или пусть аппаратные выбрать. Если вы хотите сгенерировать правильную атомную инструкцию без накладных расходов на вычитание указателя, фоновый сервер отвечает за возврат указателя обратно в правильное адресное пространство.