Посоветуйте что выбрать OpenCL/Cuda/Stream

Я собираюсь использовать gpgpu в своей дипломной работе.
В основном будут решаться большие СЛАУ, желательно с double точностью.. но СЛАУ я не собираюсь ограничиваться..

Я ещё не решил какое API я буду использовать OpenCL/Cuda/Stream.
Сейчас у меня есть ATI HD4850, но я планирую взять новую видеокарту, максимальная стоимость 15000 т.р.
Посоветуйте какая из видеокарт больше подходит для каждого из трёх API (конкретные модели), а также на каком API мне остановиться.
Или хотя бы выскажете своё мнение, буду очень признателен.
(Материнка asus P5Q3. если понадобиться купить другую - не вопрос)
Приоритетом выбора является кол-во GFlops(двойной точности) и обьём памяти.

На данный момент я больше склоняюсь к OpenCL, из-за поддержки вычислений на процессорах и кроссдевайсностью. Но я не знаю у кого сейчас поддержка OpenCL лучше у NVidia или ATI?

О конечном выборе и о результатах своей работы я отпишусь.

Forums: 

по поводу opencl это к lexa

по поводу opencl это к lexa вопрос.
знаю лишь что карточки ati показывают лучшие результаты при вычислениях с двойной точностью. У nvidia пока, насколько мне известно, только один юнит на мультипроцессор отвечает за операции с двойной точностью.

С одной стороны, двойная

С одной стороны, двойная точность на ATI уже сейчас (со всеми их проблемами с производительностью OpenCL) быстрее чем у NVidia.
Новые карты на ферми возможно эту тенденцию победят, но они первое время не будут дешевле $500.

Еще с одной стороны, средства пошаговой отладки есть пока только под CUDA (Nexus), если вы "пошаговый отладчик", то на сегодня - только так.

С третьей стороны, под NVidia есть CuBLAS т.е. для решения систем линейных уравнений можно брать готовую библиотеку, а не писать самому.

С четвертой стороны, OpenCL уже сейчас обеспечивает реальную переносимость между картами разных производителей.

Т.е. если бы я писал диплом - я бы брал OpenCL, а какую карту - не знаю, может быть и NVidia (GTX260, как разумно стоящую). Если бы в дипломе захотелось бы извернуться, то сделал бы гибридную схему с приближенным вычислением в single и уточнением в double.

Но решить за вас не могу.

Спасибо за развёрнутый

Спасибо за развёрнутый ответ.

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

По поводу уточнения single до double. Тоже над этим думал, тем более это хорошо подходит к итерационным методам.. но, для того, чтобы выбрать что эффективнее(например только ati+double или nvidia с уточнением, и т.д.) нужно делать тесты...

Думаю скорей всего буду использовать OpenCL. Но вот дилемма Ati или NVidia..
Ati на двойной точности быстрее.. А NVidia как я понял обладает более лучшем SDK(?), и вроде CUDA очень похожа на OpenCL, что даёт большой плюс.
Сначала я думал взять видеокарты обоих производителей и сравнить, но как я понял в одном компьютере они плохо уживаются.. постоянно перетыкивать не очень хочется..
Средства пошаговой отладки это конечно хорошо, но всегда можно и без них обойтись, правда это суждение основано только на опыте программирования приложений не для GPU.

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

NVIdia и ATI в одной машине

NVIdia и ATI в одной машине уживаются нормально, я про это написал вчера: http://www.gpgpu.ru/node/184
Точнее, уживаются нормально если у вас два монитора, по одному к видеокарте.

Если монитор один, то дело хуже, придется перетыкать монитор, кроме того "пошаговая отладка" (Nexus) в такой ситуации не может быть использована.

Вот лично я бы ориентировался бы на NVidia + OpenCL + гибридная схема (решение в single, уточнение в double, это в любом случае быстрее на любой архитектуре), а если бы осталось время (если защита летом, то скорее всего не останется) то исследовал бы оптимальные параметры и на карте ATI (они скорее всего будут другими, чем на NVidia).

