OpenCL в Delphi XE2 под Win64bit: можете протестировать на своей видеокарте?

@Prok12 · Регистрация: 08.11.2011

Студворк — интернет-сервис помощи студентам

Суть проблемы. Есть весьма объёмный ПК для научных расчётов: интерфейс, База данных, много-поточные расчёты (на CPU пока), графика 2D, графика 3D на OpenGL, анимация результатов в виде компрессированного avi-файла на выходе - всё делалось на Delphi XE2...XE6 и прекрасно работает. Надо подключить модуль для переноса части длительных фрагментов расчёта на GPU. Этот модуль пытались писать на Visual Studio 2015 - расширение языка C++ AMP, как DLL: работает, но криво: не на всех видеокартах, увы (текст тестовой программы выложен в той же папке для скачивания - см.ниже). Ну и вечная проблема с недостающими DLL к тому что написано на Visual Studio: даже опции компилятора /MT не помогают... На JAVA тоже сейчас мои коллеги пишут этот расчётный модуль - через OpenCL... Ну не в этом вопрос. Есть примеры использования библиотеки OpenCL.dll прямо в проектах Delphi (чехи, например, делали, есть ещё компоненты от MITOV для Delphi Seattle - http://www.mitov.com/ ). Эти несложные примеры работают, но...только при компиляции под Win32, а компилируешь проект Delphi под Win64 - затык...
-----
Сейчас вроде с помощью советов с другого форума, удалось траблы победить, и сделать работоспособный тест.
-----
Вот ссылка для скачивания: http://gofile.me/2Zesj/C0f3wb1o
========
Там в папке:
========
1) Полезная утилита GPU_Caps_Viewer_Setup_v1.25.0.0.exe для контроля состояния видеокарты, особенно, если нажать кнопку "More GPU Info".
2) OpenCL1_2_Delphi.zip - файл с примером от чехов (Университет Брно, насколько помню) использования OpenCL в Delphi, 2013год. У меня этот пример не всегда корректно работал при компиляции под Win64 (не на всех видео-картах). Но там хороший заголовочный файл CL.pas - его можно чуть доработать с учётом замечания про NativeInt с другого форума: заменить тип size_t на NativeInt для Win64, и использовать вместо моего "укороченного" MyOpenCL.pas - см.ниже.
---
3) Вложенная папка с моим вариантом программы-теста. Она самодостаточна: больше ничего не надо, ну разве что последние версии видео-драйверов установить. Тест сделан под Delphi XE2 (работает и под XE6). Причём работает, по крайней мере у меня, при компиляции под Win32, и что более важно- под Win64. Всего 2 файла:
--> MyOpenCL.pas: заголовочный; я его урезал - убрал ненужные мне функции, которые в принципе можно взять из файла чехов CL.pas + учёл полезное замечание с др.форма про NativeInt ;
--> FMain.pas : тестовая программа. Выделяет на GPU память (создаёт буферы) под 16 массивов (квадратных матриц из cl_Float, размер стороны которых задаётся на основной форме, по умолчанию 1200) и под ещё один массив, того же размера, куда пишется результат несложных вычислений на GPU. Потом запускается итерационный цикл (количество итераций - для проверки времени работы - задаётся тоже на форме, по умолчанию - 12). Внутри каждой итерации 16 массивов заполняются данными, затем запускается ещё и внутренний цикл: 30 проходов вызова расчётного ядра на GPU. Такая структура тестовой программы больше всего подходит под наши научные задачи (вычислит. гидродинамика), но пока это лишь тест. В тесте есть несложная проверка правильности вычислений - для одного из элементов матрицы-результата, поскольку НЕ все драйверы видеокарт (особенно старые) могут диагностировать Kernel Error.
--> программа на С для ядра - текстовый файл ProgramGPU.CL (должен лежать рядом с MyFirstOpenCL.exe !!): там можно глянуть, что делается с одномерными массивами на GPU. Она прямо передаётся в виде строки на GPU, там копилируется и линкуется - ну как обычно в OpenCL.
Нужная DLL для OpenCL - под Win64 или Win32 - подключится из системных папок Windows автоматически, в зависимости от компиляции.
----
4) Можете на своём сетапе запустить прямо готовый MyFirstOpenCL.exe (это версия 64bit !!) с параметрами:
--> размер стороны матрицы = 3200; лезут 17 таких матриц на видеокарту от 1Gb и выше; если не влезут (часть видео-памяти уже занята чем-то - GPU_Caps_Viewer в помощь) - возможны ошибки, которые диагностируются НЕ всем драйверами видеокарт, а только новыми;
--> количество итераций (внешний цикл) = 12.
---
5) У меня для указанного примера время счёта вышло:
NVIDIA GT-430 (1Gb, 96core, 700MHz core) - 3мин 37сек;
NVIDIA GTS-450 (1Gb, 192core, 810MHz core) - 1мин 36сек;
NVIDIA GTX-570 (1.28Gb, 480core, 732MHz core) - 42.5сек;
AMD Radeon R9 M275X (2Gb, 640core, 925MHz core) - 1мин 27сек - настольный моноблок ASUS 2702.
----
Может кто-то попробовать MyFirstOpenCL.exe на своих видеокартах??? Поскольку это проект дельфовый, никакие внешние DLL ему не нужны (кроме тех что в Windows у всех есть), как и всякие - джава-машины: только для видео-карты новые драйверы желательны.

