Сталкивался с такой проблемой, просто переустановил дрова, если интересно скачал здесь youzer.tk/load/torrent/programmy/nvidia_driver_39_s_296_10_whql/9-1-0-24

Жду кода очень.
Так как мне он очень нужен.
Не могу правится с распаралеливанием.

Проблема была в файле cutil32 и в отсутствии PATH к cl.exe на сколько я понял, точнее файл был но он был не во всех нужных каталогах

А можно поподробнее рассказать в чем была проблема?

Я тоже сейчас мучаюсь с кудой, но этот момент вроде прошел. Советую заглянуть сюдой http://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/... и написать обработчик ошибок, там при копировании памяти у тебя точно должно что то вылазить. А пока ты ничего не открыл, попробуй сделать следующее:
в строках 34-36 убиваем (void**), без него чудно должно и так работать.
строка 46 - убить (void *), оставить cudaMemcpy((void *)c, cDev,.....
а строки 41-42 убиваем полностью, не используются

Если не запустится, смотри как получать коды ошибок и их расшифровывать

по расшифровке кодов - ссылка выше
коды получаются из ф-ций работы с памятью - это результат, а после вызова ядра вызываем cudaGetLastError().
Удачи, надеюсь поможет...

и вот пример реализации

http://eigen.tuxfamily.org/index.php?title=Main_Page

http://eigen.tuxfamily.org/dox-devel/TopicLazyEvaluation.html

http://eigen.tuxfamily.org/index.php?title=FAQ#Vectorization

http://eigen.tuxfamily.org/index.php?title=Benchmark

кстати, gcc при -ftree-vectorizer-verbose=7 говорит что он вставляет рантайм проверки на алиасинг. по ассемблеру видно что в цикле проверок вообще нет. То есть по-идеи, там где критично, умный компилятор должен уметь делать проверки вне цикла. И я не против дуплицирования бинарного кода для нескольких случаев.

Интелу, зачастую, надо все разукрасить прагмами (там то, в отличие от MSVC, можно и алиасинг выключить и вообще подсказать что цикл большой).

Но про конкретный случай не уверен, конечно.

Докрутку хвоста сделает компилятор, сам, без ансамбля.
Во всяком случае, я это вижу для LLVM-based OpenCL (и AMD и Intel) и, если ничего не путаю, для ISPC тоже.

Тут то проблем нет.

ну так это же общий код.
если нужно будет написать a+b+c*d, нужно будет только новую zip_func написать.
Если нужно будет использовать хитрый итератор - тоже мало что изменится. Если захочется использовать OpenMP - только поменять код simd_transform, и весь зависимый код трогать не надо.

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

Вот немного подслащённая версия: http://ideone.com/mhI3P - там даже zip_func не надо создавать, используются lambda и bind'ы.
Если такого кода надо много писать, можно сделать специальный lambda для zip_iterator, тогда бинды будут не нужны, и код будет вида _1*_2+_3

код http://ideone.com/nFcAO

Мы ходим по кругу.

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

Плюс к этому, 90-й фортран я знаю еще хуже C++, меня учили когда 77-й был новинкой.

Поэтому если уж внешнее исполнение, то, в зависимости от размера данных/арифметической интенсивности - OpenCL и/или ispc. Ispc я нежно люблю, хотя его тоже постепенно перетяжеляют.

http://ispc.github.com/ispc.html#generating-generic-c-output
чёткая фишка

А не спасет. Одно дело параметризованный массив, а другое - какая-то левая аллокация снаружи, невыровненая и неизвестно сколько там чего.

А вот OpenCL (который векторизуется) или ispc - да, оно самое и есть.

А он как векторизует? Встроенные типы массивы?
Или именно циклы векторизует?

про столбики: я никуда не спешу - подожду.

mulps/addps
я смотрю автовекторизация в msvc появится только в 11 версии. Я в шоке, sse2 появилось в 2001 году.
Вот смотрю на код - вот прям почти подряд идут инструкции..
У меня почему-то в памяти отложилось, что MSVC векторизовывал, я наверное с одномерным sse перепутал.

но это не беда, это просто усложняет задачу - надо сделать lamda'ы которые могут быть провёрнуты как simd, которые сами будут нужные intrinsics вызывать - та ещё шаблонная магия, но топик-то интересный..

http://software.intel.com/en-us/articles/a-guide-to-auto-vectorization-w...
кстати, вот тут в pdf'ке есть:
#pragma vector nontemporal
"gives a hint to the compiler that data will not be reused, and therefore to use streaming stores that bypass cache"
может пригодится

Пытался Intel C++ cкормить некоторые из этих прагм - не работают, нужен новый..
В общем, как я понял мне нужен и новый Intel, и gcc. А у меня CentOS 5, надо что-то мэйнстримовое поставить.

а запись по 8 байт что-ли идёт? или оно группируется?

movntps: http://blog.lexa.ru/2011/09/08/opyat_pro_movntps.html

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

Вы про какую технику говорите?
Запись в память ведь идёт по 64 байта (cache-line), это где такой регистр? Или это если всё-равно на скорость записи?

А код то какой получается? Вот мы хотим два vector<float> перемножить поэлементно и третий поэлементно же прибавить.
И хотим это делать SSE, ибо реально быстрее.

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

я не против :)
желательно тогда и в блоге тоже.

