NVidia GTX 280, Tesla T10P

Как и обещали, публике предъявили новый чипсет и новые видеокарты на нем. Нас они интересуют не с точки зрения графики, а с точки зрения вычислений, так их и рассмотрим.

GeForce GTX 280

В пресс-релизах пишут про примерно терафлопс, но внутренний голос мне подсказывает, что это такая же наколка, как 500 гигафлопс у 8800GTX - это посчитано для каких-то графических операций, а реальную производительность для вычислений считали как "1 MAD (2 операции) на такт внутренней частоты на процессор". Для 8800GTX получалось 128*2*1.35GHz = 345.6 GFLOP/s, реально получить удавалось до 205 на SGEMM, а на какой-то вычислительной химии получали практически теоретическую производительность (за 300 GFLOP/s). Для 280 GTX получается, по той же формуле, 240 * 2 * 1.296 = 622 GFLOP/s.

На реальных задачах прирост производительности запросто может быть в эти самые два раза, ибо memory bandwidth тоже выросла почти вдвое, что приятно.

Двойная точность

Двойная точность объявлена, каких-то бенчмарок еще нет. На CNET пишут про 90 GFLOP/s на fp64, что не очень убедительно. Сама цифра - хорошая, только вот 300-долларовый 4-ядерный процессор от Intel показывает примерно вдвое меньше, а карта - вдвое дороже.

Впрочем, рекомендуемая сейчас техника - получить приближенное решение в одинарной точности, а потом его уточнить в двойной - теперь может применяться и прямо на GPU, без пересылки промежуточного результата обратно с карты. Эта техника интересна и для обычных CPU тоже, там одинарная точность примерно вдвое быстрее и формально и по факту, ведь упираемся мы в bandwidth памяти.

Про поддержку CUDA и про драйвера я пока не понимаю. Не видел я, чтобы была заявлена совместимость с новыми картами, но может быть плохо смотрел. Понятно, что CUDA 2.0 готовили именно под эти карты.

Вердикт: надо щупать живьем.

Tesla T10P

Кроме игровых карт, анонсированы и научные. Как карта (C1060), так и сервер (S1070), доступность осенью этого года.

Из опубликованых спецификаций ясно не очень много. В первом приближении это тот же чипсет (один в карте, 4 в сервере), памяти в нем 4 гигабайта на чип, вместо гигабайта в игровой версии (а ведь я уже предсказал видеокарты с 64-битной адресацией, недолго осталось ждать).

Интересно то, что 4 гигабайта быстрой памяти сделать, похоже, нельзя, поэтому в Тесле пропускная способность памяти в 1.4 раза меньше (и для многих приложений это аукнется)

Tags: 

Comments

90 dFLOPS CNET криво

90 dFLOPS CNET криво посчитал. реально 3(dFPU)*10(TC)*1,3(freq)*2(MAD)=78 dFLOPS
думаю реально будет меньше, т.к. архитектура заточена читать из памяти и регистров по 32 бита, плюс double могут "обидеть" в pipeline и тп, ведь для графики он не нужен. но всё равно, даже в 10 раз медленее чем float он полезен для mixed precision схем, как вы упомянули.

С чтением из общей памяти там

С чтением из общей памяти там все более-менее нормально. float2 конечно помедленнее, но терпимо.

Вот с shared - не могу сообразить, но вроде тоже особо большой беды нет.

классический пример проблемы

классический пример проблемы - AMD K8. в 128 бит SSE регистры читает медленно, 1 за 2 такта. а в 64 бит MMX регистры быстро, 2 за 1 такт, т.е. в 2 раза быстрее.
если ширина шин регистры - кэш сделана под 32 бита, то 64-бит слово может читаться медленее, чем 2 по 32 бит.

Для float4 оно действительно

Для float4 оно действительно сильно хуже (я мерял). Для float2 - мне помнится, что я тоже мерял и получил что-то хорошее, надо конечно перемерять

Может иметь смысл мерять

Может иметь смысл мерять чтение, запись и копирование раздельно. Лично у меня получаются весьма занимательные результаты.

Здравствуйте, Насколько я

Здравствуйте,
Насколько я понял вместо SHARED можно использовать регистровую память
например чтобы обработать обменные граничные условия с разных тредов одного блока (или даже другого блока) сразу в регистровой памяти не переходя в SHARED ??

обменивать данные с разных

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

Пример

Здравствуйте

есть еще такой вопрос

возможна ли такая комбинация

i = blockIdx.x * BlockDim.x + threadIdx.x

C[ i ] = A[ i -1 ] + B[ i+1 ];
для всех i пренадлежащих одному блоку