Запустить, нажать верхнюю кнопку на форме, установить размер матрицы 3200 (если 1200 отработает корректно), выбрать платформу-девайс, нажать нижнюю кнопку на форме. Записать время счёта: оно выдаётся в итоговом окне.
Нам сейчас на работу надо будет покупать новую карту 4-6Gb именно для расчётов...денег более 60тр вряд ли выделят... AMD конечно подешевле, но... Такие сопоставительные тесты на задачах, приближенных к нашим, были бы полезны!! Были бы полезны и замечания участников форума по тестовой программе.

===
К сожалению, у меня нет под рукой видео-карты (акселератора) с памятью на борту 4Gb-12Gb. Было бы интересно протестить и на таких "монстрах", повышая понемногу размер входной матрицы (требуемую память тестовая программа сообщит в начале работы): для того и писалось под Win64 - нужна большая память. Буду признателен!

@FIL · 21.01.2016, 20:58

Запустил программу на компе с процессором i7-4790K со встроенной видяхой.
Определилось 2 устройства: процессор и видяха.
При выборе видяхи вылетела ошибка.
При выборе процессора программа отработала за 1 мин. 36 сек.
Я так понимаю, что при этом вычисления ведутся в режиме эмуляции видеокарты.
Соответственно, если данный алгоритм реализовать на SSE-регистрах процессора, то вероятно он будет работать значительно быстрее.

@Prok12 · 21.01.2016, 21:41 **[ТС]**

Сообщение от FIL

При выборе видяхи вылетела ошибка.

А какая у Вас видяха: встроенная в процессор? На таких не пробовал, но видимо - нельзя.
Какой объём видеопамяти (или это просто основная память): при размере матрицы хотя бы 1200 идёт?
И главное: какое сообщение об ошибке?

Сообщение от FIL

Я так понимаю, что при этом вычисления ведутся в режиме эмуляции видеокарты.

Да, на CPU тоже может работать OpenCL: медленнее, но тоже распараллеливает

@FIL · 21.01.2016, 21:56

Название и сообщение об ошибке на скрине.
Ошибка возникает при любом размере матрицы.
Судя по тестам под видяху выделяется 1Гб оперативки.

@Prok12 · 21.01.2016, 23:40 **[ТС]**

Сообщение от FIL

Название и сообщение об ошибке на скрине.

Ах вон оно что!! Если не затруднит, уберите любым текстовым редактором первую строку в файле-программе для GPU - файл ProgramGPU.CL .
Вот эту строку:
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
Её смысл: директива использовать возможность GPU работать с арифметикой Double. Нам она, собственно, пока и не нужна. Если после этого заработает, напишите пож.!

@x128 · 22.01.2016, 10:49

Запустил на i5-6600, тоже была ошибка как у FIL, после удаления строчки с "cl_khr_fp64" процесс пошел.
результат для GPU:

OpenCL в Delphi XE2 под Win64bit: можете протестировать на своей видеокарте?

результат для CPU:

@snake32 · 22.01.2016, 12:34

Тестили как-то двойную точность. Если интересны результаты:
https://www.cyberforum.ru/vide... 42529.html

@Prok12 · 22.01.2016, 13:19 **[ТС]**

Сообщение от snake32

Тестили как-то двойную точность. Если интересны результаты:

Весьма интересно! Выходит, на двойной точности AMD Radeon HD 7970 на много опередил правда более простые модели NVIDIA. Интересно, как дело обстоит при одинарной точности?
Я никак не могу определиться с выбором новой видяхи. Ориентировочно, 4-6GB, с ценой желательно уложиться в 40тр. У всех на форумах в основном NVIDIA GeForce: AMD Radeon не котируется... Так ли он плох для OpenCL, если CUDA меня вообще не интересует?
Кстати, на чём была написана тестовая программа по Вашей ссылке?

