Войти
ПрограммированиеФорумГрафика

Vulkan API (вышел!) (543 стр)

Страницы: 1542 543 544 545559 Следующая »
#8130
(Правка: 15:41) 15:40, 1 фев. 2021

™­•-=MASTER=-•™
> а GPU можно работать со страницами памяти вручную?
sparse memory, но не факт, что это физическая память.
Есть еще device address, это адрес памяти, по которому можно читать/писать в шейдере, из него можно нарезать буферов и использовать как угодно.


#8131
12:41, 3 фев. 2021

™­•-=MASTER=-•™
Зачем тебе напрямую работать со страницами памяти? И что там с безопасностью будет, подумал? Если был бы полный доступ к памяти GPU, ты бы мог чекать память остальных процессов

#8132
12:24, 5 фев. 2021

IBets
> Зачем тебе напрямую работать со страницами памяти?
ну а как ты напишешь эффективный менеджер памяти GPU в таком случае? Тебе ж по хорошему нужно захапать сразу большой кусок и менеджить в нём... Что ты можешь сделать? Хмм.. ну можешь сымитировать страничную работу контроллера памяти, разбить в своём большом куске всё на страницы, ну там по 4096 байт например и выдавать пользователю состряпанный из своих виртуальных страниц линейный блок памяти, хотя стоп... как ты это сделаешь, ты ж не операционная система или драйвер... кхе, не сделаешь.. Ну а накой тогда вообще все эти менеджеры памяти, тупо что бы быстрее выделять память, хмм... в большинстве своём она выделяет при загрузке софта или новой карты в игре, а там где она постоянно выделяется и стирается - в своём менеджере очень скоро поймаешь проблему фрагментации, если конечно не будешь постоянно двигать блоки памяти, что тоже не всегда возможно...  Собственно, как вообще делают эффективные менеджеры памяти GPU?

#8133
12:33, 5 фев. 2021

™­•-=MASTER=-•™
> Собственно, как вообще делают эффективные менеджеры памяти GPU?

ты бредишь

#8134
12:39, 5 фев. 2021

™­•-=MASTER=-•™
У вас шизофрения товарищь. Вообще бред несете

#8135
12:44, 5 фев. 2021
какой муд. придумал шину памяти в 352 бита (1080ti)? Это спецом сделано, что бы на дешёвых картах были хардверно зашитые тормоза типа? 352бит/8(бит в байте)/4(sizeof(float)) == 11 float-ов можно кинуть за раз, размер warp-а = 32, блок можно и нужно делать каратным warp-у, то есть 32/64/128/256/512/1024, ну и где же здесь 11 float-ов?... хмм...11*33 == 33, не 32!; 11*6 = 66, не 64 и тд! А что такое число 33? А это значит, что в первом блоке из 32-х нитей (каждый пишет 1 float) будет 1 лишний float, а во втором блоке (второй в kernel-е) ты транзакции по шине отправишь с пустышками в 10 float-ов...и так далее, едрить раскодрить, поэтому в одном блоке нельзя ровно/кратно работать с блоками float-ами, что бы не ловить пенальти по шине или по размеру warp_а... то ли дело шина в TeslaA1000 HBM2/2e == 4096 бита, то есть 4096/8 = 512 float-ов можно кинуть за раз и это число кратно всем размерам warp-ов соответственно...
Какой дятел вообще эти warp_ы придумал? Нельзя было сделать тупо полностью самостоятельный нити, что бы девелоперы как огня не боялись бранчей и не уродовали математические алгоритмы...
#8136
12:46, 5 фев. 2021

™­•-=MASTER=-•™
> ты бредишь
ты процитировал вопрос "как сделать" и твоим ответом было что я брежу, может просто тебе нечего сказать, ну тогда б и молчал дурачок lol
IBets, тебя это тоже касается, можете вдвоём поцеловаться в дёсны

#8137
12:49, 5 фев. 2021

™­•-=MASTER=-•™
> Какой дятел вообще эти warp_ы придумал? Нельзя было сделать тупо полностью самостоятельный нити, что бы девелоперы как огня не боялись бранчей и не
> уродовали математические алгоритмы...
нельзя, тогда CPU с гипертредингом получаются.

#8138
13:01, 5 фев. 2021

HolyDel
> нельзя, тогда CPU с гипертредингом получаются.
ну и хорошо, был бы массивно параллельный CPU :)  Да не, реально размер warp-а напрягает, если бы он ещё был программируемым... 

#8139
13:11, 5 фев. 2021
HolyDel
> нельзя
нельзя общаться с безумцами
#8140
(Правка: 14:38) 14:37, 5 фев. 2021

™­•-=MASTER=-•™
У вас реально проблемы товарищь, советую обратиться к врачу. Из вас хлещет поток чуши

#8141
14:53, 5 фев. 2021

™­•-=MASTER=-•™
> > нельзя, тогда CPU с гипертредингом получаются.
> ну и хорошо, был бы массивно параллельный CPU :)

Я думаю что ключевой вопрос здесь в том как сделать эффективный доступ в память. Коалесинг.
Понятно что в принципе то можно было бы в конвейер просто пихать инструкции с разных потоков.
У Sun был похожий мультипроцессор, по-моему даже это было ещё до https://ru.wikipedia.org/wiki/UltraSPARC_T1

#8142
(Правка: 19:56) 17:09, 5 фев. 2021

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 будет нормальным, а не дешманским дополнением ради маркетинга.

#8143
0:15, 10 фев. 2021

Хз было или нет, вот результат опроса, который недавно был:
https://www.lunarg.com/wp-content/uploads/2021/01/2020-Vulkan-Eco… y-Results.pdf

#8144
12:47, 10 фев. 2021

Более 60% считают хорошо. 10% великолепно.
Подтасовка голосов. Не иначе !

Страницы: 1542 543 544 545559 Следующая »
ПрограммированиеФорумГрафика