(Если нет) Как лучше считать такое выражение?;
1. считать смежные ячейки все время в SHARED ?
2. или постоянно загружать в глобальную память после обмена в SHARED и считать там?
N порядка 100 000 double complex

3. уменьшить обмены таким образом

for(j=0 ; j < M ; j++ ) {
C[ i + j ] = A[ i - 1 + j ] + B[ i + 1 + j ];
}
т.е. на каждом треде последовательно выполняется часть сетки, а не одна ячейка.

спасибо

хотелось бы узнать немного

хотелось бы узнать немного больше о задаче. А и B --- два различных массива? почему бы их не сдвинуть тогда и переписать алгоритм как C[i] = A[i]+B[i]? Или всё же у вас конечные разности?

A[ i ] и B[ i ] различные

A[ i ] и B[ i ] различные массивы в расчетах конечных разностей, каждому i соответствует свой тред
но в том и дело что этот сдвиг нужно сделать очень быстро
ведь производится много итераций

еще работающей программы

еще работающей программы нет
но собираюсь ее написать
межблочный - только один выход
через глобальную память, но данных мало, только с краев блоков 1D
еще думал чтобы использовать только один блок
создать 512 тредов
раскидать данные до 1 000 000 ячеек по одному блоку потоков
всего на каждый тред по 2 обменной ячейки
т.е. 512*2 в SHARED должно поместиться
другое дело что это не эффективное использование получается
и говорят что это не эффективно даже относительно одного мультипроцессора
создать на нем один блок потоков.

Если у вас миллион точек, то

Если у вас миллион точек, то все данные в shared не влезут (если данные - одно complex-число, то на миллион точек нужно 8 мегабайт, а даже на 280GTX разделяемой памяти на все мультипроцессоры - меньше мегабайта), а значит на каждом шаге по времени вы вынуждены их сохранять в глобальную память (а на следующем шаге - оттуда читать).

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

"Чтобы прятать - блоков

"Чтобы прятать - блоков должно быть много больше, чем мультипроцессоров." <- расхожее, но не совсем точное выражение. Например, на практике приходилось наблюдать близкую к пиковой производительности при всего лишь 2х блоках на мультипроцессор. И это как раз в 7-точечном конечно-разностном шаблоне, где обращений к памяти относительно много. В FFT же или GEMM для достижения максимальной производительности достаточно 3х-4х блоков на мультипроцессор. Сильно больше имеет смысл заводить лишь для равномерного распределения нагрузки.

Спасибо приступаю к

Спасибо
приступаю к работе
думаю через пару месяцов
будут интересные анализы производительности
расскажу

возможно ли следующая

возможно ли следующая операция и является ли это опрерацией с
регистрами или это операции в глобальной памяти:

если массивы A,B находятся в глобальной памяти

B[ i ] = A[ i-1 ] + A[ i ] + A[ i +1 ] если i = blockDim.x
*blockIdx.x + threadIdx.x
будет ли эта операция параллельна
или возможно сложение только двух элементов из глобальной или общей памяти
например
B[ i ] = A[ i-1 ] + A[ i ]
если это возможно, происходят ли эти операции в регистрах и зачем
тогда использовать общую память?

Возможна. Сначала A[ i-1 ],

Возможна. Сначала A[ i-1 ], A[ i] и A[ i +1 ] загрузятся из глобальной памяти в регистры, потом сложатся, потом запишутся в глобальную память.

проще ли загружать из

проще ли загружать из глобальной памяти (её еще и больше чем общей) в регистры
с помощью
i = blockDim.x*blockIdx.x + threadIdx.x

B[ i ] = A[ i-1 ] + A[ i ] + A[ i +1 ]

или быстрее будет загрузить в общую память и там за несколько (если не
хватит памяти) циклов общитать
B[ threadIdx.x ] = A[ threadIdx.x - 1 ] + A[ threadIdx.x ] + A[ threadIdx.x +1 ]
(т.е. получается что скорость загрузки из общей памяти в регистры быстрее)

или например самому при общете в общей памяти переписать как
B[ threadIdx.x ] = A[ threadIdx.x - 1 ] + A[ threadIdx.x ]
B[ threadIdx.x ] + = A[ threadIdx.x + 1 ]

т.е. отличаются ли скорости расчетов в глобальной и общей памяти
или только скорости загрузки

одинакова ли скорость загрузки из глобальной в общую, из общей в глобальную.

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

простите за спам :)

конечно, лучше сначала