@snake32 · 22.01.2016, 14:40

Сообщение от Prok12

AMD Radeon HD 7970 на много опередил правда более простые модели NVIDIA

На момент тестов по цене 7970 была как 760 в районе 10к. А чуть позже "типа старые" 7970 (хотя "новая" R9 280X абсолютно на том же чипе) потом вообще сливали по 7,5к руб в регарде. Ну это цены до декабря 2014.

Сообщение от Prok12

Интересно, как дело обстоит при одинарной точности?

Могу сказать только за двойную точность - там радики рулят однозначно.(для себя бы за 40 взял R9 390X; 8GB; DP/SP = 1/8)
В одинарной точности нужно тестить. Там не всё так однозначно. И думаю, возможно будет зависить от конкретного алгоритма.
В плюсы к Radeon можно занести тот факт что глубина цикла в кернеле может быть очень большой, тогда как на nVidia наблюдается падение драйвера(так же как и на Intel CPU) и как я понял ограничено 64 итерациями(больше в принципе и не имеет смысла для SMID-архитектур). Так же были ситуации когда nVidia драйвер вёл себя не по стандарту, в частности я мог выделить памяти больше чем описано в доке при этом не сигнализируя об ошибке, однако, в процессе работы опять падение драйвера.
Вообщем, я склоняюсь к варианту что продукты от AMD более соответствуют стандарту OpenCL/OpenGL чем от других фирм, хотя возможно и чуть медленнее.

Сообщение от Prok12

Кстати, на чём была написана тестовая программа по Вашей ссылке?

Delphi XE2. CL.pas сам писал, так как найденные мной готовые варианты очень кривы, не удобны в использовании и ориентированы на древние версии Delphi.

@Prok12 · 22.01.2016, 18:18 **[ТС]**

Сообщение от snake32

В плюсы к Radeon можно занести тот факт что глубина цикла в кернеле может быть очень большой, тогда как на nVidia наблюдается падение драйвера

Пробовал внутренние циклы на кернеле до 50...до 1000 на разных видяхах. Драйвер падает по-разному: зависит и от общего заполнения памяти, и от видяхи тоже - Radeon R9 M275X тоже падает. Программы на C++ AMP тоже "кладут" драйвер при длительных внутренних циклах на кернеле - проверял, да и об этом многие пишут. С чем это связано: локальный перегрев шейдеров: им противопоказаны длительные операции?

@snake32 · 22.01.2016, 18:54

Сообщение от Prok12

С чем это связано: локальный перегрев шейдеров: им противопоказаны длительные операции?

Это особенности SMID-архитектур. Более сказать ничего не могу. Для них противопоказаны условия и тем более циклы.
Тут нужен другой подход: не внутри кернела крутить цикл, а сам вызов кернела крутить в цикле. Правда в этом случае узким местом становится CPU. Но можно немного сбалансировать например внутри кернела делать не более 64 итераций цикла, а потом опять перезапускать весь кернел, но уже с 65 итерации по 128 и так далее, пока не пройдёт нужное число итераций. То есть есть внешний цикл запуска кернела и внутренний непосредственно в кернеле, который за 1 вызов крутит не более 64). У меня как раз такие задачи что приходится анализировать миллионы массивов разной длины. И даже при таком подходе на унылой GTX 760 считает в десятки раз быстрее чем на CPU.

Delphi
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
procedure TCL1812_0.CompAPLoop1( ha,Zm,Vm50,Im50,gim50:CLMem );
  var cmd:CLCommandQueue;
      krn:CLKernel;
      i:CLuint;
      iter:TIVec2;
begin
  cmd := OpenCL.CommandQueue;
  krn := OpenCL.GetKernel( 'ke1812_0_APLoop1' );
 
  i := 1;
  krn.SetArg( i, ha ); Inc(i);//__global double * ha,
  /* вырезал установки кучи аргументов*/
 
  // 0-итерацию пропустил, кручу цикл [1..RangeLast.y]
  iter := ivec2( 1, OpenCL.MaxIter+1 );// OpenCL.MaxIter = 64
  while iter.x <= RangeLast.y do // RangeLast.y - максимальная длина всех массивов
  begin
    krn.SetArg<TIVec2>( 0, iter );
    cmd.NDRangeKernel( krn, nil, works, nil, nil, nil );
    cmd.Finish;
    Inc( iter.x, OpenCL.MaxIter );
    Inc( iter.y, OpenCL.MaxIter );
  end;
end;

Сам кернел:

C++
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
 