Я всё ещё в раздумьях (у меня

Я всё ещё в раздумьях (у меня много других дел, да и по диплому тоже есть работа над другими частями... да и с CUDA можно вообще на эмуляторе тестировать, так что пока ещё не слишком спешу.. ).
Думаю поступить следующим образом: подождать пока выйдет GTX480, даже если не куплю её, то думаю что она окажет влияние на цены предыдущий моделей.
Можно поступить ещё следующим образом, взять видеокарту попроще и отлаживать на ней. А перед дипломом обменять её на более производительную, чтобы получить хороший результат. Вот только я не вижу GTX260 в магазинах, GTS250 полно, а GTX260 нет(причём GTX280 285 есть).
Я даже думал вести разработку на CUDA, потом портировать на OpenCL(как я понял, это происходит вообще без проблем) и гонять на ATI. Вот только я что-то не вижу тестов производительности ATI(может не там ищу), очень хотелось бы узнать показатели HD5850, HD5870 хотя бы на SGEMM/DGEMM(было бы вообще хорошо узнать эти показатели для GTS250 и GTX260).

По поводу моей задачи: я собираюсь решать симметричные положительно определённые разряженные матрицы, которые приводятся к ленточному виду, так что использовать CUBLAS здесь будет не очень эффективно. Смотрел в сторону итерационных методов, которые прекрасно подходят под принцип уточнения решения(сначала float, а потом double), а именно метод сопряжённых градиентов(который прекрасно распараллеливаится, так как состоит из умножения матрицы на вектор, и вектора на вектор(скалярное произведение), и который можно оптимизировать под матрицы ленточного вида).
Только определился с методом(в идеале конечно лучше реализовать несколько), и как часто бывает нашёл в интернете, что кто-то уже занимался этим, даже есть исходный код - а именно OpenNL, в которой есть plugin CNC(который работает на Cuda). Но у меня есть ряд идей оптимизации конкретно под мою задачу.

C dgemm/sgemm проблема. Ну то

C dgemm/sgemm проблема. Ну то есть можно, конечно, Волковскую реализацию перенести впрямую, а вообще готового кода я не видел (хотя и не искал специально).

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

Интересно сделать следующий

Интересно сделать следующий тест:
Каждый поток считывает значение из памяти,
затем делает n-раз MAD над этим значением,
затем записывает результат.
И соответственно вычисляется GFlops/s .Можно даже построить график зависимости от N.
Чтобы bandwidth памяти играл не большую роль, нужно использовать большое N.
Теоретически при большом N можно получить реальную вычислительную производительность. Понятное дело что на реальных задачах она будет ниже, но имхо это хороший показатель для сравнения видеокарт.

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

Ну вот конкретно для

Ну вот конкретно для dgemm/sgemm важна не скорость чтения из глобальной памяти, а скорость работы локального стораджа (регистров + shared memory). Как раз с глобальной памятью у радеонов все хорошо и вычислятор тоже быстрый.

Ну меня интересует хотя бы

Ну меня интересует хотя бы скорость вычислятора. Сколько GFlops на Single и Double точности.

Вот кстати на счёт умножения матриц, во многих примерах обычно потоки считывают значения в локальную память, а затем выполняют вычисления.
С чем это связано? Для того чтобы получить coleascing при считывании строк одной из матриц? то есть допустим если бы умножали столбцы на столбцы, а не строки на столбцы, то использовать shared memory не обязательно? Или всё сложнее?
(понятное дело столбцы или строки зависит от способа хранения..)

А вы посчитайте чтения-записи

А вы посчитайте чтения-записи из глобальной памяти для разного размера блока (ну хоть 1 и 2x2) и сразу все поймете. Глобальная память - медленная, если бы она с скоростью вычислятора работала при произвольном доступе, жизнь была бы другой.

Ну и абстрактные гигафлопсы вычислятора - просто неинтересны. Ну, почти.

Меня проглючило, я почему-то

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

Вопрос на счёт абстрактных

Вопрос на счёт абстрактных гигафлопсос вычислятора:
Вот здесь http://www.intuit.ru/video/59/ на 3:38:20 (чтобы не мотать: http://img683.imageshack.us/img683/3338/gemm.png ) приводится результат профилировщика для умножения матриц с использованием shared mem, более 80% ушло на "instructions", вот вопрос чего 80%? Времени? И что входит в instructions? Если только вычисления, то флопсы вычислятора очень интересны

В пределе (все MAD, нету

В пределе (все MAD, нету тормозов на всяких местах) - вы придете к теоретическому пределу.

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

На задачах типа *GEMM у ферми, кстати, будет изрядный прирост за счет роста числа регистров и shared mem

Кстати, и не при произвольном

Кстати, и не при произвольном доступе - похоже тоже, тормозить будет память.

Даже при банальном умножении вектора на константу, слово считали, слово записали - 8 байт на операцию. Для половины терафлопса - нужно 4 терабайта в секунду. А имеем раз в 30 меньше (не говоря о том, что умножение вектора на константу выгоднее в таком случае делать на CPU, потому что мы ограничены полосой заливки-выливки).

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

По поводу sparse linear systems

По поводу sparse linear systems можно посмотреть еще по линку:
http://lattice.bu.edu/quda/ ,
и документация к нему
http://arxiv.org/abs/0911.3191.
Люди реализовали CG и BiCG алгоритмы. просто откомпилируйте и можете погонять тесты (invert_test.c, в частноти) Хотя написан пакет под конкретный вид матриц (Wilson-Dirac operator), возможно попробовать приспособить под свои задачи. По крайней мере, как пример реализации указанных алгоритмов на CUDA вполне пойдет (inv_cg_quda.cpp and inv_bicgstab_quda.cpp).
Eще мне понравился self-tuning позволяющий приложению подобрать оптимальные параметры под конкретное железо (размеры блока и тредгрида).

P.S.: да, все это написано под linux, ессно.

Вопрос к Алексею: А каким

Вопрос к Алексею: А каким образом, на Вашем форуме, можно связаться с другими участниками форума?
Я бы хотел пообщаться с J0hn более приватно, по его вопросу)

Кликаете в ник, дальше в

Кликаете в ник, дальше в "Обратная связь".

Это если пользователь зарегистрирован, конечно.

Но мне было бы приятно, если бы вопросы программирования GPU обсуждались бы публично.

Приношу свои извинения:) Но

