Тёмная тема
- /A\
- Участник
™•-=MASTER=-•™
> а GPU можно работать со страницами памяти вручную?
sparse memory, но не факт, что это физическая память.
Есть еще device address, это адрес памяти, по которому можно читать/писать в шейдере, из него можно нарезать буферов и использовать как угодно.
™•-=MASTER=-•™
Зачем тебе напрямую работать со страницами памяти? И что там с безопасностью будет, подумал? Если был бы полный доступ к памяти GPU, ты бы мог чекать память остальных процессов
IBets
> Зачем тебе напрямую работать со страницами памяти?
ну а как ты напишешь эффективный менеджер памяти GPU в таком случае? Тебе ж по хорошему нужно захапать сразу большой кусок и менеджить в нём... Что ты можешь сделать? Хмм.. ну можешь сымитировать страничную работу контроллера памяти, разбить в своём большом куске всё на страницы, ну там по 4096 байт например и выдавать пользователю состряпанный из своих виртуальных страниц линейный блок памяти, хотя стоп... как ты это сделаешь, ты ж не операционная система или драйвер... кхе, не сделаешь.. Ну а накой тогда вообще все эти менеджеры памяти, тупо что бы быстрее выделять память, хмм... в большинстве своём она выделяет при загрузке софта или новой карты в игре, а там где она постоянно выделяется и стирается - в своём менеджере очень скоро поймаешь проблему фрагментации, если конечно не будешь постоянно двигать блоки памяти, что тоже не всегда возможно... Собственно, как вообще делают эффективные менеджеры памяти GPU?
™•-=MASTER=-•™
> Собственно, как вообще делают эффективные менеджеры памяти GPU?
ты бредишь
™•-=MASTER=-•™
У вас шизофрения товарищь. Вообще бред несете
™•-=MASTER=-•™
> ты бредишь
ты процитировал вопрос "как сделать" и твоим ответом было что я брежу, может просто тебе нечего сказать, ну тогда б и молчал дурачок lol
IBets, тебя это тоже касается, можете вдвоём поцеловаться в дёсны
™•-=MASTER=-•™
> Какой дятел вообще эти warp_ы придумал? Нельзя было сделать тупо полностью самостоятельный нити, что бы девелоперы как огня не боялись бранчей и не
> уродовали математические алгоритмы...
нельзя, тогда CPU с гипертредингом получаются.
HolyDel
> нельзя, тогда CPU с гипертредингом получаются.
ну и хорошо, был бы массивно параллельный CPU :) Да не, реально размер warp-а напрягает, если бы он ещё был программируемым...
™•-=MASTER=-•™
У вас реально проблемы товарищь, советую обратиться к врачу. Из вас хлещет поток чуши
™•-=MASTER=-•™
> > нельзя, тогда CPU с гипертредингом получаются.
> ну и хорошо, был бы массивно параллельный CPU :)
Я думаю что ключевой вопрос здесь в том как сделать эффективный доступ в память. Коалесинг.
Понятно что в принципе то можно было бы в конвейер просто пихать инструкции с разных потоков.
У Sun был похожий мультипроцессор, по-моему даже это было ещё до https://ru.wikipedia.org/wiki/UltraSPARC_T1
FROL
> Я думаю что ключевой вопрос здесь в том как сделать эффективный доступ в
> память. Коалесинг.
понимаешь в чём фокус, коалескинг далеко не всегда является проблемой в случае независимости нитей. Представим, например, такое (это на CUDA, но суть от этого не меняется):
//kernel grid dim: 32 blocks and 32 threads in each block
#define samplesNum 1024
__global__ void kernel(const float* __restrict__ input, float* __restrict__ output1, float* __restrict__ output2)
{
unsigned int id = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ float samples[samplesNum];
samples[id] = input[id];
__syncthreads();
for (int i = 0; i < samplesNum; i++)
{
float value = samples[i];
if (value > 0.5f)
output1[id] = value;
else
output2[id] = value;
}
}как видишь, сперва все нити параллельно и с коалескингом скопировали необходимые данные из глобал мемори в шаред мемори, которой собрались многократно пользоваться в цикле ниже, а вот в самом цикле торчит условие, что мол если значение больше 0.5, тогда данные улетают в такой-то буфер, а если меньше - в другой...
Ну так вот, как видишь, проблем с последовательным доступом к памяти на чтение здесь нет, но здесь есть бранч в цикле, который конкретно просаживает перформанс, т.к. в случае, если одна нить warp-а пошла по первому сценарию - другие становится бесполезными и идут вместе с ней на выход, затем шедулдер натравилает варп на другой сценарий бранча... Короче, проблемы варп-ов не касаются проблем памяти, это в общем-то отдельно стоящие проблемы...
p.s.: в данном выдуманном коде смысла в шаред мемори вообще нет, показано чисто для примера, т.к. в реальном случае эти сэмплы входного буфера будут помещены в кэш автоматически.
Я вот жду новых решений в плане гибридных центральных процессоров, которые уже замаячили на горизонте: CPU + FPGA, особенно ввиду того, что AMD купила Xillinx, ведущего производителя FPGA: https://www.ixbt.com/news/2021/01/06/amd-rabotaet-nad-processorom… ana-fpga.html
пока сложно сказать, но если это будет FPGA из коробки, причём с хорошим интеропом с CPU - будет бомба, если конечно сам FPGA будет нормальным, а не дешманским дополнением ради маркетинга.
Хз было или нет, вот результат опроса, который недавно был:
https://www.lunarg.com/wp-content/uploads/2021/01/2020-Vulkan-Eco… y-Results.pdf
Более 60% считают хорошо. 10% великолепно.
Подтасовка голосов. Не иначе !