__kernel void ke1812_0_APLoop1( const int2 iter,
                        /*вырезал*/
                         )
{
    int i = get_global_id(0);
    int j = iter.x;
    int last = ilast[i];
    
    while(( j <= last )&&( j < iter.y ))// j =[1..ilast])
    {
        /*здесь расчёт*/
        ++j;
    }
}

@FIL · 22.01.2016, 19:58

Без cl_khr_fp64 получился такой результат:

@Prok12 · 23.01.2016, 09:10 **[ТС]**

Сообщение от snake32

Delphi XE2. CL.pas сам писал, так как найденные мной готовые варианты очень кривы, не удобны в использовании и ориентированы на древние версии Delphi.

Можете выложить свой вариант "заголовочного" файла CL.pas ? Или сравнить с моим, особенно в плане определений типов для OpenCL, и типов параметров вызываемых функций. С типами тут путаницы не мало, потому, видимо, у меня поначалу при компиляции под Win64 не работало!

@snake32 · 25.01.2016, 01:53

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

Delphi
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
type
  intptr_t= type NativeUINT;
  size_t  = type NativeUINT;
  size_t_3 = array[0..2] of size_t;
 
  CLchar    = AnsiChar;
  PCLchar   = PAnsiChar;
  PPCLchar  = ^PAnsiChar;
  CLstring  = AnsiString;
 
  CLshort   = int16;
 
  psize_t   = ^size_t;
  CLint     = type int32;
  CLuint    = type uint32;
  PCLuint   = ^uint32;
  CLbool    = type LongBool;
  CLulong   = type uint64;
  CLfloat   = type single;
  CLdouble  = type double;
 
  CLException = class(Exception);
 
  CLbitfield                      = type CLulong;
  CLContextProp                   = type intptr_t;
  CLDeviceType                    = type CLbitfield;
  CLCommandQueueProp              = type CLbitfield;
  CLMapFlags                      = type CLbitfield;
  CLMemFlags                      = type CLbitfield;
  PCLCommandQueueProp             = ^CLCommandQueueProp;

CLPlatformID,CLContext,CLDeviceID,CLMem, CLProgram,CLKernel,CLEvent : Pointer;//nativeint or nativeuint
как душе угодно,
флаги - обычный инт.

Добавлено через 2 часа 22 минуты
1мин 11сек на 7970 (ms:3200; iter:12) удачный код, рандом по памяти

Добавлено через 2 минуты
на таком код Fury с HBM бы затестить...

@Prok12 · 25.01.2016, 13:06 **[ТС]**

Сообщение от snake32

cmd.Finish;

Увидел в Вашем коде такую полезную строку. Насколько понимаю, она гарантирует, что расчётное ядро (например) будет выполнено ДО КОНЦА, прежде чем из буферов GPU начнут зачитываться рассчитанные им данные. Добавил в своём новом коде тоже после вызова расчётного ядра:
clEnqueueNDRangeKernel(CommandQueue,.... )
вызов:
clFinish(CommandQueue);
====
Попутно разобрался, чем всё-таки отличается clFinish(CommandQueue) от похожего "стопора" clEnqueueBarrier(CommandQueue). На нашинских форумах по этому вопросу - безсвязное мычание, на забугорных - уж больно сложно объясняют: http://stackoverflow.com/quest... d-clfinish. Из описания Khronos (ведёт OpenCL) выглядит так, что это одно и тоже, в обоих случаях - "synchronization point that ensures that all queued commands in command_queue have finished execution before the next batch of commands can begin execution"
https://www.khronos.org/regist... rrier.html
https://www.khronos.org/regist... inish.html
Объясню (пусть меня поправят) по рабоче-крестьянски. Для максимального ускорения выполнения команд, поставленных в очередь на GPU - это та самая CommandQueue - иногда эту очередь создают с флагом CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE (по умолчанию этого флага нет). То бишь, GPU разрешается выполнять поставленные в очередь к нему команды в ЛЮБОМ порядке (out-of-order), лишь бы очередь побыстрее продвигалась. Такой трюк проходит, например, если надо записать из CPU в память GPU десяток-другой больших массивов: соответственно, 10-20 OpenCL-команд, а в каком порядке их выполнять - пофиг, хоть все скопом. Но вот настало время запускать расчётное ядро-Kernel, а ему ВСЕ эти массивы нужны уже в законченном виде. Так что нельзя допустить, чтобы и команда запуска расчётного ядра "пролезла без очереди". Вот тогда перед вызовом расчётного ядра и понадобится вызов clEnqueueBarrier(CommandQueue). Это барьер разрешит выполнять в любом порядке только команды OpenCL до себя, и в любом прядке команды после себя. Но друг с дружкой мешать их в кучу не позволит. Так что, если флаг "ускорения out-of-order" при создании очереди - CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE - не был установлен, команда clEnqueueBarrier(CommandQueue) ничего не даст. Тут нужен clFinish(CommandQueue).
=====
Надеюсь, гуру по OpenCL меня поправят.