я согласен, тут это можно рассмотреть как kiss, ну с натяжкой:
в C++ есть разные low-level фишки, которые можно комбинировать друг с другом.
Но, часто комбинирование простых операций, намного менее эффективно чем одна жирная.

Вот например, есть покомпонентное умножение векторов(по-моему Hadamard или как-то так) и есть Reduce одного вектора. Скомбинировав их получим dot product.
Но, например для GPU это будет не эффективно(когда это два разных ядра). Потому что после Hadamard'а, можно воспользоваться данными загруженными из глобальной памяти в блок, и сразу на месте сделать Reduce(по блоку), то есть ядро Hadamard+Reduce. Ещё пару раз запустить Reduce(чистый) прийдётся, но это будет гораздо быстрее чем первоначальный вариант.

Ага, и на GPU передаст.
Я совершенно не шучу - PGI сделали именно это. Оно не безумно эффективно, как я понимаю, но работает.
это хорошо(я вообще не против Fortran'а), но больше напоминает matlab с gpu плагинами (или оно уже встроено?).

видел /Qparallel у intel'а, но пока особо не верю в него, так как не пробовал/не видел результатов.

Ага, и на GPU передаст.

Я совершенно не шучу - PGI сделали именно это. Оно не безумно эффективно, как я понимаю, но работает.

Параллельные нити даже Интел может запустить с /Qparallel. Даже на C/C++ может, но только в очень простых случаях, реально на практике они встречаются редко

ну как я понимаю тут главный вопрос в unroll'е?

А параллельные нити он тоже запустит?

Или может передаст на GPU при большом размере умножаемых матриц?

Низкоуровневые циклы с SIMD (и параллелизацией) налепит компилятор с фортрана. Сам. Использующий код тоже ничего не заметит.

и самое главное, если выйдет/вдруг захочется использовать какую-то новую/другую технологию распараллеливания и т.п., достаточно переписать/специализировать transform, и весь код его использующий ничего не заметит.
А в случае, когда у нас по проекту раскиданы низкоуровневые циклы по [i], это всё прийдётся править руками.

Ну подождите, итератор видит только один элемент (каждого/всех) векторов за раз.

Параллелизацию, векторизацию и прочая - придется как-то снаружи навешивать.

и a[i]*b[i]+c[i[ и все варианты, когда вместо a[i]/b[i]/c[i] константа (или две)? Безобразно много получается.

Насколько я вижу, С++ way это:
1. std::transfrom + кастомный итератор-обвёртка, который будет итерировать сразу три объекта. а дальше функтор. Это всё можно ещё подсластить шаблонной магией, и получится итератор который может итерировать произвольное колличество(compile-time) векторов одновременно. + boost::lambda (ох чувствую мне C++11 ещё не скоро светит). И я подозреваю, что такой итератор уже кем-то реализован, нужно только его имя знать.
2. a*b+c - сделать это как операцию между векторами, а не поэлементно, но действия делать не сразу, а конструировать объект специального типа, а при присвоении делать непосредственно вычисления. По-типу того, как реализованы lambda в boost.
Можно даже отложить вычисление не до присваивания, а гораздо позже - до первой необходимости знать результат.

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

Вместо нее есть restrict, но толку с него не очень много для C++.
Почему?

да блин, начитался, проникся.
везде какие-то компиляторо-зависимые опции (ну почти везде) - хорошо хоть это есть.
http://msdn.microsoft.com/en-us/library/chh3fb0k%28v=vs.71%29.aspx
и т.п.

В стандарте есть правда valarray, и компиляторы имеют право его векторизировать - но пользуются ли этой возможностью msvc, gcc, icc - нужно руками проверять.
А так как я и говорил, писать
d[i] = a[i]*b[i]+c[i]
это не C++ way, и не только из-за алиасинга.
Я за то, что векторные операции - нужно делать не по-поэлементно. тогда появляется простор для ручной оптимизации(обощённой) и т.п.
А valarray надо посмотреть, да.

В Fortran оно просто будет транслироваться в векторное.
В C - если restrict написать. Без этого - компилятор будет подозревать aliasing и не векторизует.
В C++ для операторов написать restrict просто негде, насколько я понимаю (или я чего-то не понимаю).

С ними то, что на AMD оно должно транслироваться в одну инструкцию (+цикл). На Intel AVX (трехадресные инструкции) - ну в две. И инструкция должна офигачивать 4 double за раз.
Ну, понятно, для векторов с размером не кратным 4-м должен еще как-то хвост досчитываться.

C удовольствием увидел бы такой код, произведенный C++-компилятором. Фортрановские - как я понимаю, могут.

С gil же - я быстро пролистал презентацию (слушать - никакого терпения нет) до места, где гордятся, что оно не хуже чем old plain C (по производительности). Я бы old plain C постеснялся бы гордиться.

C там в предыдущем сообщении (по ссылке). У C++ проблема с алиасингом на this

А итераторы-векторы-смарт-поинтеры - отличная штука, пока мы не хотим сделать банального:
d[i] = a[i]*b[i]+c[i]

А иногда - заметно медленнее: http://blog.lexa.ru/2011/10/07/o_semantike_c_prodolzhenie.html

Благодарственные благодарности! )

Цифр нет. Есть ощущения:

1) OpenCL-приложения работают на всех ядрах без дополнительных усилий со стороны программиста. Для C/C++ программы придется отдельно стараться (есть OpenMP, конечно)
2) OpenCL-приложения лучше векторизуются (т.к. проблема aliasing не волнует, выравнивание можно обеспечить).
Но зато clEnqueueBuffer (двойная буферизация).

