Здравствуйте! У меня вопрос немного не по теме, но я уже честно попробовал много разных способов получить ответ, в том числе официальные TBB licensing FAQ, Intel support mailing list, TBB forum, StackOverflow; пока безрезультатно.
Я разрабатываю приложение для компании. Приложение с закрытым кодом, но бесплатное, и не предполагает подсчет пользователей, какие-либо манипуляции с регистрацией и прочее. Существует ли вид лицензии Intel TBB, который компания могла бы приобрести, что позволило бы мне использовать TBB в моем приложении?
Все виды коммерческих лицензий, которые я видел до сих пор, упоминают per-user, или per-LAN модель, но платить $xxx за пользователя бесплатного приложения нас не устраивает, поэтому суть вопроса сводится вот к чему: правильно ли я понял коммерческую модель описанную выше, и если да, то есть ли возможность выкупить лицензию для неограниченного распространения одного программного продукта?
Гн Herb Sutter в своем блоге много пишет про многопоточность, атомики, и все это в свете c++11. А что делать тем, кто по каким-либо причинам не может перейти на c++11, и вынужден пользоваться c++ — может быть есть публично доступные библиотеки/врапперы, которые реализуют атомики по аналогии с std::atomic?
Только сегодня скачал триал и уже нашел потенциальную гонку данных. У меня вопрос к автору: какой инструмент (от intel?) мог бы помочь определить причину следующего поведения: есть приложение серверного типа, которое на старте потребляет 30% CPU. Далее нагрузка на процессор растет линейно, хотя нагрузка по количеству данных и соединений остается равномерной, и уже через час-два процессор не выдерживает. Утечек памяти нет, проверенно уже — какой инструмент или комбинация может помочь обнаружить причину? спасибо!
Учитывая состав аккумуляторов, и довольно нередкие сообщения о взрывах аккумуляторов, хотелось бы предположить, что он спрятан подальше от глаз. Кстати — вопрос к автору! =)
Если большая часть итеративного процесса может выполняться на регистрах и разделяемой памяти, то механизм синхронизации разбиением на кернелы может внести избыточные задержки на синхронизацию через глобальную память (очень дорогая она!).
Из вашего описания я понял, что ваш подход чем-то напоминает семпл reduce, когда каждый блок считает частичную сумму, потом один из блоков ждет когда все завершатся (читая значение атомарного счетчика выполненных блоков), и суммирует частичные суммы. Такой подход не вводит зависимостей между блоками (если реализован в точности как в семпле), и может сэкономить много времени.
В CUDA API неспроста отсутствует встроенный примитив синхронизации всех блоков (тредов) грида. Связано это с тем, что грид может (и должен, согласено Best Practices Guide) содержать блоков больше, чем SM-ов. Всвязи с этим, все неявные алгоритмы, которые вводят зависимости по данным между независимыми блоками (а значит, и SM-ами), могут приводить к простою, и снижению производительности.
Но и это еще не все. Гораздо хуже ситуация может случиться при определенных типах зависимостей, где все SM-ы получат по блоку на исполнение, и будут ждать некоторое условие, которое может произойти, а может и нет, в зависимости от того, в каком порядке SM-планировщик будет получать работу. Такие сценарии крайне тяжело отлаживать, поэтому рекомендуется отдавать предпочтение самым простым и доступным способам синхронизации.
Касаемо разделения (I) & (II) на независимые кернелы, это не самый плохой вариант. При использовании асинхронного API и cuda-канала, отличного от нуля, оверхед на запуск кернела будет только в самом первом запуске. Все остальные кернелы будут запланированы из внутренней очереди в драйвере. Это конечно при условии, что вам не требуется между (I) и (II) вставлять операции копирования памяти.
В заключение, отмечу, что в CUDA 5.0 и в самой продвинутой карте семейства Kepler доступна технология запуска кернела _изнутри_ другого кернела. Это в разы увеличивает гибкость программирования на CUDA, но к сожалению, доступность железа пока отстает.
Рекомендую не опираться на алгоритмы основанные на гистерезисе без лишней надобности (non-maximum suppression в canny edge detector'е).
Хоть вы и говорите что ваши объекты могут «быть совершенно непохожими на окружности», попробуйте сделать так:
1. Возьмите экземпляр объекта высокого разрешения (шаблон), примените к нему преобразование (о нем позже)
2. Для каждого квадратного окна входного изображения А. уменьшите его до размера шаблона и Б. посчитайте над ним тоже самое преобразование, а затем корреляцию результата и преобразованного шаблона. Если корреляция велика, то перед вами экземпляр объекта
Преобразование может быть например таким:
1. Перевод RGB-Luma
2. OutPix = dx*dx + dy*dy, где dx — это результат применения оператора Собеля по горизонтали (например с ядром 3х3), а dy — по вертикали. Таким образом в каждом пикселе будет записан квадрат производной по направлению, т.е. грубо говоря сила границы.
Есть такое demo compo — TMDC (Text Mode Demo Compo). Туда принимают работы, которые для вывода используют только консоль. Основной девиз — «Can you make textmode look good?». Больше всего мне понравилась одна из invitro tmdc5: www.youtube.com/watch?v=aYyz6FphY04
У них на сайте залежи демок. К сожалению, компо сейчас чуть менее чем совсем мертвое, потому что все присланные работы занимают места в top 10 в порядке убывания зрелищности. Правда, возможно, благодаря таким топикам интерес к этой части сцены будет подогрет =)
Все же оптимизацию CUDA-программы надо начинать не с конфликтов банков, а с определением паттерна доступа к глобальной памяти. Затем приоритетным этапом является определение мест нерационального бранчинга, и только после этого можно сосредоточиться на тюнинге на уровне регистров, банков и их конфликтов. Это потому, что если вы вычистите все конфликты, но не будут соблюдены правила объединенных запросов к глобальной памяти, то быстродействие программы будет отличаться от максимально-возможного в несколько раз.
Теперь немного замечаний по тексту:
1. Конфликт банков на Fermi гораздо сложнее вызвать, особенно при работе с маленькими типами char и short. Можно любым числом потоков адресовать один банк (разные его байты, но в рамках одного слова).
2. При необходимости обрабатывать один байт на поток на архитектурах до Fermi можно использовать т.н. bit-twiddling hack, который заключается в подмене threadIdx.x на такую пермутацию, которая позволяет обходить конфликт банков. Идея заключается в произведении циклического сдвига в младших 4 (Например для пермутации линейного блока из 64 потоков в группы по 16):
3. Счетчик warp serialize показывает именно количество сериализаций варпов, случившихся в железе по факту исполнения. Но складывается он не только из конфликта банков. Например, любое ветвление (и в частности те, про которые пишется в branching и divergent branching) вызывает одну сериализацию. Также есть менее значительные (подвластные программисту) явления, вызывающие нарастание этого счетчика. Вообще, счетчики профилировщика рекомендуется оценивать в динамике, а не конкретные их значения. Т.е. лучще уменьшать плохие счетчики (uncoalesced, divergent branch) и увеличивать хорошие (occupancy, coalesced, cache hit rate)
Вот бы была гуглокарта, на которой чекбоксами можно смотреть зоны покрытия — сильно пригодилось бы для дачников. Сколько ни ездил на дачу к друзьям в ближнее подмосковье — ни Йота ни Скайлинк не ловили…
Всегда хотел узнать как формализуется решение задачи поиска по фразе как в после — «Столица островного государства на улице имени одного из аэропортов столицы»
Первое что приходит на ум — искать все островные гос-ва, выписывать в список их столицы, дальше аэропорты, потом искать пересечения в списках…
Я разрабатываю приложение для компании. Приложение с закрытым кодом, но бесплатное, и не предполагает подсчет пользователей, какие-либо манипуляции с регистрацией и прочее. Существует ли вид лицензии Intel TBB, который компания могла бы приобрести, что позволило бы мне использовать TBB в моем приложении?
Все виды коммерческих лицензий, которые я видел до сих пор, упоминают per-user, или per-LAN модель, но платить $xxx за пользователя бесплатного приложения нас не устраивает, поэтому суть вопроса сводится вот к чему: правильно ли я понял коммерческую модель описанную выше, и если да, то есть ли возможность выкупить лицензию для неограниченного распространения одного программного продукта?
Спасибо!
Из вашего описания я понял, что ваш подход чем-то напоминает семпл reduce, когда каждый блок считает частичную сумму, потом один из блоков ждет когда все завершатся (читая значение атомарного счетчика выполненных блоков), и суммирует частичные суммы. Такой подход не вводит зависимостей между блоками (если реализован в точности как в семпле), и может сэкономить много времени.
Но и это еще не все. Гораздо хуже ситуация может случиться при определенных типах зависимостей, где все SM-ы получат по блоку на исполнение, и будут ждать некоторое условие, которое может произойти, а может и нет, в зависимости от того, в каком порядке SM-планировщик будет получать работу. Такие сценарии крайне тяжело отлаживать, поэтому рекомендуется отдавать предпочтение самым простым и доступным способам синхронизации.
Касаемо разделения (I) & (II) на независимые кернелы, это не самый плохой вариант. При использовании асинхронного API и cuda-канала, отличного от нуля, оверхед на запуск кернела будет только в самом первом запуске. Все остальные кернелы будут запланированы из внутренней очереди в драйвере. Это конечно при условии, что вам не требуется между (I) и (II) вставлять операции копирования памяти.
В заключение, отмечу, что в CUDA 5.0 и в самой продвинутой карте семейства Kepler доступна технология запуска кернела _изнутри_ другого кернела. Это в разы увеличивает гибкость программирования на CUDA, но к сожалению, доступность железа пока отстает.
if (threadIdx.x < halfData) sMins[j] = (sMins[j] < sMins[j+halfData]) ? sMins[j] : sMins[j+halfData]; __syncthreads();
if (threadIdx.x < halfData/2) sMins[j] = (sMins[j] < sMins[j+halfData/2]) ? sMins[j] : sMins[j+halfData/2]; __syncthreads();
//...
//после того как число в условии станет 32 можно уже не делать __syncthreads
if (threadIdx.x < 32) sMins[j] = (sMins[j] < sMins[j+32]) ? sMins[j] : sMins[j+32];
if (threadIdx.x < 16) sMins[j] = (sMins[j] < sMins[j+16]) ? sMins[j] : sMins[j+16];
if (threadIdx.x < 8) sMins[j] = (sMins[j] < sMins[j+8]) ? sMins[j] : sMins[j+8];
if (threadIdx.x < 4) sMins[j] = (sMins[j] < sMins[j+4]) ? sMins[j] : sMins[j+4];
if (threadIdx.x < 2) sMins[j] = (sMins[j] < sMins[j+2]) ? sMins[j] : sMins[j+2];
if (threadIdx.x < 1) sMins[j] = (sMins[j] < sMins[j+1]) ? sMins[j] : sMins[j+1];
результат в sMins[0]
удачи, отличная статья!
Хоть вы и говорите что ваши объекты могут «быть совершенно непохожими на окружности», попробуйте сделать так:
1. Возьмите экземпляр объекта высокого разрешения (шаблон), примените к нему преобразование (о нем позже)
2. Для каждого квадратного окна входного изображения А. уменьшите его до размера шаблона и Б. посчитайте над ним тоже самое преобразование, а затем корреляцию результата и преобразованного шаблона. Если корреляция велика, то перед вами экземпляр объекта
Преобразование может быть например таким:
1. Перевод RGB-Luma
2. OutPix = dx*dx + dy*dy, где dx — это результат применения оператора Собеля по горизонтали (например с ядром 3х3), а dy — по вертикали. Таким образом в каждом пикселе будет записан квадрат производной по направлению, т.е. грубо говоря сила границы.
У них на сайте залежи демок. К сожалению, компо сейчас чуть менее чем совсем мертвое, потому что все присланные работы занимают места в top 10 в порядке убывания зрелищности. Правда, возможно, благодаря таким топикам интерес к этой части сцены будет подогрет =)
Теперь немного замечаний по тексту:
1. Конфликт банков на Fermi гораздо сложнее вызвать, особенно при работе с маленькими типами char и short. Можно любым числом потоков адресовать один банк (разные его байты, но в рамках одного слова).
2. При необходимости обрабатывать один байт на поток на архитектурах до Fermi можно использовать т.н. bit-twiddling hack, который заключается в подмене threadIdx.x на такую пермутацию, которая позволяет обходить конфликт банков. Идея заключается в произведении циклического сдвига в младших 4 (Например для пермутации линейного блока из 64 потоков в группы по 16):
__device__ DEVICEINLINE int permuteThreads8u(int x)
{
return (x >> 4) + ((x & 0xF) << 2);
}
3. Счетчик warp serialize показывает именно количество сериализаций варпов, случившихся в железе по факту исполнения. Но складывается он не только из конфликта банков. Например, любое ветвление (и в частности те, про которые пишется в branching и divergent branching) вызывает одну сериализацию. Также есть менее значительные (подвластные программисту) явления, вызывающие нарастание этого счетчика. Вообще, счетчики профилировщика рекомендуется оценивать в динамике, а не конкретные их значения. Т.е. лучще уменьшать плохие счетчики (uncoalesced, divergent branch) и увеличивать хорошие (occupancy, coalesced, cache hit rate)
Первое что приходит на ум — искать все островные гос-ва, выписывать в список их столицы, дальше аэропорты, потом искать пересечения в списках…
Folding@Home,
vReveal,
+
на офсайте есть каталог с кучей коммерческих прог
developer.download.nvidia.com/pix/tools/Compute_Debugger.png
developer.download.nvidia.com/pix/tools/Compute_Analyzer.png
developer.download.nvidia.com/pix/tools/Graphics_Debugger.png
developer.download.nvidia.com/pix/tools/Graphics_Inspector.png