@snake32 · 25.01.2016, 14:19

Prok12, я далёк от гуру OpenCL. И пока сам до конца не разобрался, поэтому навтыкал clFinish, что бы уж наверняка ошибок out-of-order не было, и без этого дел пока хватает.
Кстати, есть ещё одна команда синхронизации clFlush. Отличие между с clFinish в том что с clFinish CPU явно ждёт выполнения всех команд, то есть эта ф-ия является блокирующей. Тогда как clFlush и clEnqueueBarrier сразу возвращают управление(imediately/nonblock).
А вот какое различие между clFlush и clEnqueueBarrier для меня загадка

@Prok12 · 26.01.2016, 10:53 **[ТС]**

Сообщение от snake32

А вот какое различие между clFlush и clEnqueueBarrier для меня загадка

С этим вроде разобрался:
1) clFlush - более "мягкий" аналог clFinish: clFlush "подталкивает" очередь GPU на выполнение одной команды, но не гарантирует (как clFlush), что вся очередь команд будет выполнена до конца. Иногда и команды clFlush достаточно, если только одна из наших OpenCL команд не блокирующая - как к примеру вызов расчётного ядра clEnqueueNDRangeKernel. Блокирующие команды - это те, которые пока не выполнятся до конца, дельфовый код (или на ином языке) дальше не пойдёт (CPU ждёт). Скажем, команды clEnqueueWriteBuffer(CommandQueue, ... , CL_TRUE,....) именно с флагом CL_TRUE - блокирующие.
2) Команда clEnqueueBarrier как и clEnqueueWaitForEvents интересны только тем, кто работает с экзотичной очередью CommandQueue - out-of-order: беспорядочное выполнение команд на GPU, для доп. ускорения - см. пост выше. До такой экзотики я ещё не дорос.

3) Команды clWaitForEvents мне вполне удаётся использовать: до какой-той команды (скажем, опять же запуска ядра на счёт) создаём своё событие командой clCreateUserEvent, сама команда (ей этот MyUserEvent передаётся, как параметр) сбрасывает после своего окончания флаг события. После команды ставим "стопор" clWaitForEvents : его параметр - указатель на массив ожидаемых событий. Ну примерно так же делается при синхронизации обычных выч.потоков на CPU (объектов TThread в Delphi). Короче, можно вместо всех этих событий, не мороча себе голову, тупо использовать clFinish или в ряде случаев - clFlush.
=====
Пока не смог врубиться, что такое "рабочие группы нитей (потоков)" на GPU, и нафига они нужны?
Размер этих групп я пока никак не задавал, но уже заметил, что ежели размер матрицы в тестовом примере по ссылке выше кратен 64 или 128, то выполнение идёт быстрее.

@snake32 · 26.01.2016, 12:11

Сообщение от Prok12

Скажем, команды clEnqueueWriteBuffer(CommandQueue, ... , CL_TRUE,....) именно с флагом CL_TRUE - блокирующие.

У команды clEnqueueReadBuffer тоже есть блокирующий флаг, но почему-то на nVidia у меня (год назад) этот флаг вообще ничего не давал. Как будто всегда CL_FALSE. И как следствие, попытка прочитать эти данные на CPU вызывало ошибку(хотя на AMD всё по стандарту). Поэтому опять юзал всемогущий clFinish.

Сообщение от Prok12

Пока не смог врубиться, что такое "рабочие группы нитей (потоков)" на GPU, и нафига они нужны?

Я их тоже пока не использую. Сложновато. Как я понял, идея групп в том, что потоки/нити внутри группы могут использовать общую память(кэш), объём которой в килобайтах(для CL 1.0 минимум 16Кб, более точно размер CL_DEVICE_LOCAL_MEM_SIZE). Теоретически, за счёт этого кэша можно значительно ускорить выполнение.

Добавлено через 4 минуты

Сообщение от Prok12

clFlush "подталкивает" очередь GPU на выполнение одной команды, но не гарантирует (как clFlush)

Тут наверное очепятка: clFinish

Что значит "подталкивает"? Это ж не человек которому можно дать пенделя.

