CUDA 3.1 beta

На девелоперском сайте NVIdia (увы, нужен логин туда), выложили CUDA 3.1 Beta.

Из реально вкусного, там 16-way kernel concurrency, плюс обновления для CUBLAS/CUFFT.

Но:

  • Примеров (GPU Computing SDK) пока нету, только сам CUDA Toolkit.
  • Драйвера, поддерживающие все это дело - только для Tesla. Наверное, можно их поставить обычным хаком (прописав PCI IDs в setup.inf), но я не стал. Update: драйвера 257.15 для GTX480 появились на сайте NVidia в разделе с бета-драйверами.
Релиз обещают "в следующем месяце", проще подождать.

Update: в форум запостили ссылку на TPB со всем этим счастьем. Update2: С драйверами 197.xx не работает.

content classify: 

Tags: 

Comments

Так ведь в  Fermi

Так ведь в  Fermi многозадачность теперь есть

Эту гадость я познал

Эту гадость я познал самостоятельно - SDK выложили (на nvdeveloper) только вчера или позавчера, на 10 дней позже остального.

Работает ли оно с десктопными драйверами 197-й версии - не проверял пока.

Проверил - не

Проверил - не работает.

Нагибать туда тесловские драйвера пока не хочу пробовать.

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

Кто установил, посмотрите, в ReleaseNotes к профилировщику для устройств Compute Capability 2.0 какие счётчики есть(особенно интересуют coalescing), там отдельная секция в описании.

Profiler counters for GPUs

Profiler counters for GPUs with compute capability 2.0

# branch : Number of branches taken by threads executing a kernel. This counter will be incremented by one if at least one thread in a warp takes the branch.
# divergent branch : Number of divergent branches within a warp. This counter will be incremented by one if at least one thread in a warp diverges (that is, follows a different execution path) via a data dependent conditional branch. The counter will be incremented by one at each point of divergence in a warp.
# sm cta launched : Number of threads blocks launched on a multiprocessor.
# local load : Number of executed local load instructions per warp on a multiprocessor.
# local store : Number of executed local store instructions per warp on a multiprocessor.
# gld request : Number of executed global load instructions per warp on a multiprocessor.
# gst request : Number of executed global store instructions per warp on a multiprocessor.
# shared load : Number of executed shared load instructions per warp on a multiprocessor.
# shared store : Number of executed shared store instructions per warp on a multiprocessor.
# instructions issued : Number of instructions issued including replays
# instructions executed : Number of instructions executed, do not include replays
# warps launched : Number of warps launched on a multiprocessor.
# threads launched : Number of threads launched on a multiprocessor.
# active cycles : Number of cycles a multiprocessor has at least one active warp.
# active warps : Accumulated number of active warps per cycle. For every cycle it increments by the number of active warps in the cycle which can be in the range 0 to 48.
# l1 global load hit : Number of global load hits in L1 cache
# l1 global load miss : Number of global load misses in L1 cache
# l1 local load hit : Number of local load hits in L1 cache
# l1 local load miss : Number of local load misses in L1 cache
# l1 local store hit : Number of local store hits in L1 cache
# l1 local store miss : Number of local store misses in L1 cache
# l1 shared bank conflicts : Number of shared bank conflicts

И при запуске - видны только они

Только с текущим десктопным драйвером (197-м) все это как-то не работает, примеры ругаются примерно так:

d:/bld_sdk10_x64.pl/devtools/SDK10/Compute_3.1/SDK10/Compute/C/src/bandwidthTest/bandwidthTest.cu(602) : cudaSafeCall() Runtime API error : CUDA driver version is insufficient for CUDA runtime version.

Т.е. нужны драйвера 256+ (255+ на Linux). Можно, вероятно, пробовать натянуть драйвера от Теслы (которые есть), пишут что работает, но это же вторую видеокарту ставить и все такое...

Спасибо! Вроде вот этих

Спасибо!
Вроде вот этих счётчиков не было:
# l1 local load hit : Number of local load hits in L1 cache
# l1 local load miss : Number of local load misses in L1 cache
# l1 local store hit : Number of local store hits in L1 cache
# l1 local store miss : Number of local store misses in L1 cache
# l1 shared bank conflicts : Number of shared bank conflicts

Вот только что значит "local load hit" и т.п.? shared mem, которая физически находится вместе с L1, кэшируется в L1?

Не, local memory - это

Не, local memory - это никакая не Shared.

5.3.2.2 (от Programming Guide 3.0):
...
The local memory space resides in device memory, so local memory accesses have same high latency and low bandwidth as global memory....

И local - это локальные (для треда) данные, которые компилятор счел неправильным класть в регистры.

Ясно, спасибо! Раньше

Ясно, спасибо! Раньше встречал local, но видимо мозг закинул эту инфу в очень далёкую корзину.

Я надеюсь ядер которые