Приношу свои извинения:) Но вопрос не напрямую связан с программированием GPU. Почему то, ни на ник J0hn, ни на Ваш, ни на мой кликнуть мышью не представляется возможным:(((

Если пользователь снял галку

Если пользователь снял галку "personal contact form", то так и будет.

Тогда остается надеяться, что ваш корреспондент откликнется по комментарию в форуме

Я здесь :)

Я здесь :)
Отправил ПМ

HD4850

Немного протестировал как обстоят дела с умножением матриц на HD4850(512MB Sapphire), версия SDK - 2.01, драйвера - 10-3, ОС Windows 2008R2X64, мп Asus P5Q3. тестировал OpenCL версию умножения матриц(из примеров SDK, мне кажется они CL код для матричного умножения поменяли по сравнению с 2.00 . Теперь есть два ядра - одно использует local storage, другое нет.. использовать LS или нет C++ код определяет автоматически(в случае 4850 соответственно нет). Я могу ошибаться, может это уже было в 2.00).

Размер матриц 2048x2048.

Приближённый результат для Single точности - 98GFlop/s

Для Double - 28GFlop/s
(я просто заменил в CL коде float4 на double4 и добавил
"#pragma OPENCL EXTENSION cl_khr_fp64 : enable".
в C++ коде заменил cl_float на cl_double.
Тесты прошли успешно(ключ -e))

Интересно узнать результат для 5870

ATI OpenCL

Есть ли для ATI что-то типа "Best Practices" от NVidia?
Готовые советы/правила как следует использовать глобальную память, регистры, ветвления и т.п.
Или хотя бы описание архитектуры железок, чтобы можно было вывести для себя правила что хорошо, а что плохо?

OpenCL ATI

Спасибо, как раз то что нужно.
Бегло просмотрел - там есть глава про отладку, надо будет попробовать.

Хмм..

http://www.ixbt.com/video3/gf100-2-part1.shtml
"Важно отметить, что несмотря на значительно ускоренные вычисления с двойной точностью, на которые способна архитектура Fermi, игровые решения на базе чипа GF100 намеренно приторможены и исполняют такие расчёты медленнее, чем теоретически могут. Производительность 64-битных вычислений в GeForce GTX 480 искусственно снижена вчетверо. В случае GTX 480 до 168 гигафлоп вместо возможных 672. "

http://forums.nvidia.com/index.php?showtopic=164417&pid=1028870&mode=thr...

Проделки маркетологов?

Или маркетологов или вот

Или маркетологов или вот GTX480 прямо сейчас продаются, а теслы еще нет.

Но это вполне официально:
http://forums.nvidia.com/index.php?showtopic=165055

ECC, в 4 раза быстрее на double и все такое. Вероятно, Теслы хреново покупали, если они от GTX280 отличались только количеством памяти.

"Вероятно, Теслы хреново

"Вероятно, Теслы хреново покупали, если они от GTX280 отличались только количеством памяти."
Так вот я про то же... смысл было покупать C1060, когда за приблизительно такую же сумму можно купить комп с тремя/четыремя gtx480 (м.п. типа Asus P6T7), при том что для теслы тоже мат плата нужна...
Теперь видно маркетологи решали дать людям смысл...

Наконец-то

Наконец-то решился и заказал GTX285(XFX GX-285X-ZWBA, немного оверколченная).
Далее перечислю пункты почему, собственно, я остановился на GTX285.
- Скорость выбранного метода решения(метод сопряжённых градиентов) сильно зависит от пропускной способности памяти: самая ресурсоёмкая операция на одной итерации - умножение матрицы на вектор. Кэшировать в локальной памяти из глобальной почти нечего(матрица разряженная, кэшировать вектор на который умножается матрица особо не получится. попробую использовать текстурную память). Coalescing достигается не трудно.
Можно сравнить ПСП видеокарт(референсных):
GTX480 - 177.4 GB/s
GTX285 - 159.0 GB/s
HD5870 - 153.6 GB/s
Если критерием является ПСП, то наиболее привлекательной по цене является GTX285.
Также можно сравнить с ПСП современных CPU: Core I7 930 - 25.6 GB/s ...
Есть задумка по поводу несколько другого подхода, в котором ПСП будет играть меньшую роль, но до сдачи диплома наверное не успею.
- Махинации NVidia с ECC и занижением скорости Double вычислений (GTX480), хоть это не так важно для моего проекта, всё же хотелось бы иметь более мощную платформу для будущих проектов.
- GTX480 по множеству обзоров является очень шумной, энергозатратной и соответственно горячей. Хотя этот пункт не является решающим для меня.
- GTX480 неясно когда поступит в продажу + время доставки в мой город, может быть уже очень поздно.
- GTX480 будет иметь очень завышенную цену, за эти деньги можно купить почти две GTX285 и получить большую суммарную ПСП(задача не требуют пересылок между видеокартами/вычислительными модулями, прекрасно масштабируется хоть на кластер.. кстате у меня в планах добавить в код что-то типа MPI, чтобы создать кластер из CPU/GPU)
-по поводу ATI, ИМХО NVidia имеет на данный момент лучшую поддержку OpenCL/Cuda (разные инструменты, документация, даже много материала на русском(хотя предпочитаю читать первоисточник)). (я пробовал запустить ядро, в котором матрица и вектор хранятся с разной точностью на HD4850 - одно из ядер вообще не захотело собираться, другое даёт неправильный результат. Хотя на CPU(OpenCL от ATI) и на Tesla C1060 всё запустилось и правильно посчиталось, может это из-за того что серия 4xxx плохо поддерживает OpenCL, если интересно могу запостить код в отдельную тему)
- этот пункт конечно не по данной теме, но имеет некоторый вес - после долгого застоя стерео очки опять начали появляется и у NVidia есть уже готовые решения(жалко что шлемы очень медленно развиваются). Эта тема мне интересна и также полезна для моей работы, я даже как-то делал трекер движений на инфракрасных светодиодах.

Ну с учетом того, что с

Ну с учетом того, что с производительностью двойной точности кинули - GTX2xx действительно выглядят поудачне по price/performance. С другой стороны, 400-я серия должна быть попроще в деле достижения хорошей (не предельной) производительности за счет кэшей,

Уговорили

Уговорили, кэши так кэши :)
А если серьёзно, не прошло и двух дней с момента заказа мной GTX285, как в том же магазине появилась GTX480 доступной для заказа, причём небольшая партия уже отправилась и приедет эта партия даже раньше моей GTX285.
В итоге пошёл в магазин и перезаказал(повезло, что они сделали перезаказ без всяких проблем).
В итоге жду Zotac GTX480 ZT-40101-10P. Цена(19990) оказалась намного ниже той, что я ожидал.