Добавлено через 24 минуты
А если очередь out-of-order то тогда clFlush какую команду выполнит? Особенно если в очереди N команд, то какую именно: Первую, последнюю или среднюю?

@Prok12 · 26.01.2016, 19:57 **[ТС]**

Сообщение от snake32

Что значит "подталкивает"? Это ж не человек которому можно дать пенделя.

Полагаю, это как с обработкой сообщений в очереди Windows API:
1) Application.HandleMessage <=> одно сообщение обработать <=> аналог clFlush
2) Application.ProcessMessages <=> всю очередь сообщений обработать <=> аналог clFinish

Сообщение от snake32

Тут наверное очепятка: clFinish

Точно, опечатка!
===
Сейчас ломаю голову, почему скорость работы OpenCL-задач сильно зависит от малейшего изменения длины векторов и кол-ва задаваемых выч.потоков-нитей (GlobalThreads) ? Эмпирически установил (уточняю сейчас), что максимальная скорость - когда длина массивов из clFloat (4байта) и количество потоков кратно 32 (или 16??). Читал, что все данные, передаваемые в функции OpenCL, должны быть вручную выровнены по длине 16байт (именно байт). Вот что чехи (Брно) пишут в комментах к своему заголовочному файлу CL.pas:
{* Vector types
*
* Note: OpenCL requires that all types be naturally aligned.
* This means that vector types must be naturally aligned.
* For example, a vector of four floats must be aligned to
* a 16 byte boundary (calculated as 4 * the natural 4-byte
* alignment of the float). The alignment qualifiers here
* will only function properly if your compiler supports them
* and if you don't actively work to defeat them. For example,
* in order for a cl_float4 to be 16 byte aligned in a struct,
* the start of the struct must itself be 16-byte aligned.
*
* Maintaining proper alignment is the user's responsibility.
*}

@snake32 · 27.01.2016, 10:41

Сообщение от Prok12

скорость работы OpenCL-задач сильно зависит от малейшего изменения длины векторов

Сильно это на сколько? И какой CL-код исполняется?

Сообщение от Prok12

должны быть вручную выровнены по длине 16байт

Пример, как это сделать есть?

Новые блоги и статьи Все статьи Все блоги /
Новый ноутбук volvo 07.12.2025 Всем привет. По скидке в "черную пятницу" взял себе новый ноутбук Lenovo ThinkBook 16 G7 на Амазоне: Ryzen 5 7533HS 64 Gb DDR5 1Tb NVMe 16" Full HD Display Win11 Pro	Музыка, написанная Искусственным Интеллектом volvo 04.12.2025 Всем привет. Некоторое время назад меня заинтересовало, что уже умеет ИИ в плане написания музыки для песен, и, собственно, исполнения этих самых песен. Стихов у нас много, уже вышли 4 книги, еще 3. . .	От async/await к виртуальным потокам в Python IndentationError 23.11.2025 Армин Ронахер поставил под сомнение async/ await. Создатель Flask заявляет: цветные функции - провал, виртуальные потоки - решение. Не threading-динозавры, а новое поколение лёгких потоков. Откат?. . .	Поиск "дружественных имён" СОМ портов Argus19 22.11.2025 Поиск "дружественных имён" СОМ портов На странице: https:/ / norseev. ru/ 2018/ 01/ 04/ comportlist_windows/ нашёл схожую тему. Там приведён код на С++, который показывает только имена СОМ портов, типа,. . .	Сколько Государство потратило денег на меня, обеспечивая инсулином. Programma_Boinc 20.11.2025 Сколько Государство потратило денег на меня, обеспечивая инсулином. Вот решила сделать интересный приблизительный подсчет, сколько государство потратило на меня денег на покупку инсулинов. . . .
Ломающие изменения в C#.NStar Alpha Etyuhibosecyu 20.11.2025 Уже можно не только тестировать, но и пользоваться C#. NStar - писать оконные приложения, содержащие надписи, кнопки, текстовые поля и даже изображения, например, моя игра "Три в ряд" написана на этом. . .	Мысли в слух kumehtar 18.11.2025 Кстати, совсем недавно имел разговор на тему медитаций с людьми. И обнаружил, что они вообще не понимают что такое медитация и зачем она нужна. Самые базовые вещи. Для них это - когда просто люди. . .	Создание Single Page Application на фреймах krapotkin 16.11.2025 Статья исключительно для начинающих. Подходы оригинальностью не блещут. В век Веб все очень привыкли к дизайну Single-Page-Application . Быстренько разберем подход "на фреймах". Мы делаем одну. . .	Фото: Daniel Greenwood kumehtar 13.11.2025	Расскажи мне о Мире, бродяга kumehtar 12.11.2025 — Расскажи мне о Мире, бродяга, Ты же видел моря и метели. Как сменялись короны и стяги, Как эпохи стрелою летели. - Этот мир — это крылья и горы, Снег и пламя, любовь и тревоги, И бескрайние. . .