Я надеюсь ядер которые используют мало регистров это никогда не каснётся.

Проверил свои ядра (через

Проверил свои ядра (через clGetDeviceInfo(CL_PROGRAM_BINARIES)), .local не наблюдается.
Надо сделать автоматическую проверку, чтобы если что big warrning.

Хм.. в самом загруженном ядре

Хм.. в самом загруженном ядре увидел:
.reg .f32 %f<2>;
.reg .f64 %fd<17>;
.reg .pred %p<5>;
.reg .s32 %r<46>;
.reg .s64 %rl<2>;

что-то как-то слишком много - компилятор совсем не стесняется.
Я правильно понимаю, что это требует больше 80 регистров? И это значит, что на Compute Capability 2.0, не получиться запустить 512 потоков в блоке? У меня сейчас это ядро запускается с 512 потоками в блоке, и результат получается адекватный.
Видимо придётся переписать код, в ассемблерном стиле - сразу объявить переменные для всех промежуточных результатов, и каждую операцию писать в новой строке.. бред какой-то..
И что значит:
"Since PTX supports virtual registers, it is quite common for a compiler frontend to generate
a large number of register names."
Что за виртуальные регистры? У одного регистра разные имена(тогда всё ясно, но как тогда для OpenCL посчитать количество реально используемых регистров)?

посмотрел сколько реально

посмотрел сколько реально используется регистров с помощью ptxas(интересно, можно ли это посмотреть с помощью OpenCL, но при беглом осмотре - вроде нет) - всё нормально, около 10.
Но вот что ещё заметил:
у PTX кода который возвращает clGetProgramInfo(CL_PROGRAM_BINARIES), есть такие директивы:
.version 1.5
.target sm_13, texmode_independent

Может быть sm_13, из-за того, что я не использовал никаких фишек Fermi. И может даже драйвер OpenCL потом компилирует это для sm_20 (--gpu-name sm_20 для ptxas). Может быть у устройств 1.x и 2.0, вообще бинарный код различается и оно 100% компилируется для sm_20

У OpenCL оборудование может

У OpenCL оборудование может быть очень разное и это уверенно прячут от пользователя (ну вот на HD4xxx сделали же эмуляцию shared-памяти через global)

Это extension. Если от этого

Это extension.

Если от этого места что-то зависит, то не проще ли сразу на CUDA делать?

Да, конечно это extension,

Да, конечно это extension, можно extension'ом было бы сделать и информацию компилятора.

Немного о том почему иногда лучше OpenCL чем Cuda(даже если не планируется использовать не NVidia карты):
Как я понял в Cuda нет возможности использовать код ядер, на языке C, сгенерированных на лету(я так понял это можно сделать с PTX. Если вдруг это можно и для Cuda C, буду только рад).
Зачем может потребоваться генерация ядер на лету:
1. в зависимости от железки, оптимальный код ядер будет разный, но это не так страшно - нетрудно собрать пять разных ядер.
2. В зависимости о размерности и структуры(для разных разряженных матриц, оптимально запускать разное количество потоков на строку при умножении на вектор. Разные ядра нужны для разной reduce в shared mem) задачи оптимальный код тоже будет разный(как и размер блока и т.п.). Можно сделать генератор ядер, который может сделать много разных вариантов(также в OpenCL, можно удобно определять свои define для кода - clBuildProgram ), и протестировать их. Бывают задачи, для которых требуется запустить много раз одни и те же ядра(например те же итерационные решатели), было бы неплохо перед началом долгих итераций, подобрать оптимальное ядро.

Кстате, DirectCompute вроде

Кстате, DirectCompute вроде тоже компилирует вычислительные шейдеры на лету.
Правда в силу не кроссплатформености, для меня DirectCompute не интересна.. Конечно если приложение использует DirectX, удобно будет пользоваться DirectCompute.
Кстате, NBody для DirectCompute, на данный момент больше всего flops показывают (при большом N, около 790GFlop/S)...

Тут ещё вот что, на данный

Тут ещё вот что, на данный момент, если важно количество используемых регистров, то нет большой разницы OpenCL* или Cuda - и там и там есть возможность посмотреть во время разработки программы.
Интересует именно возможность посмотреть во время исполнения программы(можно конечно таскать с программой ptxas, но это большой костыль.. к тому же врядли ptxas - redistibutable).
Возможные применения этому которые приходят в голову: планирование размера блока и оптимизация ядер по количеству регистров, для умещения большего колличество потоков на SM. В общем вреда от этого не будет.

*просмотр PTX кода для OpenCL ядер, могут, как я понял, убрать в будующем: "Currently, the PTX intermediate representation can be obtained by calling clGetProgramInfo() with CL_PROGRAM_BINARIES and can be passed to clCreateProgramWithBinary() to create a program object, but this will likely not be supported in future versions.".
Я думаю это связано с тем, что OpenCL компилирует код на целевой машине, а CL_PROGRAM_BINARIES используются обычно для кэширования компиляции, поэтому быстрее для пользователей будет кэшировать бинарный код, а не PTX.

Правильно, как я понимаю,

Правильно, как я понимаю, таскать с собой вообще не ptxas, а decuda.

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

На данный момент ведь NVidia

На данный момент ведь NVidia OpenCL выдаёт PTX код, его ведь надо ещё в код девайса перевести, а ptxas ещё и выдаёт колличество используемых регистров. decuda ведь не компилирует PTX?

Вы правы. А я,

Вы правы. А я, соответственно, нет (остался мыслями в Cuda 2.x).

.cubin и, соответственно,

.cubin и, соответственно, decuda. И это казалось как-то надежнее, чем PTX рассматривать.

а разве

а разве через Stream KernelAnalyzer нельзя посмотреть?

А Stream KernelAnalyzer

А Stream KernelAnalyzer удобно таскать вместе программой или может у AMD он как extension OpenCL?
Его легко использовать из программы, во время её работы?
Да, безусловно можно посмотреть, что он сгенерирует для всех известных ему чипов. Но вот если программы уже у конечных пользователей, у которых новый чип, которого не было на момент создания программы..

Что-то можно. Но на

Что-то можно. Но на конкретном устройстве. На 4xxx, как мне кажется, вообще ничего нельзя.

Вообще, производительность OpenCL на AMD пока оставляет желать, тот же SGEMM на егонном ассемблере в разы быстрее, чем на OpenCL.

С другой стороны, если про 4xxx вообще забыть (кроме маков, увы, там 5xxx еще нету), то производительностим для моих задач пока более чем хватает, ну и ладно.

Ну да, компилятор не

Ну да, компилятор не стесняется, а потом драйвер делает оптимизацию.

Попробовал 257.15 beta(я так

Попробовал 257.15 beta(я так понимаю если у меня OpenCL, то кроме драйверов ничего не надо устанавливать) для GTX480.
В итоге результаты у моей программы стали другими - стало медленней сходится. Даже не дождался, у меня порог стоит 30000 итераций.. обычно на той же задаче, для той же точности хватает около 15000 итераций(примерно столько же надо AnSys'у).
Мне такие сюрпризы не нужны, вернулся к 197.75(что у них за система нумерации драйверов?).
Нету времени разбираться в чём дело, но если такая же ситуация будет в релизе драйвера, придётся писать багрепорт с test-case'ом

Если kernel-ы - в исходниках,

Если kernel-ы - в исходниках, а не в ptx, то только драйвера меняются.

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

Kernel'ы полностью в

Kernel'ы полностью в исходниках(OpenCL), exe'шники идентичные использовал.
Для выявления ядра, которое даёт сбой, надо дампить все промежуточные результаты танцуя с драйверами и выявлять место, где появится первое расхождение. Кстати, сначала(первые тысяч пять итераций) сходится очень похоже на запуск с нормальными драйверами(невязка практически одинаковая, может даже одинаковая). Нужно ещё попробовать разные варианты - SingleDouble, SingleSingle, DoubleDouble.
Если будет время, то конечно отправлю баг репорт с test-case'ом, если времени не будет, то постараюсь просто баг репорт.

Отправил bug-report без

Отправил bug-report без test-case. Я думаю они и так найдут в чём дело - скорей всего они что-то химичили с точностью

К вопросу о драйверах:

К вопросу о драйверах:
257.19 - приложили к свежему Нексусу (Win only)
258.19 - приложили к OpenCL 1.1

Если есть логин на девелоперский сайт, то можно помацать на предмет вашей баги.

вышли 257.21 WHQL, если есть

вышли 257.21 WHQL, если есть минутка - то ваше мнение про их глюки весьма интересно (т.к. это WHQL-драйвера с CUDA 3.1 и хотелось бы понять, можно ли уже кастомеров просить обновляться)

Я у себя проблем не вижу пока, но у меня все просто совсем.

257.21

Немного потестил - сходимость такая же, как и была на предыдущей whql. Цифирки(норма вектора невязки) совпадают полностью(все значащие цифры в double), что намекает на то, что вычисления остались такими же(по результату), если были бы другие - через 15k итераций наверняка бы проявилась разница. Количество итераций такое же.

Ускорение GFlops на SPMV 1.11
Ускорение итераций 1.09
Ускорение ПСП (на ядре DAEqB, http://www.gpgpu.ru/node/180#comment-1507) 1.03

Судя по значительно ускорившемуся GFlops на SPMV, для которого главное сейчас ПСП, и незначительно ускорившемуся ПСП, double вычисления стали ощутимо быстрее(я думаю >10%). Ядер для bogoFlops не делал..

Приятно получить такой бонус ничего не делая, конечно заново строить графики в пояснялке не буду - не успею..