конечно, лучше сначала загрузить A в shared memory, потом считать. В любом случае время выполнения программы будет ограничено пропускной способностью памяти (ПСП). Если A читать один раз, а не три, то потребление ПСП будет в 2 раза меньше. Соответственно, работать будет до 2х раз быстрее. Вопрос в том, как обрабатывать точки на границе блока.

Хранить B в shared memory смысла не вижу.

почему бы не хранить в памяти

почему бы не хранить в памяти уже в сдвинутом виде? тогда в этом кернеле все доступы к памяти будут выровнены, будет работать на пике пропускной способности памяти. Иначе --- на GTX280 будет работать в полтора---два раза медленнее пика.

я так понимаю есть какие-то другие компоненты приложения которые наоборот будут тормозить если массивы в памяти сдвинуть?

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

регистры vs shared

Насколько я понимаю, передать данные через регистры между тредами не получится никак. Поэтому регистры - для локальных переменных, shared memory - для обмена внутри блока, global - для обмена между блоками.

В форуме нвидии выложены

В форуме нвидии выложены ссылки на Cuda 2.0 beta 2, там есть ссылки на доки и спецификации.

Из всего того, что я прочёл, так и не понял, поменялось ли что-то в собственно эффективности железа и отношению чипа к ветвлению в кернеле. Вдвое больше регистров, 240 cores - это всё замечательно; double при просадке скорости в 10 раз - академически любопытно; но, если кроме вышеназванного ничего принципиально не поменялось, то новости не такие хорошие. Огорчает, что объём shared memory не поменялся, 16K не для всех задач комфортны.

никто и не собирается дальше

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

Эффективность по ветвлениям

Эффективность по ветвлениям скорее всего сильно не поменялась.

А с shared - я не думаю, что это будет меняться, на это уже заточены все программы, что теперь переписывать отдельно под G80, отдельно под G200 ?
Вот register file увеличивать можно почти прозрачно.

если нужна высокая

если нужна высокая производительность, shared memory лучше вообще использовать поменьше. Лучшие алгоритмы работают через регистры, на GT200 их под 2MB.

Про регистры

Я несколько упустил, как именно надо писать кернел для того, чтобы он жил только в регистрах ?

У меня например задача - мне нужен стек, причём раздельный стек для каждого треда в блоке. Я его завожу в shared mem ... есть другой путь ?

GPU имеет SIMD архитектуру,

GPU имеет SIMD архитектуру, что значит оперирует с векторами, а не со скалярами. Скалярные треды - это лишь програмная модель популярная в графических задачах. Эту модель можно использовать и для программирования SSE юнитов на ядрах Core2 (и Pentium), считая что один SSE юнит исполняет 4 скалярных треда в одинарной точности. Навряд ли вы будете в таком случае использовать 4 стека на ядро. Точно так же, редко кто использует 32 стека на векторном процессоре с длиной вектора 32.

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

Вот такая задача

За конкретные советы буду весьма благодарен.

Есть выражение, записанное в обратной польской записи (далее RPN), например такое:
(X + X) * (X + 2), что в RPN будет: X X + X 2 + *

У меня есть множество выражений такого вида (все разные), мне надо каждое из них обсчитать на множестве значений переменной X (для обсчёта и нужен стек).

То есть если считать, что размер множества выражений равен 1 (выражение единственное), а размер множества значений переменной равен 3, то мне надо получить 3 числа, каждое из которых будет результатом обсчёта выражения с подставленным конкретным значением переменной X.

В реальной жизни - выражений не одно, а много (десятки тысяч), возможных значений переменной X не 3, а тысячи.

Сейчас мой кернел работает по такой схеме:

__global__ void Kernel(int* pExpressions, int nNumExpressions,
float* pX, float* pOut, int nNumTestCases)
{
const int nThreadIdx = __mul24(blockDim.x, blockIdx.x) + threadIdx.x;
const int nThreadIdxInBlock = threadIdx.x;
const int nNumThreads = __mul24(blockDim.x, gridDim.x);

for (int nItem = nThreadIdx; ; nItem += nNumThreads)
{
int nExpression = nItem / nNumTestCases;
if (nExpression >= nNumExpressions)
break;

int nTestCase = nItem % nNumTestCases;
pOut[nExpression * nNumTestCases + nTestCase] = EvaluateTestCase(
pExpressions[nExpression], pX[nTestCase], nThreadIdxInBlock);
}
}

Внутри EvaluateTestCase заводится свой стек, живущий в shared mem, для каждого вызова EvaluateTestCase стек свой - посколько EvaluateTestCase вызывается для разных значений pX. Кол-во выражений и значений переменной кратно 32, каждый варп работает по одному выражению (divergence по идее нет).

Это всё щастье работает в два раза медленнее, чем на Athlon X2 4800+, считающем обоими ядрами. После соотнесения частот и прочих характеристик я прихожу к выводу, что эта задача будет считаться на GTX280 в 4-5 раз быстрее, чем на Core2Quad 3ГГц, что меня совсем не радует. По идее - задача исключительно параллелится, но пока что эффективность удручает.

Забыл сказать

Экспериментирую я сейчас на 8500GT, всего 16 конвейеров и всё прочее медленное.

Я думаю тут рано переходить

Я думаю тут рано переходить на использование регистров. У вас чтение выражения из памяти в EvaluateTestCase не coaslesced, значит на порядок медленнее, чем могло бы быть. Для контроля достигнутой скорости доступа к памяти посчитайте общее количество байт ваш кернел читает и пишет в память, и поделите на общее время исполнения. Получится какая-то скорость в ГБ/c. Она должна быть близка к пиковой пропускной способности, обычно 80-90% пропускной способности шины (12.8ГБ/с на 8500GT). У вас будет гораздо меньше --- есть над чем работать.

Другое узкое место --- вы читаете одни и те же данные в соседних тредах независимо, то есть потребляете на один или несколько порядков больше пропускной способности памяти, чем надо. Одно решение: читать выражение через кеш (текстурную память, лучше линейную, 1D). Недостаток: у кэша большая задержка, несколько сот циклов. Другое решение: прочитать сначала выражение в shared memory, потом уже читать оттуда из разных тредов. В этом случае надо будет убедиться, что все треды одного блока работают с одним выражением. Можно использовать паддинг если это не так.

После того, как доступ в память эффективно организован и производительности ограничена пропускной способностью исполнения инструкций, будет имеет смысл оптимизировать доступ к shared memory и регистрам.

Совершенно согласен по сути

Однако - я не очень понимаю, как это бороть.

1. Не-coalesced чтение в EvaluateTestCase. Да - оно в цикле, который проходит по всему выражению int-by-int, но я хоть убей не могу понять, как от этого избавиться. По thread id я считаю номер выражения для вычисления и номер тест кейса для подстановки, как в это втиснуть ещё и сам обсчёт - просто не знаю.

2. Читать в shared mem я не могу - она почти всё занята стеками (я запускаю максимально возможное кол-во потоков (кратное 32) на блок в зависимости от ожидаемой макс глубины стека, вот почему я где-то выше сетовал про 32K памяти в shared mem.

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

Кстати, я вылизал кернел настолько, что он стал примерно в полтора раза быстрее моего атлона (был в два раза медленнее), но это всё суть ловля блох, не глобальный прорыв, а рытьё инета пока не дало результатов по параллельному обсчёту постфиксных выражений.

заменить не-coalesced чтение

заменить не-coalesced чтение на чтение из текстуры вовсе не сложно. Надо только привязать текстуру, скажем tExpressions, к участку памяти где хранятся выражения --- cudaBindTexture(0, tExpressions, pExpressions, size), и вместо pExpressions[index] писать tex1Dfetch(tExpressions,index). Я думаю одно это может дать сильный прирост, поскольку устраняет не-coalesced доступ и снижат нагрузку на шину памяти.

А зачем вы запускаете максимальное количество потоков на блок? И сколько это получается? Какова максимальная глубина стека?

Примерно так

Дело в том, что в выражении могут быть вызовы функций, описанных в этом же выражении (задача из области Genetic Programming, поэтому буду называть такие функции так, как их принято называть в этой области - Automatically Defined Functions или ADFs).

Пример такого выражения:
X X + 2 - ADF0 arg0 arg0 *

Это означает вот что: ADF0(X + X - 2), где ADF0 имеет один аргумент arg0 и возвращает его квадрат (arg0 * arg0). Таких функций в выражении может быть скажем до пяти, аргументов у каждой может быть скажем тоже до пяти, кроме того, ADF со старшим номером может вызывать ADF с младшим номером (ADF2(ADF1(ADF0(arg0)))) сколько угодно много. Обращать внимание на возможность наличия ADFs не надо - они не меняют задачу принципиально - но объясняют, почему на некоторых выражениях максимальная глубина стека может быть 60-70 элементов.

Я могу варьировать макс глубину стека перед запуском кернела, предварительно попробовав на зуб выражения на одном тест кейсе на CPU, далее, в зависимости от требуемой shared mem на поток (которая определяется максимальной возможной глубиной), потоков может быть от 32 до 160. Пытаюсь я запустить как можно больше потоков потому, что при 32 потоках в текущей имплементации всё работает существенно медленнее, чем при 160.

Моё мыло: babylon73@mail.ru. Если Вы бросите туда тестовое сообщение, то я могу отослать весь вычислительный код, он не сильно здоровый, но даст полную картину. Обсуждение же можно (и думаю нужно) продолжить тут - может быть полезным многим.

ещё замечание --- основная

ещё замечание --- основная разница в функциональности между shared memory и регистрами, это 1) shared memory доступна из разных скалярных тредов исполняющихся на одном мультипроцессоре (внутри блока). Вас эта функциональность не интересует, так как стек раздельный для каждого треда. 2) shared memory позволяет индексированный доступ. Нужно вам это или нет зависит от того, что вы с этим стеком собираетесь делать.