@Prok12 90 / 16 / 1 Регистрация: 08.11.2011 Сообщений: 96

	OpenCL в Delphi XE2 под Win64bit: можете протестировать на своей видеокарте? 21.01.2016, 13:04. Показов 15373. Ответов 44 Метки нет (Все метки) Суть проблемы. Есть весьма объёмный ПК для научных расчётов: интерфейс, База данных, много-поточные расчёты (на CPU пока), графика 2D, графика 3D на OpenGL, анимация результатов в виде компрессированного avi-файла на выходе - всё делалось на Delphi XE2...XE6 и прекрасно работает. Надо подключить модуль для переноса части длительных фрагментов расчёта на GPU. Этот модуль пытались писать на Visual Studio 2015 - расширение языка C++ AMP, как DLL: работает, но криво: не на всех видеокартах, увы (текст тестовой программы выложен в той же папке для скачивания - см.ниже). Ну и вечная проблема с недостающими DLL к тому что написано на Visual Studio: даже опции компилятора /MT не помогают... На JAVA тоже сейчас мои коллеги пишут этот расчётный модуль - через OpenCL... Ну не в этом вопрос. Есть примеры использования библиотеки OpenCL.dll прямо в проектах Delphi (чехи, например, делали, есть ещё компоненты от MITOV для Delphi Seattle - http://www.mitov.com/ ). Эти несложные примеры работают, но...только при компиляции под Win32, а компилируешь проект Delphi под Win64 - затык... ----- Сейчас вроде с помощью советов с другого форума, удалось траблы победить, и сделать работоспособный тест. ----- Вот ссылка для скачивания: http://gofile.me/2Zesj/C0f3wb1o ======== *Там в папке:* ======== 1) Полезная утилита GPU_Caps_Viewer_Setup_v1.25.0.0.exe для контроля состояния видеокарты, особенно, если нажать кнопку "More GPU Info". 2) OpenCL1_2_Delphi.zip - файл с примером от чехов (Университет Брно, насколько помню) использования OpenCL в Delphi, 2013год. У меня этот пример не всегда корректно работал при компиляции под Win64 (не на всех видео-картах). Но там хороший заголовочный файл CL.pas - его можно чуть доработать с учётом замечания про NativeInt с другого форума: заменить тип size_t на NativeInt для Win64, и использовать вместо моего "укороченного" MyOpenCL.pas - см.ниже. --- 3) Вложенная папка с моим вариантом программы-теста. Она самодостаточна: больше ничего не надо, ну разве что последние версии видео-драйверов установить. Тест сделан под Delphi XE2 (работает и под XE6). Причём работает, по крайней мере у меня, при компиляции под Win32, и что более важно- под Win64. Всего 2 файла: --> MyOpenCL.pas: заголовочный; я его урезал - убрал ненужные мне функции, которые в принципе можно взять из файла чехов CL.pas + учёл полезное замечание с др.форма про NativeInt ; --> FMain.pas : тестовая программа. Выделяет на GPU память (создаёт буферы) под 16 массивов (квадратных матриц из cl_Float, размер стороны которых задаётся на основной форме, по умолчанию 1200) и под ещё один массив, того же размера, куда пишется результат несложных вычислений на GPU. Потом запускается итерационный цикл (количество итераций - для проверки времени работы - задаётся тоже на форме, по умолчанию - 12). Внутри каждой итерации 16 массивов заполняются данными, затем запускается ещё и внутренний цикл: 30 проходов вызова расчётного ядра на GPU. Такая структура тестовой программы больше всего подходит под наши научные задачи (вычислит. гидродинамика), но пока это лишь тест. В тесте есть несложная проверка правильности вычислений - для одного из элементов матрицы-результата, поскольку НЕ все драйверы видеокарт (особенно старые) могут диагностировать Kernel Error. --> программа на С для ядра - текстовый файл ProgramGPU.CL (должен лежать рядом с MyFirstOpenCL.exe !!): там можно глянуть, что делается с одномерными массивами на GPU. Она прямо передаётся в виде строки на GPU, там копилируется и линкуется - ну как обычно в OpenCL. Нужная DLL для OpenCL - под Win64 или Win32 - подключится из системных папок Windows автоматически, в зависимости от компиляции. ---- 4) Можете на своём сетапе запустить прямо готовый MyFirstOpenCL.exe (это версия 64bit !!) с параметрами: --> размер стороны матрицы = 3200; лезут 17 таких матриц на видеокарту от 1Gb и выше; если не влезут (часть видео-памяти уже занята чем-то - GPU_Caps_Viewer в помощь) - возможны ошибки, которые диагностируются НЕ всем драйверами видеокарт, а только новыми; --> количество итераций (внешний цикл) = 12. --- 5) У меня для указанного примера время счёта вышло: NVIDIA GT-430 (1Gb, 96core, 700MHz core) - 3мин 37сек; NVIDIA GTS-450 (1Gb, 192core, 810MHz core) - 1мин 36сек; NVIDIA GTX-570 (1.28Gb, 480core, 732MHz core) - 42.5сек; AMD Radeon R9 M275X (2Gb, 640core, 925MHz core) - 1мин 27сек - настольный моноблок ASUS 2702. ---- Может кто-то попробовать MyFirstOpenCL.exe на своих видеокартах??? Поскольку это проект дельфовый, никакие внешние DLL ему не нужны (кроме тех что в Windows у всех есть), как и всякие - джава-машины: только для видео-карты новые драйверы желательны. Запустить, нажать верхнюю кнопку на форме, установить размер матрицы 3200 (если 1200 отработает корректно), выбрать платформу-девайс, нажать нижнюю кнопку на форме. Записать время счёта: оно выдаётся в итоговом окне. Нам сейчас на работу надо будет покупать новую карту 4-6Gb именно для расчётов...денег более 60тр вряд ли выделят... AMD конечно подешевле, но... Такие сопоставительные тесты на задачах, приближенных к нашим, были бы полезны!! Были бы полезны и замечания участников форума по тестовой программе. === К сожалению, у меня нет под рукой видео-карты (акселератора) с памятью на борту 4Gb-12Gb. Было бы интересно протестить и на таких "монстрах", повышая понемногу размер входной матрицы (требуемую память тестовая программа сообщит в начале работы): для того и писалось под Win64 - нужна большая память. Буду признателен! 0

@FIL Модератор 3492 / 2614 / 742 Регистрация: 19.09.2012 Сообщений: 7,977
	21.01.2016, 20:58
	Запустил программу на компе с процессором i7-4790K со встроенной видяхой. Определилось 2 устройства: процессор и видяха. При выборе видяхи вылетела ошибка. При выборе процессора программа отработала за 1 мин. 36 сек. Я так понимаю, что при этом вычисления ведутся в режиме эмуляции видеокарты. Соответственно, если данный алгоритм реализовать на SSE-регистрах процессора, то вероятно он будет работать значительно быстрее. Миниатюры 0

@FIL Модератор 3492 / 2614 / 742 Регистрация: 19.09.2012 Сообщений: 7,977
	21.01.2016, 21:56
	Название и сообщение об ошибке на скрине. Ошибка возникает при любом размере матрицы. Судя по тестам под видяху выделяется 1Гб оперативки. Миниатюры 0

@x128 240 / 218 / 46 Регистрация: 17.04.2010 Сообщений: 526
	22.01.2016, 10:49
	Запустил на i5-6600, тоже была ошибка как у FIL, после удаления строчки с "cl_khr_fp64" процесс пошел. результат для GPU: результат для CPU: 1

@snake32 3502 / 1685 / 236 Регистрация: 26.02.2009 Сообщений: 8,368 Записей в блоге: 6
	22.01.2016, 12:34
	Тестили как-то двойную точность. Если интересны результаты: https://www.cyberforum.ru/vide... 42529.html 0

@FIL Модератор 3492 / 2614 / 742 Регистрация: 19.09.2012 Сообщений: 7,977
	22.01.2016, 19:58
	Без cl_khr_fp64 получился такой результат: Миниатюры 1

@snake32 3502 / 1685 / 236 Регистрация: 26.02.2009 Сообщений: 8,368 Записей в блоге: 6
	25.01.2016, 14:19
	Prok12, я далёк от гуру OpenCL. И пока сам до конца не разобрался, поэтому навтыкал clFinish, что бы уж наверняка ошибок out-of-order не было, и без этого дел пока хватает. Кстати, есть ещё одна команда синхронизации clFlush. Отличие между с clFinish в том что с clFinish CPU явно ждёт выполнения всех команд, то есть эта ф-ия является блокирующей. Тогда как clFlush и clEnqueueBarrier сразу возвращают управление(imediately/nonblock). А вот какое различие между clFlush и clEnqueueBarrier для меня загадка 0