Другими словами "случаи бывают всякие"

Большое спасибо всем за присланные исходники.
C:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\bin - он по этому пути какой-то файл cubin создаёт?

Можно мне тоже исходник на vpdiman@mail.ru . Никак не могу компилировать

В CUDA 4.1 они похваляются новым компилятором, на llvm, от которого все становится быстрее :)

Код точно становится другим - поэтому профайлер в руки, без него ничего сказать невозможно

В процессе работы понял что похоже копирования на самом деле нет, просто для процессора эти функции вызываются для порядка так скажем))

Ну так API такое, как же без буферов то?

Да, не вполне разумно, но с другой стороны - это дает гарантию, что в буферах никто не будет копаться в процессе работы kernel

Если вы имеете в виду HPL, то он живет поверх (потенциально любого) Lapack.

А что тогда значит следующее высказывание:

"Также GPU содержит ряд потоковых мультипроцессоров, каждый из которых способен одновременно выполнять 768 (1024 - для более поздних моделей) нитей. При этом количество потоковых мультипроцессоров зависит от модели GPU. Так, GTX 280 содержит 30 потоковых мультпроцессоров..."?

Всем спасибо за помощь, в моём случае проблема заключалась в конфиге компилятора nvidia. По каким-то причинам при ранее приходилось переписывать файлик nvcc.profile, вернул всё по-умолчанию, нормально заработало.