RV770

AMD в своих материалах говорит про 1,2 Тфлопс в режиме одинарной точности и про 240 в режиме двойной для Radeon HD 4870.

(если i = blockDim*blockIdx.x

(если i = blockDim*blockIdx.x + threadIdx.x)
a[ i ] + = a[ i + 1 ]
примерно в 10 раз быстрее, чем
a[ i + 1 ] + = a[ i ]
1блок 512потоков

не знаю точно в чем разница

в первом случае не-coalesced

в первом случае не-coalesced a[i+1] только читается а во втором ещё и пишется

сомневаюсь что запускать всего один блок имеет смысл с точки зрения производительности

т.е. не-coalesced это когда

т.е. не-coalesced это когда пишется i потоком в i + 1 поток в случае если a[ i + 1 ] = a[ i ]

в вашем случае я думаю да к

в вашем случае я думаю да

к слову, ваша программа навряд ли будет выдавать полезный результат так как вы пишете в тот же массив откуда читаете. На практике вы бы использовали что-то навроде b[i]=a[i]+a[i+1] или хотя бы барьер поставили между чтением и записью. (Локального барьера будет мало если количество блоков больше одного.)

Здравствуйте,подскажите

Здравствуйте,подскажите пожалуйста, что делать?
_global__ void test(N)
{
for(int i=1; i<N; i++)
__syncthreads();
}
N=100 000 000;
test(N);
cudaThreadSynchronize();
вот такая чтука вылетает примерно через 5 секунд работы
для любого моего кернел
не только для простой синхронизации
спасибо

Ну да, известная фишка

Ну да, известная фишка Windows-драйверов - если девайс к которому подключен монитор занят собой более 5 секунд, то срабатывает watchdog. В висте - еще хуже.

Т.е. нужно писать короткие kernels (и перезапускать их в цикле) или же считать на видеокарте, которая не является дисплеем с точки зрения системы.

в процессе написания

в процессе написания программы с использованием комплексных чисел при помощи библиотеки
возникла проблема
при компиляции комплексной функции float complex a[idx]=cexpf( I + 1.0 );
в __global__ com(float complex *a) функции
возникает проблема
INve.cu(63): error: calling a host function from a __device__/__global__ function is only allowed in device emulation mode

подскажите пожалуйста, что можно сделать в этом случае?
использование библиотеки очень желательно

Ну раз complex.h нету - можно

Ну раз complex.h нету - можно же ручками?

float complex
cexpf (float complex z)
{
float a, b;
float complex v;

a = REALPART (z);
b = IMAGPART (z);
COMPLEX_ASSIGN (v, cosf (b), sinf (b));
return expf (a) * v;
}
(это из встроенных функций gcc)

Привет интересует такой

Привет
интересует такой вопрос
есть ли какой-нибудь стандартный метод реализации последовательной части алгоритма на GPU
например копирование на хост
если выделяю одну треду if
то доступна область глобальной памяти только для соответствующего блока
Спасибо

Подскажите пожалуйста как

Подскажите пожалуйста
как работает такая штука

for( int i = 1 ; i<str ; i + + ){
if( idx = = i){ do smth; }
}

Я позволил себе чуть

Я позволил себе чуть подредактировать - чтобы знак "меньше" не воспринимался тегом.

На NVidia нормально работает - тред у которого совпало - выполняет код, а у остальных conditional execution эффективно заменяет оный код NOP-ами.

есть ли какая-нибудь

есть ли какая-нибудь возможность использовать регистры
таким образом

a=Z[idx]
b=Y[idx]

и потом для каждого idx посчитать с использованием a и b
т.е. для каждого idx будут свои a и b
вот эта штука не работает

for( j = idx - 1 ; j < idx+1; j ++)
{

if( j == idx ){

X[ j+1]=a+b*X[ j ]

}

}

Привет как использовать

Привет

как использовать DeviceToDevice
для копирования из одного GPU на другое GPU при использовании двух процессов?

Спасибо

Pages