Так делать нельзя… обнулиться может уже после того, как какие то другие блоки добавят свое значение…
согласен, не подумал об этом, поправил в кодах, добавил UPD
Все равно остались такие же места в temp_scale() и tstat_nose(), только они в конце, а не в начале:
if (!blockIdx.x)
if (!threadIdx.x)
{
md->engKin = md->teKin;
}
Также никто не гарантирует, что этот код в блоке 0 выполнится после всех остальных блоков, скорее даже наоборот, он выполнится до того, как большинство других блоков начнет исполняться и там вычислится:
k = sqrt(md->teKin / md->engKin);
Вообще не стоит ориентироваться на номер блока и выполнять в зависимости от него какой то код — не могу даже придумать случая, где это не было бы ошибкой.
По поводу оптимизаций можно много всего написать. Но мне кажется, здесь имеет смысл оптимизировать только вычисление сил, т.к. оно должно занимать практически все время вычислений. И здесь бы я сделал все таки обработку каждой частицы в отдельном потоке, а ячейки использовал только для поиска соседей. Т.е. в каждом потоке обрабатывается одна частица, но для нее перебираются не все соседи, а только частицы в той ячейке, в которой она находится и в соседних. И избавился бы от atomic'ов — для этого сохранял бы силы только для обрабатываемой частицы — да в этом случае для каждой пары частиц компонент сил будет вычисляться дважды, но это ничто по сравнению с избавлением от atomic'ов. Вычислительные операции на gpu это практически самая дешевая операция — лучше что-то лишний раз посчитать, чем читать из глобальной памяти, а тем более с использованием atomic'ов — одной из наиболее тяжелых операций.
if (blockIdx.x == 0)
{
md->engElecField = 0;
}
...
if (!threadIdx.x) // 0th thread
{
atomicAdd(&md->engElecField, shEngElField);
}
Так делать нельзя. Синхронизации между блоками не предусмотрено, и никто не гарантирует, что первый блок выполнится первым. Соответственно обнулиться может уже после того, как какие то другие блоки добавят свое значение. Соответственно сумма скорее всего вычисляется неправильно. Во втором проходе то же самое.
Обнулять нужно в отдельном kernel ну или через cudaMemset() заранее.
В put_periodic():
atomicAdd(&md->specAcBox[type].y, 1);
y — это описка? Везде вроде x.
Входной параметр step равен количеству частиц (атомов) обрабатываемых одним потоком.
Очень плохая идея. Если step > 1, получается плохой патерн доступа к памяти в warp'е — потоки обращаются к памяти не последовательно, а с шагом в step, т.е. warp будет читать не один кусок памяти целиком, а несколько маленьких кусочков — по одному на каждый поток. Это может замедлять выполнение в несколько раз.
В чем смысл то? Если классическим образом каждый символ кодировать таким же количеством бит, сколько уровней в дереве, по объему памяти получится то же самое, а все операции будут за O(1).
В данном алгоритме в любом случае будет использована вся таблица, но коллизий будет больше. В стандартных алгоритмах, будет использоваться не вся.
Это был просто пример данных, на которых такой подход — использование напрямую значения ключа в качестве хэша, будет работать неэффективно. Причем это достаточно реалистичный пример — например, такое может быть, если ключи являются адресами в памяти.
В таком случае возможное значение ключа будет вообще только одно.
Я писал о другом, причем в предположении, что размер хэш-таблицы — степень двойки, как здесь.
Допустим размер хэш таблицы 16, тогда для четных ключей возможны только 8 значений хэшей: 0, 2, 4, 8, 10, 12, 14, т.е. будет использоваться только половина таблицы. Если ключи кратны 4, то хэшей будет только 4 и использовать 1/4 таблицы и т.д.
Ну так если у нас, например, все ключи четные или кратны какому то еще числу, являющемуся степенью двойки? Тогда и хэши будут кратны этому числу и будет все печально.
Ну так скорость зависит от реализации доступа, а не от способа хранения. Для тех же 5 трит на 8 бит, триты можно доставать делением, а можно через lookup-таблицу. Понятное дело, что делением будет скорее всего в разы медленнее чем через lookup-таблицу. А способ хранения более оптимальный вряд ли можно придумать. Чисто математически понятно, что больше не выжать.
Если почитать про реализацию, по приведенным ссылкам (конкретно — в BitMagic), то можно увидеть, что за O(1) находится блок из 65536 бит, в котором потом осуществляется дальнейший поиск. Накладные расходы по памяти у них на один такой блок кроме самих бит блока: 4 байта (это для 32-битных индексов, для 64-битных должно быть соответственно 8) + 2 * 2 байта для ускорения поиска в самих блоках.
Задача понятна. Но генерируемая таким образом последовательность (назовем ее X), не является последовательностью, генерируемой указанным ГПСЧ, а является последовательностью, полученной из первых чисел последовательностей, генерируемых разными ГПСЧ, инициализированными случайным образом. Соответственно нельзя утверждать, что X обладает свойствами последовательностей, генерируемых данным ГПСЧ. Тогда какой смысл в его использовании вообще?
Я думаю, если сгенерировать тестовую последовательность описанным способом (без сохранения состояния), то указанный в статье тест она пройдет не лучше, чем последовательность, сгенерированная хэш функцией, используемой для генерации состояний. Поэтому я и написал, что проще использовать сразу результат хэш функции. Хотя возможно у этой хэш функции случайность вообще плохая и присутствуют какие то явные патерны, тогда первое число, генерируемое ГПСЧ, действительно может обладать большей случайностью, но все же вся последовательность не будет обладать свойствами, обеспечиваемыми этим генератором.
Насколько я понимаю, при использовании этого (как и многих других) ГПСЧ, каждый вызов next() должен использовать состояние, сохраненное предыдущим вызовом. Только в этом случае можно говорить о какой то случайной последовательности и, соответственно, ее свойствах.
Здесь же при каждом вызове RND генерируется новое состояние с помощью этого макроса SEED, которое потом никуда не сохраняется и соответственно при следующем вызове xorshift256_next() никак не используется. Т.е. каждое новое случайное число определяется новым SEED и зависит от него, а не от свойств ГПСЧ.
А какой смысл использования ГСПЧ, если функция xorshift256_next не сохраняет состояние? Т.е. получается, что каждый следующий вызов RND никак не зависит от предыдущего — при каждом новом вызове просто инициализируется новый ГПСЧ, который генерирует одно число. Не проще ли напрямую использовать хэш от __COUNTER__ и __TIME__?
Имеются ввиду процессоры, обладающие кэшами различного уровня и конвейером.
Линейный доступ выигрывает из-за эффективного использования кэша за счет префетчинга и из-за отсутствия сбоев конвеера за счет эффективного предсказания переходов. Но если идет поиск постоянно в одном наборе данных небольшого размера, то скорее всего он и так будет закэширован.
Если речь идет о простых типах, тех же интах, то, как я написал, тут выигрыш у бинарного поиска начинается где то на сотнях элементов. Логарифм 100 — будет в районе 7, т.е. сравнений будет как минимум раз в 10 меньше при бинарном поиске, для 200 меньше уже в 20 раз и т.д. Т.е. линейный поиск должен быть во столько раз эффективнее на одно сравнение чтобы выигрывать. Ну видимо где то в этом районе сейчас граница и проходит. Потому что по опыту получаются именно такие размерности, когда бинарный поиск начинает выигрывать.
Если же говорить о более сложных типах — строках или каких то классах, имеющих свою сложную функцию сравнения, то само время выполнения функции сравнения становится уже сравнимым с оверхедом от отсутствия кэширования/сбоев конвейера и эффективность линейного поиска на одну операцию сравнения становится выше уже не в десятки, а в несколько раз или даже просто на несколько процентов. Поэтому в этом случае граница и смещается в область десятков элементов, а не сотен.
Бинарный поиск на современных процессорах выигрывает уже на сотнях элементов. Это для простых типов. При увеличении размера элемента или увеличении сложности функции сравнения, бинарный поиск становится еще эффективнее. В общем случае бинарный поиск выгоднее уже на десятках элементов.
Ну SDK можно использовать любой. Если реально в коде используется только 1.1, то проблем не должно быть. А если используются какие то вызовы из 1.2, то тут уже никакой SDK не поможет, если драйвер эту версию не поддерживает.
У Nvidia доступна вся память, но одним куском можно выделить 2G максимум, это да. Через CUDA такого ограничения нет.
У меня со сборкой собственного OpenCL кода под винду проблем нет. Использую OpenCL SDK, который идет в составе CUDA. А вот если нужно собрать что то стороннее, не знаю. Но например libx264 с поддержкой кодирования в OpenCL у меня вроде собирался под виндой из коробки, но они там вообще сторонний SDK не используют. Как я понял, просто сами загружают OpenCL.dll и находят там все нужные символы.
Ну у меня сегодня не получилось проверить. Зарегистрировался, написано что тестовый доступ для конфигураций с ценой меньше 3000р. Пока игрался с конфигурацией, чтобы уложиться в эту цену, не заметил, что видеокарта пропала. В итоге создал без видеокарты. Когда заметил, удалил, но новую с тестовым доступом создать уже не получилось. Завтра может попрошу кого-нибудь зарегистрировать новый аккаунт и проверить.
Но в любом случае, даже если там 2G видеопамяти всего, этого все равно мало. Даже для десктопного решения это не много, а для серверного точно не вариант. Ну а если выделение памяти там реально не контролируется, то непонятно как этим пользоваться — в любом момент кто то другой может использовать всю видеопамять и ты об этом никак не узнаешь.
Потестировать конечно можно, чтобы понять будет ли работать там в принципе. Но с максимумальным объемом видеопамяти в 1Gb использовать такой сервер на практике бессмысленно. 1Gb это объем самых бюджетных видеокарт и на них в нашей задаче как раз в основном все упирается в видеопамять.
Все равно остались такие же места в temp_scale() и tstat_nose(), только они в конце, а не в начале:
Также никто не гарантирует, что этот код в блоке 0 выполнится после всех остальных блоков, скорее даже наоборот, он выполнится до того, как большинство других блоков начнет исполняться и там вычислится:
Вообще не стоит ориентироваться на номер блока и выполнять в зависимости от него какой то код — не могу даже придумать случая, где это не было бы ошибкой.
По поводу оптимизаций можно много всего написать. Но мне кажется, здесь имеет смысл оптимизировать только вычисление сил, т.к. оно должно занимать практически все время вычислений. И здесь бы я сделал все таки обработку каждой частицы в отдельном потоке, а ячейки использовал только для поиска соседей. Т.е. в каждом потоке обрабатывается одна частица, но для нее перебираются не все соседи, а только частицы в той ячейке, в которой она находится и в соседних. И избавился бы от atomic'ов — для этого сохранял бы силы только для обрабатываемой частицы — да в этом случае для каждой пары частиц компонент сил будет вычисляться дважды, но это ничто по сравнению с избавлением от atomic'ов. Вычислительные операции на gpu это практически самая дешевая операция — лучше что-то лишний раз посчитать, чем читать из глобальной памяти, а тем более с использованием atomic'ов — одной из наиболее тяжелых операций.
Так делать нельзя. Синхронизации между блоками не предусмотрено, и никто не гарантирует, что первый блок выполнится первым. Соответственно обнулиться может уже после того, как какие то другие блоки добавят свое значение. Соответственно сумма скорее всего вычисляется неправильно. Во втором проходе то же самое.
Обнулять нужно в отдельном kernel ну или через cudaMemset() заранее.
В put_periodic():
y — это описка? Везде вроде x.
Очень плохая идея. Если step > 1, получается плохой патерн доступа к памяти в warp'е — потоки обращаются к памяти не последовательно, а с шагом в step, т.е. warp будет читать не один кусок памяти целиком, а несколько маленьких кусочков — по одному на каждый поток. Это может замедлять выполнение в несколько раз.
Ну и это только то, что заметил в первой части :)
Это был просто пример данных, на которых такой подход — использование напрямую значения ключа в качестве хэша, будет работать неэффективно. Причем это достаточно реалистичный пример — например, такое может быть, если ключи являются адресами в памяти.
Я писал о другом, причем в предположении, что размер хэш-таблицы — степень двойки, как здесь.
Допустим размер хэш таблицы 16, тогда для четных ключей возможны только 8 значений хэшей: 0, 2, 4, 8, 10, 12, 14, т.е. будет использоваться только половина таблицы. Если ключи кратны 4, то хэшей будет только 4 и использовать 1/4 таблицы и т.д.
Я думаю, если сгенерировать тестовую последовательность описанным способом (без сохранения состояния), то указанный в статье тест она пройдет не лучше, чем последовательность, сгенерированная хэш функцией, используемой для генерации состояний. Поэтому я и написал, что проще использовать сразу результат хэш функции. Хотя возможно у этой хэш функции случайность вообще плохая и присутствуют какие то явные патерны, тогда первое число, генерируемое ГПСЧ, действительно может обладать большей случайностью, но все же вся последовательность не будет обладать свойствами, обеспечиваемыми этим генератором.
Здесь же при каждом вызове RND генерируется новое состояние с помощью этого макроса SEED, которое потом никуда не сохраняется и соответственно при следующем вызове xorshift256_next() никак не используется. Т.е. каждое новое случайное число определяется новым SEED и зависит от него, а не от свойств ГПСЧ.
Линейный доступ выигрывает из-за эффективного использования кэша за счет префетчинга и из-за отсутствия сбоев конвеера за счет эффективного предсказания переходов. Но если идет поиск постоянно в одном наборе данных небольшого размера, то скорее всего он и так будет закэширован.
Если речь идет о простых типах, тех же интах, то, как я написал, тут выигрыш у бинарного поиска начинается где то на сотнях элементов. Логарифм 100 — будет в районе 7, т.е. сравнений будет как минимум раз в 10 меньше при бинарном поиске, для 200 меньше уже в 20 раз и т.д. Т.е. линейный поиск должен быть во столько раз эффективнее на одно сравнение чтобы выигрывать. Ну видимо где то в этом районе сейчас граница и проходит. Потому что по опыту получаются именно такие размерности, когда бинарный поиск начинает выигрывать.
Если же говорить о более сложных типах — строках или каких то классах, имеющих свою сложную функцию сравнения, то само время выполнения функции сравнения становится уже сравнимым с оверхедом от отсутствия кэширования/сбоев конвейера и эффективность линейного поиска на одну операцию сравнения становится выше уже не в десятки, а в несколько раз или даже просто на несколько процентов. Поэтому в этом случае граница и смещается в область десятков элементов, а не сотен.
У меня со сборкой собственного OpenCL кода под винду проблем нет. Использую OpenCL SDK, который идет в составе CUDA. А вот если нужно собрать что то стороннее, не знаю. Но например libx264 с поддержкой кодирования в OpenCL у меня вроде собирался под виндой из коробки, но они там вообще сторонний SDK не используют. Как я понял, просто сами загружают OpenCL.dll и находят там все нужные символы.
Но в любом случае, даже если там 2G видеопамяти всего, этого все равно мало. Даже для десктопного решения это не много, а для серверного точно не вариант. Ну а если выделение памяти там реально не контролируется, то непонятно как этим пользоваться — в любом момент кто то другой может использовать всю видеопамять и ты об этом никак не узнаешь.