Буфер переменной длины PTX в разделяемой памяти

Я пытаюсь реализовать ядро ​​глобального сокращения в PTX, которое использует общую память для сокращения в блоке потока (как и все примеры CUDA C). В CUDA C on есть возможность определить массив переменной длины в разделяемой памяти с

extern __shared__ float sdata[];

Как я могу получить эквивалент в PTX?

Что не кажется подходящим, так это массив фиксированной длины, например

.shared .f32 sdata[ LENGTH ];

Так как я хочу, чтобы ядро ​​можно было повторно использовать для разных длин входного массива.

Что я мог бы сделать, это определить одну переменную

.shared .f32 sdata;

и используйте его как базовый адрес массива. В надежде, что она выделена в начале разделяемой памяти. Затем я мог бы получить доступ к элементу массива, например

ld.shared.f32 %r4,[sdata + <offset>]

Также это выглядит немного забавно, потому что sdata определяется как float. Но на самом деле это адрес поплавка. В этом смысле приведенная выше строка действительно верна. Однако я не уверен, что это гарантированно правильно, скажем, до тех пор, пока смещение не превышает размер общей памяти, указанный при запуске ядра.

Руководство PTX не объясняет буферы переменной длины в разделяемой памяти.

Кто-нибудь знает, как реализовать буфер переменной длины в PTX?


shared-memory buffer cuda ptx
person ritter    schedule 26.10.2012    source источник
comment
Я не делал этого сам, но кажется, что вы можете создать простой пример cuda C, который делает то, что вы хотите, а затем скомпилировать и сохранить вывод ptx (nvcc -ptx ...), чтобы увидеть, как он выглядит в коде ptx. .   -  person Robert Crovella    schedule 26.10.2012
comment
Нравится идея! Я сделал то, что вы предложили. Оказывается, существует недокументированный указатель константы sdata, который отмечает начало разделяемой памяти.   -  person ritter    schedule 26.10.2012
comment
Не могли бы вы опубликовать решение в качестве ответа, @Frank?   -  person harrism    schedule 29.10.2012

Ответы (2)


arrow_upward
1
arrow_downward

Это работает. Однако это не идеальное решение, поскольку оно вводит переменную связи extern.

.version 2.3
.target sm_20
.extern .shared .align 4 .b8 sdata[];
.entry func (.param .s32 param0,...)
{
 //
 // Base addresses
 mov.u64 w2,sdata;  // shared memory
 ld.shared.s32 i9,[w2+0];
}
person ritter    schedule 29.10.2012

arrow_upward
-1
arrow_downward

В CUDA C on есть возможность определить массив переменной длины в разделяемой памяти с

extern __shared__ float sdata[];

Это не массив переменной длины в обычном смысле этого термина — это просто синтаксис для доступа к динамически ограниченному объему разделяемой памяти, который устанавливается при запуске ядра.

Тот факт, что компилятор CUDA вводит определение .extern, является, TBH, неудачной деталью реализации, которую nVIDIA, к сожалению, представила как часть синтаксиса CUDA.

person einpoklum    schedule 05.02.2020