Я рекомендую сделать dll на C++, в которой реализовать функцию вызова кода куды, и в программе C# подключить эту библиотеку. И работать будет быстрее. Возможно CUDA .Net еще сырая, да и документации по ней мало.

А как его напрямик в массив скомпилировать? Все дело в том, что при компиляции создается файл с расширением cubin - это двоичный файл, который и нужно загрузить с помощью LoadModule. Ну можно еще его загрузить из ресурса, спрятать внутри программы или передать еще каким-нибудь хитрым способом, вот для этого, видимо и сделали возможность передачи последовательности байт. Но здесь видимо что-то не так компилируется, или загружается не то, то надо. А создается устройство с инициализацией? Примерно, вот так: CUDA cuda = new CUDA(0, true);?

Открыть файл cubin с помощью потока, считать массив байт ну и подсунуть его:

FileStream S = new FileStream(FileName, FileMode.Open, FileAccess.Read);

byte[] cubin = new byte[S.Length];

S.Read(cubin, 0, cubin.Length);

S.Close();

CUmodule module = cuda.LoadModule(cubin);

Я последнюю версию 4.0 скачал с официального сайта http://developer.nvidia.com/cuda-toolkit-40

Но Compute Capability зависит от видеокарты, а не от установленной версии CUDA API, если, конечно, причина именно в Compute Capability. Я свой пример только что проверил - компилируется и работает.

Даже и не знаю в чем причина. Может CUDA .Net ориентирован на куду версии 2.

Отправил. Надеюсь, поможет.

Процесс (kernel) - в первом приближении один (во втором - на Fermi можно несколько kernels одновременно пускать, единица гранулярности - мультипроцессор).

"Одновременных сложений" - по числу Cuda cores, 336 на обычных GTX460 и 288 на SE.

смотрите что в документации пишут на CUDA 4.1

"Starting with CUFFT version 4.1, CUFFT library is thread safe"

Значит, в предыдущих версиях - не thread safe. Обновляйтесь до 4.1

1. V 4.0
2. Девайсы не переключал. Видел пример по работе с двумя видеокартами, попробую...
3. Geforce 310M/512Мб. ОС Windows7, Core i3-330M, 2.17GHz...

Потому что при случайном локальном доступе - дает.
А при последовательном оптимальном для глобальной - отнимает.

Цитата оттуда: Карта оснащена интерфейсами DVI, HDMI и DisplayPort. Цену изделия производитель не приводит.

При использовании Driver API - все есть в драйвере.
При использовании Runtime API - cudart.dll нужной битности.
При использовании библиотек (cuBLAS, NPP и т.п.) - DLL-ки от них

какая самая бюджетная карта с поддержкой cuda?

помогите пожалуйста, в общем у меня прога производит простое скалярное произведение векторов, и все работает, но если в код вставляю __syncthreads(); то он пишет идентификатор "syncthreads" не определен, при компиляции он его съедает почему то и считает правильно, но без syncthreads он тоже правильно считает, хотя теоретически если нити в блоке не одновременно досчитают, то будет ошибка. в общем помогите разобраться

Но вопросов этот простенький пример порождает много:

- компиляция C-кода делается g++, хотя бы потому, что C & C++ образуют по-разному внешние имена для линковки:

- стоит поменять g++ на gcc на линковке и всё рассыпается ... из-за дефаултных подключаемых библиотек (а как их дефаулт - посмотреть?):

- ну и в подтверждение того, что это одно и то же (почти) приложение: