Свежие комментарии

Empty call на 160 микросекунд - это что-то очень странное. Люди получают цифры в районе 10-30 usec. Подозреваю, что вы неправильно меряете.

С другой стороны, гонять данные на видеокарту, чтобы посчитать 2+2 или даже просто сумму элементов вектора - категорически невыгодно. Вы же этот вектор туда льете со скоростью, ну три гигабайта в секунду, ну может 4 (надо, кстати, на новой материнке померять). А на CPU это будет молотить со скоростью памяти т.е. в разы выше.

Конечно, всякий анроллинг и "рекурсию" делать через макроподстановки (неважно, препроцессором или темплейтером) - удобно.

А что касается CuBLAS - так это паттерн такой. Берем 200 мегабайтные матрицы и умножаем, умножаем, умножаем...

dpp с точкой черного работает вполне прилично, что для недодержаных кадров весьма критично.

Не знаю. Если бы знал, думаешь купил бы я маковский ноут ?
Конвертор был не единственным критерием, но довольно существенным.

Я боюсь, что мой совет по конвертору тебя не устроит. Но в последних статьях он есть (в каждой).

Исходники CuBLAS я тоже смотрел. Тоже не понравилось. Те же макросы.
Может быть каждая из CuBLAS функций и эффективна, но сама концепция порочна. Программу для GPU нельзя собирать как кубики из CuBLAS. Она по определению будет не эффективной.
Оказывается только вызов пустой глобал функции стоит 0.16 мсек. Имхо, нвидиа разработала CuBLAS на скорую руку, причем в Си стиле и предназначает в основном для использования в старых фортран-пакетах.

Не в CUDA, а вообще вычислителями.

Но насколько я помню исходники CuBLAS, там все в макросах (я туда сунул нос и сбежал)

Ну да, такие вещи обычно на препроцессоре делаются, но темплейтер ничем не хуже (считать ли параметризацию с типом и авто-fallback удобнее или нет - дело вкуса)

В тот-то и фокус что не руками. Это все можно поручить компилятору. Вот так

// f = 2^log2(N) ÷
template
struct Exp2Log2{ enum { value = 2*Exp2Log2::value };};
template<>
struct Exp2Log2<1> { enum { value = 1 };};

// RecursivSum, " "
// sdata blockSize
template
__device__ void RecursivSum(__shared__ myType *sdata)
{
if (threadIdx.x < blockSize-blockSizeAligned)
sdata[threadIdx.x] += sdata[threadIdx.x + blockSizeAligned];
__syncthreads();
// instance , blockSize, blockSizeAligned/2:
if (blockSizeAligned >= 2) //
RecursivSum(sdata);
};

// g_idata n, g_odata
template
__global__ void VectorSum(const myType *g_idata, myType *g_odata, const unsigned int n)
{
__shared__ int sdata[blockSize];
sdata[threadIdx.x] = 0;
for (unsigned int i = threadIdx.x; i < n; i += blockSize)
sdata[threadIdx.x] += g_idata[i];
__syncthreads();
RecursivSum::value>(sdata);
if (threadIdx.x == 0) g_odata[0] = sdata[0];
}

Добрый день Алексей,
Затронутая в статье проблема относится к классу проблем редукции. Интересное решение дал Mark Harris в SC07_CUDA_5_Optimization_Harris.pdf.
Развивая эти идеи я пришел к возможности реализации некоторых из методов метапрограммирования в CUDA, в частности удалось пирамиду представить в виде квази-рекурсии (здесь слово квази употреблено поскольку настоящей рекурсии нет, а используются новые instance той же функции). Детали смотрите в http://www.gamedev.ru/code/forum/?id=80317

Я не понимаю, почему 1-2 $(наверное все-таки %) рынка? CUDA работает на всех Geforce 8xxx/9xxx, это что-нибудь в районе трети рынка новых машин (и 2/3 от машин с отдельной видеокартой). Т.е. уже можно.

Что касается перформанса, то умножение матриц на 8800GTX - 100Gflops на их библиотеке, народ вытаскивает до, если не путаю, 170. 2xWoodcrest 5140 (и то же самое C2Q) у меня давали 48 гигафлопсов на single precision.

А дальше все от задачи. Есть класс задач, который ложится на CUDA очень хорошо и там выигрыши в десятки раз. Есть те, которые плохо - там в разы или менее. Причем, то что ляжет на C2Q (возможно параллельное исполнение без блокировок) - ляжет и на Geforce

Не забывай еще, что отдавая что-то считаться на видеокарту - ты разгружаешь CPU и его можно занять чем-то полезным.

Какой троллинг, я просто не понимаю Илью и пытыюсь понять, что ж он доказывает-то. Со статьей та же проблема, но мы от нее ушли.

Илья высказал некоторую мысль про бесполезность усиления ("аппаратного увеличения ISO"), что разница между аппаратным усилением и грамотной софтверной конвертацией с плавающей точкой 0.3 EV в худшем случае, кое-где уже усилитель выкинули и т.д. Я предложил протестировать ISO 100 vs ISO 1600, Илья вроде бы не возражал, но как дошло до дела, после пары итераций на тему "а ты сам проверь" почему-то выплыла тема про сравнение изменения на 1 стоп, да еще и со странным цитированием :) Мы куда идем?

То, что правильный конвертор можем выиграть полстопа, а то и чуть больше перед неправильным - тут и спорить не о чем - верю на слово. Даже тестировать нет смысла - глазом все не так хорошо видно, как на 4 стопах, только время потратим на обсуждение методик тестирования. Только это никак отменяет аппаратный усилитель, который в моей камере аж на 4 стопа усиливает сигнал.

Секундочку, о каких текстах идет речь?

Мы же только что беседовали о бесполезности усилителя перед АЦП и возможностью поднять снимок с ISO 100 -4 EV до эквивалента ISO 1600? :)

> А зачем так много?

Анатолий, там сказано - "На каждый стоп усиления...Разница между умножением с плавающей точкой в raw-сонверторе и аналоговым усилением в камере по замерам даёт 0.3eV по шумам - в худшем известном мне случае." Вот мне и интересно, что конкретно происходит применительно к Вашей камере.

Я Вас прошу не домысливать текстов, и затем не спорить с Вашими же представлениями о том, что авторы хотели сказать.

Что за "горячая экспозиция"?

А зачем так много? Мы ж хотим просто недодержанный ISO 100 поднять до ISO 1600? :)