Merge pull request #1942 from SunSerega/Open-CL-GL

Модуль для OpenCL и простенький пример к нему

InstallerSamples/Open[CL,GL]/OpenCL/MatrMlt.cl

@@ -0,0 +1,29 @@
+
+
+
+__kernel void MatrMltMatr(__global double* A, __global double* B, __global double* C, __global int* gW)
+{
+	int cX = get_global_id(0);
+	int cY = get_global_id(1);
+	int W = *gW;
+	
+	double sum = 0.0;
+	for (int i=0; i<W; i++)
+		sum += A[i + cX*W] * B[cY + i*W];
+	
+	C[cX + cY*W] = sum;
+}
+
+__kernel void MatrMltVec(__global double* C, __global double* V, __global double* V2, __global int* gW)
+{
+	int i = get_global_id(0);
+	int W = *gW;
+	
+	double sum = 0.0;
+	for (int j=0; j<W; j++)
+		sum += C[j + i*W] * V[j];
+	
+	V2[i] = sum;
+}
+
+

InstallerSamples/Open[CL,GL]/OpenCL/MatrMlt.pas

@@ -0,0 +1,134 @@
+uses OpenCL;
+uses System;
+uses System.Runtime.InteropServices;
+
+const
+  MatrW = 4; // можно поменять на любое положительное значение
+  
+  MatrL = MatrW*MatrW;
+  MatrByteSize = MatrL*8;
+
+procedure MatrToUnmanagedArr(m: array[,] of real; mem: IntPtr);
+begin
+  var a := new real[m.Length];
+  Buffer.BlockCopy(m,0, a,0, a.Length*8); // вообще не хорошо, лучше через GCHandle, тогда копировать не надо ничего
+  Marshal.Copy(a,0, mem,a.Length);
+end;
+
+procedure UnmanagedArrToMatr(mem: IntPtr; m: array[,] of real);
+begin
+  var a := new real[m.Length];
+  Marshal.Copy(mem,a, 0,a.Length);
+  Buffer.BlockCopy(a,0, m,0, a.Length*8);
+end;
+
+begin
+  var ec: ErrorCode;
+  
+  // Инициализация
+  
+  var platform: cl_platform_id;
+  cl.GetPlatformIDs(1, @platform, nil).RaiseIfError;
+  
+  var device: cl_device_id;
+  cl.GetDeviceIDs(platform, DeviceTypeFlags.Default, 1, @device, nil).RaiseIfError;
+  
+  var context := cl.CreateContext(nil, 1, @device, nil, nil, @ec);
+  ec.RaiseIfError;
+  
+  var command_queue := cl.CreateCommandQueue(context, device, CommandQueuePropertyFlags.NONE, ec);
+  ec.RaiseIfError;
+  
+  // Чтение и компиляция .cl файла
+  
+  {$resource MatrMlt.cl}
+  var prog_str := System.IO.StreamReader.Create(GetResourceStream('MatrMlt.cl')).ReadToEnd;
+  var prog := cl.CreateProgramWithSource(
+    context,
+    1,
+    new string[](prog_str),
+    new UIntPtr[](new UIntPtr(prog_str.Length)),
+    ec
+  );
+  ec.RaiseIfError;
+  
+  cl.BuildProgram(prog, 1, @device, nil, nil, nil).RaiseIfError;
+  
+  var MatrMltMatrKernel := cl.CreateKernel(prog, 'MatrMltMatr', ec);
+  ec.RaiseIfError;
+  
+  var MatrMltVecKernel := cl.CreateKernel(prog, 'MatrMltVec', ec);
+  ec.RaiseIfError;
+  
+  // Подготовка параметров
+  
+  writeln('Матрица A:');
+  var A := MatrRandomReal(MatrW,MatrW,0,1).Println;
+  writeln;
+  var Amem := Marshal.AllocHGlobal(MatrByteSize);
+  MatrToUnmanagedArr(A,Amem);
+  var AmemObj := cl.CreateBuffer(context, MemoryFlags.READ_WRITE or MemoryFlags.USE_HOST_PTR, new UIntPtr(MatrByteSize), Amem, ec);
+  ec.RaiseIfError;
+  
+  writeln('Матрица B:');
+  var B := MatrRandomReal(MatrW,MatrW,0,1).Println;
+  writeln;
+  var Bmem := Marshal.AllocHGlobal(MatrByteSize);
+  MatrToUnmanagedArr(B,Bmem);
+  var BmemObj := cl.CreateBuffer(context, MemoryFlags.READ_WRITE or MemoryFlags.USE_HOST_PTR, new UIntPtr(MatrByteSize), Bmem, ec);
+  ec.RaiseIfError;
+  
+  writeln('Вектор V:');
+  var V := ArrRandomReal(MatrW);
+  V.Println;
+  writeln;
+  var Vmem := Marshal.AllocHGlobal(MatrW*8);
+  Marshal.Copy(V,0, Vmem,MatrW);
+  var VmemObj := cl.CreateBuffer(context, MemoryFlags.READ_WRITE or MemoryFlags.USE_HOST_PTR, new UIntPtr(MatrW*8), Vmem, ec);
+  ec.RaiseIfError;
+  
+  var CmemObj := cl.CreateBuffer(context, MemoryFlags.READ_WRITE, new UIntPtr(MatrByteSize), nil, @ec);
+  ec.RaiseIfError;
+  
+  var V2memObj := cl.CreateBuffer(context, MemoryFlags.READ_WRITE, new UIntPtr(MatrW*8), nil, @ec);
+  ec.RaiseIfError;
+  
+  var MatrWParam := MatrW;
+  var WmemObj := cl.CreateBuffer(context, MemoryFlags.READ_WRITE or MemoryFlags.USE_HOST_PTR, new UIntPtr(4), @MatrWParam, @ec);
+  ec.RaiseIfError;
+  
+  // Выполнение C := A*B
+  
+  cl.SetKernelArg(MatrMltMatrKernel, 0, new UIntPtr(UIntPtr.Size), AmemObj).RaiseIfError;
+  cl.SetKernelArg(MatrMltMatrKernel, 1, new UIntPtr(UIntPtr.Size), BmemObj).RaiseIfError;
+  cl.SetKernelArg(MatrMltMatrKernel, 2, new UIntPtr(UIntPtr.Size), CmemObj).RaiseIfError;
+  cl.SetKernelArg(MatrMltMatrKernel, 3, new UIntPtr(UIntPtr.Size), WmemObj).RaiseIfError;
+  
+  cl.EnqueueNDRangeKernel(command_queue, MatrMltMatrKernel, 2, nil,new UIntPtr[](new UIntPtr(MatrW),new UIntPtr(MatrW)),nil, 0,nil,nil).RaiseIfError;
+  
+  // Выполнение V2 := C*V
+  
+  cl.SetKernelArg(MatrMltVecKernel, 0, new UIntPtr(UIntPtr.Size), CmemObj).RaiseIfError;
+  cl.SetKernelArg(MatrMltVecKernel, 1, new UIntPtr(UIntPtr.Size), VmemObj).RaiseIfError;
+  cl.SetKernelArg(MatrMltVecKernel, 2, new UIntPtr(UIntPtr.Size), V2memobj).RaiseIfError;
+  cl.SetKernelArg(MatrMltVecKernel, 3, new UIntPtr(UIntPtr.Size), WmemObj).RaiseIfError;
+  
+  cl.EnqueueNDRangeKernel(command_queue, MatrMltVecKernel, 1, nil,new UIntPtr[](new UIntPtr(MatrW)),nil, 0,nil,nil).RaiseIfError;
+  
+  // Чтение и вывод результата
+  
+  cl.EnqueueReadBuffer(command_queue, CmemObj,  0, new UIntPtr(0), new UIntPtr(MatrByteSize), Amem, 0,nil,nil).RaiseIfError;
+  cl.EnqueueReadBuffer(command_queue, V2memObj, 0, new UIntPtr(0), new UIntPtr(MatrW*8),      Vmem, 0,nil,nil).RaiseIfError;
+  
+  cl.Finish(command_queue).RaiseIfError;
+  
+  writeln('Матрица С = A*B:');
+  UnmanagedArrToMatr(Amem,A);
+  A.Println;
+  writeln;
+  
+  writeln('Вектор V2 = C*V:');
+  Marshal.Copy(Vmem,V,0,MatrW);
+  V.Println;
+  
+end.

InstallerSamples/Open[CL,GL]/OpenCL/SimpleAddition.cl

@@ -0,0 +1,11 @@
+
+
+
+__kernel void TEST(__global int* message)
+{
+	int gid = get_global_id(0);
+	
+	message[gid] += gid;
+}
+
+

InstallerSamples/Open[CL,GL]/OpenCL/SimpleAddition.pas

@@ -0,0 +1,64 @@
+uses OpenCL;
+uses System;
+uses System.Runtime.InteropServices;
+
+//Описания всех подпрограмм найдёте в справке по OpenCL:
+//www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_API.html
+
+begin
+  var ec: ErrorCode;
+  
+  // Инициализация
+  
+  var platform: cl_platform_id;
+  cl.GetPlatformIDs(1, @platform, nil).RaiseIfError;
+  
+  var device: cl_device_id;
+  cl.GetDeviceIDs(platform, DeviceTypeFlags.Default, 1, @device, nil).RaiseIfError;
+  
+  var context := cl.CreateContext(nil, 1, @device, nil, nil, @ec);
+  ec.RaiseIfError;
+  
+  var command_queue := cl.CreateCommandQueue(context, device, CommandQueuePropertyFlags.NONE, ec);
+  ec.RaiseIfError;
+  
+  // Чтение и компиляция .cl файла
+  
+  {$resource SimpleAddition.cl} // эта строчка засовывает SimpleAddition.cl внутрь .exe, чтоб он не нужен был для запуска .exe
+  var prog_str := System.IO.StreamReader.Create(GetResourceStream('SimpleAddition.cl')).ReadToEnd;
+  var prog := cl.CreateProgramWithSource(
+    context,
+    1,
+    new string[](prog_str),
+    new UIntPtr[](new UIntPtr(prog_str.Length)),
+    ec
+  );
+  ec.RaiseIfError;
+  
+  cl.BuildProgram(prog, 1, @device, nil, nil, nil).RaiseIfError;
+  
+  // Подготовка и запуск программы на GPU
+  
+  var kernel := cl.CreateKernel(prog, 'TEST', ec); // Обязательно то же имя что у карнела из .cl файла. И регистр важен!
+  ec.RaiseIfError;
+  
+  var mem := Marshal.AllocHGlobal(40);
+  Marshal.Copy(ArrFill(10,1),0,mem,10);
+  var memobj := cl.CreateBuffer(context, MemoryFlags.READ_WRITE or MemoryFlags.USE_HOST_PTR, new UIntPtr(40), mem, ec); // USE_HOST_PTR значит что нужно скопировать память из mem в memobj
+  ec.RaiseIfError;
+  
+  cl.SetKernelArg(kernel, 0, new UIntPtr(UIntPtr.Size), memobj).RaiseIfError;
+  
+  cl.EnqueueNDRangeKernel(command_queue, kernel, 1, nil,new UIntPtr[](new UIntPtr(10)),nil, 0,nil,nil).RaiseIfError;
+  
+  cl.Finish(command_queue).RaiseIfError;
+  
+  // Чтение и вывод результата
+  
+  cl.EnqueueReadBuffer(command_queue, memobj, 1, new UIntPtr(0), new UIntPtr(40), pointer(mem), 0,nil,nil).RaiseIfError;
+  
+  var res := new integer[10];
+  Marshal.Copy(mem,res,0,10);
+  res.Println;
+  
+end.

ReleaseGenerators/RebuildStandartModules.pas

@@ -1,21 +1,22 @@
-uses 
-   PABCSystem, PABCExtensions, __RedirectIOMode, __RunMode, ABCButtons, ABCHouse, ABCObjects,
-   ABCSprites, CRT, DMCollect, DMTaskMaker, DMZadan, Drawman, 
-   DrawManField, Events, FilesOperations, 
-   GraphABC,
-   GraphABCHelper,  
-   GraphWPFBase,
-   Graph3D,
-   GraphWPF,
-   WPFObjects,
-   NumLibABC,
-   IniFile, PointerTools, PointRect, PT4, PT4MakerNetX, Robot, RobotField,
-   RobotTaskMaker, RobotZadan, Sockets, Utils, VCL, Timers, PT4Exam, PT4TaskMakerNET, RBDMUtils,
-   Collections, Arrays, Core, FormsABC, MPI, ClientServer, OpenGL, Speech, Sounds, Countries,
-   ABCDatabases,BlockFileOfT,Controls
-   ;
-   
-begin
-  writeln(cos(pi));
-  readln;
+uses 
+   PABCSystem, PABCExtensions, __RedirectIOMode, __RunMode, ABCButtons, ABCHouse, ABCObjects,
+   ABCSprites, CRT, DMCollect, DMTaskMaker, DMZadan, Drawman, 
+   DrawManField, Events, FilesOperations, 
+   GraphABC,
+   GraphABCHelper,  
+   GraphWPFBase,
+   Graph3D,
+   GraphWPF,
+   WPFObjects,
+   NumLibABC,
+   IniFile, PointerTools, PointRect, PT4, PT4MakerNetX, Robot, RobotField,
+   RobotTaskMaker, RobotZadan, Sockets, Utils, VCL, Timers, PT4Exam, PT4TaskMakerNET, RBDMUtils,
+   Collections, Arrays, Core, FormsABC, MPI, ClientServer, OpenGL, Speech, Sounds, Countries,
+   ABCDatabases,BlockFileOfT,Controls,
+   OpenCL
+   ;
+   
+begin
+  writeln(cos(pi));
+  readln;
 end.