Да, в первый месяц 20к

Да, в первый месяц 20к (плюс-минус) - похоже нормальная цена. Хотя на текущий день я не вижу их реально в наличии, похоже только *едут*

Эта теория (и маркет и

Эта теория (и маркет и прайс.ру) - мне известна. А на практике - я мониторю, очень хочу такую пощупать живьем (хотя плюс-минус месяц ничего не решают у меня)

Только есть магазины и магазины. Одни пишут "в наличии" когда у поставщика оно появилось в базе (даже не на складе), а вторые - честные.
Вот все честные(*) пишут сейчас в духе "на дальнем складе в Подмосковье, срок поставки 5-10 дней". Означает это, скорее всего, что оно уже все тут, но еще не растаможено и лежит на складе временного хранения. Может завтра растаможится, а может во вторник-среду.

(*)Сноска к звездочке: я, естественно, не всех "честных" просмотрел, а пяток тех про кого мне известно что а) цены разумные б) не имеют привычки врать про наличие на складе и прочие сроки поставки.

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

Tesla C2050

Появилась небольшая возможность протестировать Tesla C2050,
nbody -n=32768 (ECC включён)
signle - 520GFlop/s (на GTX480 - 655)
double - 210GFlop/s (на GTX480 - 113)

dtod - 76GB/s (ECC включён)

Протестировал скорость передачи своими ядрами:
ECC: 95.79 GiB/s
non-ECC: 109.63 GiB/s

На GTX480 это же ядро показало 136.9 GiB/s. Отношение ПСП по спецификациям: 1.23. Отношение намеренного(без ECC): 1.25

//если получится ещё что-нибудь протестировать, обновлю это сообщение

Вот что вычитал в "NVIDIA

Вот что вычитал в "NVIDIA CUDA Programming Guide Version 3.0":
"Code compiled for devices with compute capability 2.0 and greater may make use of C++ classes, as long as none of the member functions are virtual (this restriction will be removed in some future release)."

Вот интересно future release чего? Новая Compute Capability или новые драйвера/тулкит?
И если будут виртуальные функции, получается что такие функции не будут инлайниться?

И еще вопрос, ответ на который мне не попадался, где хранится код программы на устройстве? Гарвардская архитектура или фон-Нейманаская?
И как это зависит от major версии Compute Capability (вроде на устройствах 2.X, можно запускать несколько разных ядер одновременно.)

В future release чего - не

В future release чего - не знаю, возможно и тулкита.

Но начиная с Fermi, на девайсе есть нормальный вызов функций и рекурсия, если я правильно все понял.

GTX480

Сегодня получил GTX480, успел сделать пару тестов по моей задачи - на многих тестах(они ещё не сильно оптимизированные, проверял один и тот же код OpenCL) быстрее чем C1060 в 1.8 раза. Если посмотреть на отношение ПСП 1.74, сразу видно - как и ожидалось узким местом является ПСП.

А если несложно, можно

А если несложно, можно результаты bandwidthTest от C1060 и от новой карты. По идее, неважно, OCL-евский вариант или CUDA, там небольшая разница должна быть.

