Я думаю space все таки значит космос, хотя бы потому что все начинается с запуска ракеты. А ещё фраза Силии в конце «Ты бросил меня на Земле!».
Я увидел сюжет так: девушка увлекалась роботами, а парень хотел стать «крутым в космосе» (хотя не исключаю что все таки речь идет не о космосе), он ее бросил, она из мести создала армию роботов и модифицировала свое тело, превратившись в робота. Всё закончилось истреблением человечества. Через 40 лет команда, во главе с тем самым парнем поймала Силию и раз за разом пытается перезаписать ей память, наделив «светлыми» воспоминаниями, чтобы остановить уничтожение мира…
Из этих графиков я могу предположить, что CUDA как-то иначе распределяет потоки, т.к. в OpenCL размер блока влияет на время отнюдь не линейно и даже не монотонно. Возможно, именно поэтому вопросу выбора размера блока в CUDA уделяется так мало внимания.
Хотя так же можно предположить, что в моём алгоритме присутствуют дополнительные факторы, которые способствуют такой зависимости (например те же ветвления). В любом случае, мне кажется очень маловероятным, что причиной таких отличий является модель видеокарты. Особенно учитывая, что во многих статьях можно увидеть подобные результаты. Но я всё же попробую запустить Вашу «пустышку» у себя.
Смысле заключается в том, что размер блока очень сильно влияет на производительность, сильнее чем число обращений к глобальной памяти (имеется в виду, когда их немного, как в Вашем случае). Есть ещё один важный фактор — кратность размеру блока (8 в Вашем случае, а у меня например 48). Я конечно понимаю, что этот алгоритм не будет работать если просто удалить обращение к глобальной памяти, я предлагаю набросать «пустышку», которая либо вообще не будет обращаться к глобальной памяти, либо число этих обращений не будет зависеть от числа блоков, и посмотреть как будет изменяться производительность в зависимости от числа блоков.
Про дедлок, я сказал что ПОДОБНАЯ конструкция может привести к дедлоку. Просто представьте, что все процессоры заняты веткой, в которой крутится цикл, а счётчик убавить получается некому.
Статья должна выйти в ближайшем номере (я надеюсь), в дополнении от ВМК. Название: «АНАЛИЗ ВЛИЯНИЯ РАЗМЕРА РАБОЧЕЙ ГРУППЫ НА ПРОИЗВОДИТЕЛЬНОСТЬ OPENCL-РЕАЛИЗАЦИИ ВЫЧИСЛИТЕЛЬНОГО АЛГОРИТМА НА ПРИМЕРЕ МЕТОДА ГАУССА РЕШЕНИЯ СЛАУ». Если интересно, код уже опубликован: github.com/Dem0n3D/solecl правда инструкцию по сборке пока не писал, а там используется обёртка QtOpeCL.
>например, для моей Tesla блок может содержать максимум 512 потоков
Уж не знаю на какой странице такое написано, но это крайне неправдовподобно, т.к. даже для моей GT555M ограничение составляет 1024, такие данные следует запрашивать непосредственно у карты, с помощью спец. функций.
>Однако, чем больше становится блоков, тем метод MKL отстает по производительности все меньше. Для 32-х блоков он даже
незначительно обыгрывает метод SKL. Связано это с тем, что чем больше блоков, тем больше больше потоков (имеющих threadIdx.x == 0) читают переменную count из медленной глобальной памяти.
Необоснованно. Попробуйте удалить обращение к глобальной памяти и провести ту же серию экспериментов, скорее всего картина будет та же.
>Если же рассматривать изменение относительной производительности в зависимости от числа потоков в блоке, при постоянном количестве самих блоков, то тоже можно заметить некоторую закономерность. Но тут работают неизвестные автору эффекты, связанные с синхронизацией потоков в блоке, управлением warp`ами в SM.
Эта особенность непосредственно связана с предыдущим замечанием. Я недавно проводил исследование на эту тему, статья скоро должна быть опубликована в «Вестнике ННГУ». Если коротко, для максимальной производительности, размер блока должен быть кратен 32 (для всех NVIDIA) и числу ядер SM, а также превышать общее число ядер на всей карте (На самом деле исследование проводилось на OpenCL, но сомневаюсь что на CUDA результат будет другим).
И ещё замечание: такие вещи как
if (threadIdx.x == 0)
while (count != 0);
делать категорически нельзя. Если внимательно прочитать (1), то там описана работа планировщика с условными конструкциями. А именно: он сначала запускает все потоки, которые идут по одной ветке, и лишь после их завершения запускает потоки, которые идут по другой. А это значит, что потенциально подобная конструкция может привести к дедлоку.
По-моему в письме все как раз по теме, билайн предлагает установить запрет на сервисы стоимостью от 50р и т.д., про то то услуга платная там не сказано, даже наоборот «через бесплатный номер 0858».
Я за себя могу ответить: на аспирантскую стипендию :)
У меня 0.6 ставки ассистента и з.п. у меня меньше стипендии (3к), и при этом меня еще лекции заставляют читать. Да, ставка ассистента у нас в Оренбургском Государственном Университете — 6000р.
Мне вот интересно — большинство этих организаций составляют университеты/институты? Т.к. у нас в гос. университете заниматься наукой ой как невыгодно :(
Я увидел сюжет так: девушка увлекалась роботами, а парень хотел стать «крутым в космосе» (хотя не исключаю что все таки речь идет не о космосе), он ее бросил, она из мести создала армию роботов и модифицировала свое тело, превратившись в робота. Всё закончилось истреблением человечества. Через 40 лет команда, во главе с тем самым парнем поймала Силию и раз за разом пытается перезаписать ей память, наделив «светлыми» воспоминаниями, чтобы остановить уничтожение мира…
Хотя так же можно предположить, что в моём алгоритме присутствуют дополнительные факторы, которые способствуют такой зависимости (например те же ветвления). В любом случае, мне кажется очень маловероятным, что причиной таких отличий является модель видеокарты. Особенно учитывая, что во многих статьях можно увидеть подобные результаты. Но я всё же попробую запустить Вашу «пустышку» у себя.
Про дедлок, я сказал что ПОДОБНАЯ конструкция может привести к дедлоку. Просто представьте, что все процессоры заняты веткой, в которой крутится цикл, а счётчик убавить получается некому.
Статья должна выйти в ближайшем номере (я надеюсь), в дополнении от ВМК. Название: «АНАЛИЗ ВЛИЯНИЯ РАЗМЕРА РАБОЧЕЙ ГРУППЫ НА ПРОИЗВОДИТЕЛЬНОСТЬ OPENCL-РЕАЛИЗАЦИИ ВЫЧИСЛИТЕЛЬНОГО АЛГОРИТМА НА ПРИМЕРЕ МЕТОДА ГАУССА РЕШЕНИЯ СЛАУ». Если интересно, код уже опубликован: github.com/Dem0n3D/solecl правда инструкцию по сборке пока не писал, а там используется обёртка QtOpeCL.
Уж не знаю на какой странице такое написано, но это крайне неправдовподобно, т.к. даже для моей GT555M ограничение составляет 1024, такие данные следует запрашивать непосредственно у карты, с помощью спец. функций.
>Однако, чем больше становится блоков, тем метод MKL отстает по производительности все меньше. Для 32-х блоков он даже
незначительно обыгрывает метод SKL. Связано это с тем, что чем больше блоков, тем больше больше потоков (имеющих threadIdx.x == 0) читают переменную count из медленной глобальной памяти.
Необоснованно. Попробуйте удалить обращение к глобальной памяти и провести ту же серию экспериментов, скорее всего картина будет та же.
>Если же рассматривать изменение относительной производительности в зависимости от числа потоков в блоке, при постоянном количестве самих блоков, то тоже можно заметить некоторую закономерность. Но тут работают неизвестные автору эффекты, связанные с синхронизацией потоков в блоке, управлением warp`ами в SM.
Эта особенность непосредственно связана с предыдущим замечанием. Я недавно проводил исследование на эту тему, статья скоро должна быть опубликована в «Вестнике ННГУ». Если коротко, для максимальной производительности, размер блока должен быть кратен 32 (для всех NVIDIA) и числу ядер SM, а также превышать общее число ядер на всей карте (На самом деле исследование проводилось на OpenCL, но сомневаюсь что на CUDA результат будет другим).
И ещё замечание: такие вещи как
if (threadIdx.x == 0)
while (count != 0);
делать категорически нельзя. Если внимательно прочитать (1), то там описана работа планировщика с условными конструкциями. А именно: он сначала запускает все потоки, которые идут по одной ветке, и лишь после их завершения запускает потоки, которые идут по другой. А это значит, что потенциально подобная конструкция может привести к дедлоку.
У меня 0.6 ставки ассистента и з.п. у меня меньше стипендии (3к), и при этом меня еще лекции заставляют читать. Да, ставка ассистента у нас в Оренбургском Государственном Университете — 6000р.