ReleaseGenerators/sect_Core.nsh

@@ -127,6 +127,7 @@
     File ..\bin\Lib\Speech.pcu
     File ..\bin\Lib\Sounds.pcu
     File ..\bin\Lib\BlockFileOfT.pcu
+    File ..\bin\Lib\OpenCL.pcu
 
     File ..\bin\Lib\PABCRtl.dll
     File ..\bin\Lib\PABCRtl32.dll
@@ -189,6 +190,7 @@
     ${AddFile} "Speech.pcu"
     ${AddFile} "Sounds.pcu"
     ${AddFile} "BlockFileOfT.pcu"
+    ${AddFile} "OpenCL.pcu"
     ${AddFile} "PABCRtl.dll"
     ${AddFile} "HelixToolkit.Wpf.dll"
     ${AddFile} "HelixToolkit.dll"
@@ -260,6 +262,7 @@
     File ..\bin\Lib\Speech.pas
     File ..\bin\Lib\Sounds.pas
     File ..\bin\Lib\BlockFileOfT.pas
+    File ..\bin\Lib\OpenCL.pas
 	File ..\bin\Lib\__RedirectIOMode.vb
 	File ..\bin\Lib\VBSystem.vb
 	
@@ -317,6 +320,7 @@
     ${AddFile} "Speech.pas"
     ${AddFile} "Sounds.pas"
     ${AddFile} "BlockFileOfT.pas"
+    ${AddFile} "OpenCL.pas"
 	${AddFile} "__RedirectIOMode.vb"
     ${AddFile} "VBSystem.vb"