Nexus работает на новой, кстати?

Запускал без параметров

Запускал без параметров Cuda'овский bandwidthTest
Device 0: GeForce GTX 480
Quick Mode
Host to Device Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4382.5
Device to Host Bandwidth, 1 Device(s), Paged memory
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 3454.8
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 111898.8

Device to Device Bandwidth - как я понимаю считывание идёт с девайса, и запись на девайс(нет времени смотреть код).
Если надо с какими-нибудь дополнительными аргументами, пишите(может буфер больше).
По поводу C1060, прямо сейчас у меня нет к ней доступа. Если будет ещё доступ, обязательно запущу.

Кстате, а как Nexus получить? Я по-моему месяц назад регистрировался на его сайте, но мне не пришло никаких писем.

Надо будет ещё GEMV протестировать.. Кстате, у меня сейчас Core I7 930, и его простетитить на GEMV.

В теории - надо бы. 32

В теории - надо бы. 32 мегабайта со скоростью 120 тысяч мегабайт в секунду - это уже суб-микросекундные времена (собственно, ~250 наносекунд). Надо на полпамяти пускать.

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

Вот commandline для 400Mb:
bandwidthTest.exe --dtod --mode=range --start=409600000 --end=409600000 --increment=1

На GTX280 получается 121Gb/sec, что вполне в ожидавшемся диапазоне. На старой тесле должно быть заметно меньше, там памяти больше, но она медленнее.

Хм..

Хм.. вот результат
Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
409600000 112156.3

На GTX280 121Gb/sec этот тот вывод который тест выдал, или ещё умножить на 2(чтение запись)?

Тот вывод, который тест

Тот вывод, который тест вывел, ничего не умножал. Оно там внутри теста уже умножается само:

  1.     //calculate bandwidth in MB/s
  2.     bandwidthInMBs = 2.0f * (1e3f * memSize * (float)MEMCOPY_ITERATIONS) /                                          (elapsedTimeInMs * (float)(1 << 20));

280-ка не разогнаная никак, драйвера - последние WHQL. CUDA 3.0, Win7/x64

Возможно, cudaMemcpy(....DeviceToDevice) еще не оптимизирована (по идее, это вообще часть драйвера).

Ок, драйвера сегодня скачивал

Ок, драйвера сегодня скачивал WHQL для Windows7X64(для WS2008R2 вроде нет отдельных), хотя у меня стоит WindowsServer2008R2(R2 только в x64 варианте существует). Как я понял у W7 и WS2008R2 одно и тоже ядро, и драйвера должны подходить.. завтра наверное ещё линукс поставлю.

Да, линукс может и помочь. По

Да, линукс может и помочь.

По опыту предыдущих раз (8800, 280), проходит где-то полгода, пока перформанс стабилизируется. И это не только трансфера касается (что естественно, компилятор то, особенно в случае OpenCL - в драйвере)

OpenCL

Вот результат oclBandwidthTest с теме же параметрами

Device to Device Bandwidth, 1 Device(s)
Transfer Size (Bytes) Bandwidth(MB/s)
409600000 109985.9

Насчёт CUDA - там ведь тоже есть некоторый транслятор/компилятор из PTX кода(если конечно используется PTX), но, конечно, оптимизации могут быть в самом компиляторе который PTX генерирует.

Ну в этом же тесте нет

Ну в этом же тесте нет никакого пользовательского кода, а есть только вызовы стандартных функций CUDA/OCL. Имплементированы эти функции в драйвере (считая сюда же и DLL-ки, которые идут с драйвером). Возможно - неоптимально.

Я как куплю - поизучаю вот в этом вот духе: http://www.gpgpu.ru/articles/nvidia_8800gtx_propusknaja_sposobnost__pamj...
Т.е. bandwidth на чтение (и запись - тоже) но уже не средствами драйвера, а средствами CUDA (и OCL).

Что касается трансляции - я почти уверен, что PTX/cubin драйвер тоже транслирует не в лоб, иначе бы не было таких просадок в скоростях на некоторых версиях драйверов.

Pages