Мне видится в этом тезисе внутреннее противоречие и действие каких-то психологических защитных механизмов, если угодно. Вот смотрите. Все признают, что, в данном случае IBM, производит на свет всякие крутые поделки и генерирует серьезные, часто научные, идеи. Возможности, «способности» и авторитет этой фирмы никто не ставит под сомнение. Обратимся теперь к противоположной стороне — Apple. Говорится, что они очень умело упаковывают идеи, и что это, может быть, еще сложнее и за это надо еще больше уважать. Во-первых на этот тезис можно посмотреть с кислой миной хотя бы потому, что есть такая штука мода, которая заставляет бежать и покупать новый телефон за 100к рублей. Ну, а если взглянуть немного по-другому…
Упаковывают… Обычно этим занимается упаковщик у конвейера. То есть, важную роль играет аккуратность рук и правила по упаковке, которым можно обучить практически любого. Перенося пример на Apple, «аккуратность рук» рук сохраняется, а вместо «правил упаковки» можно подставить внимательное наблюдение за «трендами», да и вообще влияние на них самих. Это, конечно, тоже работа, но совсем их другой области, нежели тема поста выше. Сравнивать две фирмы через сравнение этих двух слабо пересекающихся областей несколько ошибочно. Можно, разве что, через приносимую в $ прибыль. То есть, сравнивать успешность компаний. А это уже не то, что пытались сравнить изначально. Это провоцирует потенциально бесконечные холивары. Предположим теперь, что упаковывание — это вовсе не раскладывание по пакетикам и завязывание бантиков, которые сейчас модны, а серьезная, т.н. «креативная», работа: надо много чего проанализировать, подобрать цвета, формы и тому подобное. Но это же и есть производство идей! То есть, на самом деле уже на сравнивают не белое с круглым, вещи из одной области. Но почему-то люди иде по упаковке называют просто упаковкой… Закрадывается подозрение, что они так поступают специально, но неосознанно: черное делают круглым, вызывая проблему сравнения. А потсупают так потому, что понимают — упаковать вовсе не круче, чем придумать.
Лично я считаю, что деятельность в роде работы Apple более культурно и социально зависима, без человека она очень мало значит (в том смысле, что все человеки пропали, а их наследие оценивает кто-то другой). Да и об очередной модели телефона сами люди уже забудут через 5-10 лет. Деятельность же IBM оказывается более фундаментальной, более важной, определяющей ход развития множества областей жизни человека. Результаты этой деятельности более безальтернативны.
В код влезать лучше не буду, а подожду статьи. Думаю, такая «пустышка» годна. По вертикальной оси абсолютное время, что выдает cudaEventElapsedTime(...):
Особо отмечу, то запускалось на Tesla с 30-ю SM. Так что, двукратное «увеличение крутизны» при переходе от 30 к 31 блокам понятно: одному SM приходится выполнять два одинаковых блока. Следовательно, он тратит в ~2 раза больше времени. А вся эскадра плывет со скоростью самого медленного корабля…
А тут запускался 1 блок, но число потоков в нем менялось шагом 1. Ступеньки, как я понимаю, обусловлены размерностью warp`ов. Кстати, интересно, что происходит внутри: формально появляются warp`ы с числом потоков меньше 32. Красная пунктирная кривая отражает просто пропорциональную зависимость. Из этого графика становится понятно, что лучше блоки делать не «толще» 256-ти потоков.
Код одного большого ядра разделяется на части в два отдельных. Они будут последовательно вызываться хостом. Одно ядро делает этап (i), второе — этап (ii). Делается это, как писали коллеги выше, для:
Простоты (кода, отладки);
Возможности использовать другие типы структур памяти (правда, наверно, в моей задаче это не поможет; надо еще почитать, чтобы ответить точно);
Синхронизации всех потоков запущенного grid`а, что как раз и необходимо в первую очередь для моего алгоритма.
1) В качестве подтверждения своих слов о максимальном количестве потоков в блоке, привожу значения некоторых полей структуры, которую дала мне cudaGetDeviceProperties(...). Судя по интернетам, у Вас ядро GF106 с поддержкой версии 2.1. Выглядит поновее того, что доступно для моего пользования.
2) А как можно удалить обращения к глобальной памяти, если именно там хранится переменная count, инкрементируемая «отработавшими» блоками? Именно она и считывается в циклах. Чем больше потоков, тем больше циклов, тем больше считываюний…
3) Про кратность 32-м и превышение количеством числа ядер говорится в «CUDA C Best Practices Guide». Но я правда не знаю причин тому, и не понимаю, как это следует из Вашего замечания №2. Но в любом случае с интересом ознакомлюсь с Вашей работой! Судя по сайту «Вестника ННГУ», скачать электронный вариант можно. У Вас в каком номере?
4) Как раз по невнимательности подробности обработки условных конструкций пропустил. Но как из того, что сперва выполняется одна ветка, а уже после — другая, может следовать дедлок? Каждый поток threadIdx.x == 0 непременным образом сделает "+1" к счетчику. Поэтому и все threadIdx.x == 0 выйдут из цикла.
Да, пока писал пост и, затем, читал комментарии, стал все больше с клонятся к разбиению ядра на части. Однако меня продолжает смущать такой момент. Мое большое ядро делает несколько тысяч последовательных преобразований X->P->X, используя при этом одинаковый для всех потоков массив, из которого выбираются весовые коэффициенты. Его размер около 5KB, и его копии хранятся в разделяемой памяти каждого блока. Если из одного ядра, делать много меленьких, то каждый запуск оных должен же сопровождаться записью этого массива в разделяемую память? Если я прав, то получается очень много лишних обращений к разделяемой памяти.
Открою страшную тайну: я решил доморощенным образом вычислять преобразование Фурье. Весьма вероятно именно поэтому мое ядро очень прожорливо до регистров, и на каждом SM я запускаю по одному блоку в 256 потоков. И grid из 30*256 потоков по кускам (из 30*256 точек) определяет функцию-образ, создавая как раз все те печальные ситуации, о которых писал Toshas.
Касательно содержания статей в журналах, то, думаю, зависит от степени популярности журнала. В суровом научном введение будет меньше, в научно-популярном — шире. Вообще же, бытует мнение, что хорошая научная статья должна во многом быть самодостаточной, во введение рассказывать о некоторых необходимых основах.
Упаковывают… Обычно этим занимается упаковщик у конвейера. То есть, важную роль играет аккуратность рук и правила по упаковке, которым можно обучить практически любого. Перенося пример на Apple, «аккуратность рук» рук сохраняется, а вместо «правил упаковки» можно подставить внимательное наблюдение за «трендами», да и вообще влияние на них самих. Это, конечно, тоже работа, но совсем их другой области, нежели тема поста выше. Сравнивать две фирмы через сравнение этих двух слабо пересекающихся областей несколько ошибочно. Можно, разве что, через приносимую в $ прибыль. То есть, сравнивать успешность компаний. А это уже не то, что пытались сравнить изначально. Это провоцирует потенциально бесконечные холивары. Предположим теперь, что упаковывание — это вовсе не раскладывание по пакетикам и завязывание бантиков, которые сейчас модны, а серьезная, т.н. «креативная», работа: надо много чего проанализировать, подобрать цвета, формы и тому подобное. Но это же и есть производство идей! То есть, на самом деле уже на сравнивают не белое с круглым, вещи из одной области. Но почему-то люди иде по упаковке называют просто упаковкой… Закрадывается подозрение, что они так поступают специально, но неосознанно: черное делают круглым, вызывая проблему сравнения. А потсупают так потому, что понимают — упаковать вовсе не круче, чем придумать.
Лично я считаю, что деятельность в роде работы Apple более культурно и социально зависима, без человека она очень мало значит (в том смысле, что все человеки пропали, а их наследие оценивает кто-то другой). Да и об очередной модели телефона сами люди уже забудут через 5-10 лет. Деятельность же IBM оказывается более фундаментальной, более важной, определяющей ход развития множества областей жизни человека. Результаты этой деятельности более безальтернативны.
Особо отмечу, то запускалось на Tesla с 30-ю SM. Так что, двукратное «увеличение крутизны» при переходе от 30 к 31 блокам понятно: одному SM приходится выполнять два одинаковых блока. Следовательно, он тратит в ~2 раза больше времени. А вся эскадра плывет со скоростью самого медленного корабля…
А тут запускался 1 блок, но число потоков в нем менялось шагом 1. Ступеньки, как я понимаю, обусловлены размерностью warp`ов. Кстати, интересно, что происходит внутри: формально появляются warp`ы с числом потоков меньше 32. Красная пунктирная кривая отражает просто пропорциональную зависимость. Из этого графика становится понятно, что лучше блоки делать не «толще» 256-ти потоков.
2) А как можно удалить обращения к глобальной памяти, если именно там хранится переменная count, инкрементируемая «отработавшими» блоками? Именно она и считывается в циклах. Чем больше потоков, тем больше циклов, тем больше считываюний…
3) Про кратность 32-м и превышение количеством числа ядер говорится в «CUDA C Best Practices Guide». Но я правда не знаю причин тому, и не понимаю, как это следует из Вашего замечания №2. Но в любом случае с интересом ознакомлюсь с Вашей работой! Судя по сайту «Вестника ННГУ», скачать электронный вариант можно. У Вас в каком номере?
4) Как раз по невнимательности подробности обработки условных конструкций пропустил. Но как из того, что сперва выполняется одна ветка, а уже после — другая, может следовать дедлок? Каждый поток threadIdx.x == 0 непременным образом сделает "+1" к счетчику. Поэтому и все threadIdx.x == 0 выйдут из цикла.
Открою страшную тайну: я решил доморощенным образом вычислять преобразование Фурье. Весьма вероятно именно поэтому мое ядро очень прожорливо до регистров, и на каждом SM я запускаю по одному блоку в 256 потоков. И grid из 30*256 потоков по кускам (из 30*256 точек) определяет функцию-образ, создавая как раз все те печальные ситуации, о которых писал Toshas.
Под операцией копирования между (i) и (ii) Вы имеете в виду обмен хост-девайс?
Касательно содержания статей в журналах, то, думаю, зависит от степени популярности журнала. В суровом научном введение будет меньше, в научно-популярном — шире. Вообще же, бытует мнение, что хорошая научная статья должна во многом быть самодостаточной, во введение рассказывать о некоторых необходимых основах.