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

Обновление POCGL (в основном OpenCLABC)
2020-08-09 18:08:20 +03:00 · 2020-08-09 18:08:20 +03:00 · f697fc8be6
parent c0906732a5 ce0ff7113a
commit f697fc8be6
21 changed files with 11634 additions and 5988 deletions
--- a/OpenCL/OpenCL/MatrMlt.cl
+++ b/OpenCL/OpenCL/MatrMlt.cl
@ -1,11 +1,10 @@



-__kernel void MatrMltMatr(__global double* A, __global double* B, __global double* C, __global int* gW)
+__kernel void MatrMltMatr(__global double* A, __global double* B, __global double* C, int W)
 {
 	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++)
@ -14,10 +13,9 @@ __kernel void MatrMltMatr(__global double* A, __global double* B, __global doubl
 	C[cX + cY*W] = sum;
 }

-__kernel void MatrMltVec(__global double* C, __global double* V1, __global double* V2, __global int* gW)
+__kernel void MatrMltVec(__global double* C, __global double* V1, __global double* V2, int W)
 {
 	int i = get_global_id(0);
-	int W = *gW;
 	
 	double sum = 0.0;
 	for (int j=0; j<W; j++)
--- a/OpenCL/OpenCL/MatrMlt.pas
+++ b/OpenCL/OpenCL/MatrMlt.pas
@ -19,8 +19,7 @@ begin
  cl.GetPlatformIDs(1, platform, IntPtr.Zero).RaiseIfError;
  
  var device: cl_device_id;
-  cl.GetDeviceIDs(platform, DeviceType.DEVICE_TYPE_DEFAULT, 1,device,IntPtr.Zero).RaiseIfError;
-//  cl.GetDeviceIDs(platform, DeviceType.DEVICE_TYPE_ALL, 1,device,IntPtr.Zero).RaiseIfError;
+  cl.GetDeviceIDs(platform, DeviceType.DEVICE_TYPE_ALL, 1,device,IntPtr.Zero).RaiseIfError;
  
  var context := cl.CreateContext(nil, 1,device, nil,IntPtr.Zero, ec);
  ec.RaiseIfError;
@ -31,7 +30,9 @@ begin
  // Чтение и компиляция .cl файла
  
  {$resource MatrMlt.cl}
-  var prog_str := System.IO.StreamReader.Create(GetResourceStream('MatrMlt.cl')).ReadToEnd;
+  var prog_str := System.IO.StreamReader.Create(
+    System.Reflection.Assembly.GetCallingAssembly.GetManifestResourceStream('MatrMlt.cl')
+  ).ReadToEnd;
  var prog := cl.CreateProgramWithSource(
    context,
    1,
@ -80,15 +81,13 @@ begin
  ec.RaiseIfError;
  
  var MatrWParam := MatrW;
-  var W_buf := cl.CreateBuffer(context, MemFlags.MEM_READ_WRITE or MemFlags.MEM_USE_HOST_PTR, new UIntPtr(sizeof(integer)),new IntPtr(@MatrWParam), ec);
-  ec.RaiseIfError;
  
  // Выполнение C := A*B
  
  cl.SetKernelArg(MatrMltMatrKernel, 0, new UIntPtr(UIntPtr.Size), A_buf).RaiseIfError;
  cl.SetKernelArg(MatrMltMatrKernel, 1, new UIntPtr(UIntPtr.Size), B_buf).RaiseIfError;
  cl.SetKernelArg(MatrMltMatrKernel, 2, new UIntPtr(UIntPtr.Size), C_buf).RaiseIfError;
-  cl.SetKernelArg(MatrMltMatrKernel, 3, new UIntPtr(UIntPtr.Size), W_buf).RaiseIfError;
+  cl.SetKernelArg(MatrMltMatrKernel, 3, new UIntPtr(sizeof(integer)), @MatrWParam).RaiseIfError;
  
  var k1_ev: cl_event;
  cl.EnqueueNDRangeKernel(command_queue, MatrMltMatrKernel, 2, nil,new UIntPtr[](new UIntPtr(MatrW),new UIntPtr(MatrW)),nil, 0,nil,k1_ev).RaiseIfError;
@ -98,7 +97,7 @@ begin
  cl.SetKernelArg(MatrMltVecKernel, 0, new UIntPtr(UIntPtr.Size),  C_buf).RaiseIfError;
  cl.SetKernelArg(MatrMltVecKernel, 1, new UIntPtr(UIntPtr.Size), V1_buf).RaiseIfError;
  cl.SetKernelArg(MatrMltVecKernel, 2, new UIntPtr(UIntPtr.Size), V2_buf).RaiseIfError;
-  cl.SetKernelArg(MatrMltVecKernel, 3, new UIntPtr(UIntPtr.Size),  W_buf).RaiseIfError;
+  cl.SetKernelArg(MatrMltVecKernel, 3, new UIntPtr(sizeof(integer)), @MatrWParam).RaiseIfError;
  
  var k2_ev: cl_event;
  cl.EnqueueNDRangeKernel(command_queue, MatrMltVecKernel, 1, nil,new UIntPtr[](new UIntPtr(MatrW)),nil, 1,k1_ev,k2_ev).RaiseIfError;
@ -113,7 +112,7 @@ begin
  A.Println;
  Writeln;
  
-  cl.EnqueueReadBuffer(command_queue, V2_buf, Bool.NON_BLOCKING, new UIntPtr(0), new UIntPtr(VecByteSize), V1[0], 1,k2_ev,IntPtr.Zero).RaiseIfError;
+  cl.EnqueueReadBuffer(command_queue, V2_buf, Bool.BLOCKING, new UIntPtr(0), new UIntPtr(VecByteSize), V1[0], 1,k2_ev,IntPtr.Zero).RaiseIfError;
  Writeln('Вектор V2 = C*V1:');
  V1.Println;
  
--- a/OpenCL/OpenCL/SimpleAddition.pas
+++ b/OpenCL/OpenCL/SimpleAddition.pas
@ -15,15 +15,8 @@ begin
  cl.GetPlatformIDs(1, platform, IntPtr.Zero).RaiseIfError;
  
  var device: cl_device_id;
-  cl.GetDeviceIDs(platform, DeviceType.DEVICE_TYPE_DEFAULT, 1,device,IntPtr.Zero).RaiseIfError;
-  
-  // DEVICE_TYPE_DEFAULT это, обычно, GPU
-  // Но, к примеру, в ноутбуке его может не быть
-  // Тогда надо хоть для чего то попытаться инициализировать
-  // DEVICE_TYPE_ALL выберет первое любое устройство, поддерживающее OpenCL
-//  cl.GetDeviceIDs(platform, DeviceType.DEVICE_TYPE_ALL, 1,device,IntPtr.Zero).RaiseIfError;
-  // Если всё ещё пишет что устройств нет - обновите драйверы,
-  // потому что даже встроенные видеокарты поддерживают OpenCL
+  cl.GetDeviceIDs(platform, DeviceType.DEVICE_TYPE_ALL, 1,device,IntPtr.Zero).RaiseIfError;
+  // Если пишет что устройств нет - обновите драйверы
  
  var context := cl.CreateContext(nil, 1,device, nil,IntPtr.Zero, ec);
  ec.RaiseIfError;
@ -34,7 +27,9 @@ begin
  // Чтение и компиляция .cl файла
  
  {$resource SimpleAddition.cl} // эта строчка засовывает SimpleAddition.cl внутрь .exe, чтоб не надо было таскать его вместе с .exe
-  var prog_str := System.IO.StreamReader.Create(GetResourceStream('SimpleAddition.cl')).ReadToEnd;
+  var prog_str := System.IO.StreamReader.Create(
+    System.Reflection.Assembly.GetCallingAssembly.GetManifestResourceStream('SimpleAddition.cl')
+  ).ReadToEnd;
  var prog := cl.CreateProgramWithSource(
    context,
    1,
--- a/OpenCL/OpenCLABC/0Простейшие
+++ b/OpenCL/OpenCLABC/0Простейшие
@ -1,11 +1,10 @@



-__kernel void MatrMltMatr(__global double* A, __global double* B, __global double* C, __global int* gW)
+__kernel void MatrMltMatr(__global double* A, __global double* B, __global double* C, int W)
 {
 	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++)
@ -14,10 +13,9 @@ __kernel void MatrMltMatr(__global double* A, __global double* B, __global doubl
 	C[cX + cY*W] = sum;
 }

-__kernel void MatrMltVec(__global double* C, __global double* V1, __global double* V2, __global int* gW)
+__kernel void MatrMltVec(__global double* C, __global double* V1, __global double* V2, int W)
 {
 	int i = get_global_id(0);
-	int W = *gW;
 	
 	double sum = 0.0;
 	for (int j=0; j<W; j++)
--- a/OpenCL/OpenCLABC/0Простейшие
+++ b/OpenCL/OpenCLABC/0Простейшие
@ -19,9 +19,9 @@ begin
    // Вообще лучше прекомпилировать .cl файл (загружать в переменную ProgramCode)
    // И сохранять с помощью метода ProgramCode.SerializeTo
    // А полученный бинарник уже подключать через $resource
-    var code := new ProgramCode(
+    var code := new ProgramCode(Context.Default,
      System.IO.StreamReader.Create(
-        GetResourceStream('MatrMlt.cl')
+        System.Reflection.Assembly.GetExecutingAssembly.GetManifestResourceStream('MatrMlt.cl')
      ).ReadToEnd
    );
    
@ -47,7 +47,7 @@ begin
    
    var V2 := new Buffer(VecByteSize);
    
-    var N := new Buffer(4);
+    var W := KernelArg.FromRecord(MatrW);
    
    // (запись значений в параметры - позже, в очередях)
    
@ -55,14 +55,14 @@ begin
    
    var Calc_C_Q :=
      code['MatrMltMatr'].NewQueue.AddExec2(MatrW, MatrW, // Выделяем ядра в форме квадрата, всего MatrW*MatrW ядер
-        A.NewQueue.AddWriteArray(A_Matr) as CommandQueue<Buffer>,
-        B.NewQueue.AddWriteArray(B_Mart) as CommandQueue<Buffer>,
+        A.NewQueue.AddWriteArray2&<real>(A_Matr),
+        B.NewQueue.AddWriteArray2&<real>(B_Mart),
        C,
-        N.NewQueue.AddWriteValue(MatrW) as CommandQueue<Buffer>
+        W
      ) as CommandQueue<Kernel>;
    
    var Otp_C_Q :=
-      C.NewQueue.AddReadArray(A_Matr) as CommandQueue<Buffer> +
+      C.NewQueue.AddReadArray2&<real>(A_Matr) +
      HPQ(()->
      lock output do
      begin
@ -74,13 +74,13 @@ begin
    var Calc_V2_Q :=
      code['MatrMltVec'].NewQueue.AddExec1(MatrW,
        C,
-        V1.NewQueue.AddWriteArray(V1_Arr) as CommandQueue<Buffer>,
+        V1.NewQueue.AddWriteArray1&<real>(V1_Arr),
        V2,
-        N // значение записывается в Calc_C_Q, тут можно использовать уже готовое
+        W
      ) as CommandQueue<Kernel>;
    
    var Otp_V2_Q :=
-      V2.NewQueue.AddReadArray(V1_Arr) as CommandQueue<Buffer> +
+      V2.NewQueue.AddReadArray1&<real>(V1_Arr) +
      HPQ(()->
      lock output do
      begin
--- a/OpenCL/OpenCLABC/0Простейшие
+++ b/OpenCL/OpenCLABC/0Простейшие
@ -4,23 +4,24 @@ begin
  
  // Чтение и компиляция .cl файла
  
-  var prog := new ProgramCode(ReadAllText('SimpleAddition.cl'));
+  var prog := new ProgramCode(Context.Default, ReadAllText('SimpleAddition.cl'));
  
  // Подготовка параметров
  
-  var A := new Buffer( 10 * sizeof(integer) ); // Буфер достаточно размера, чтоб хранить 10 элементов типа "integer"
+  var A := new Buffer( 10 * sizeof(integer) ); // Буфер достаточного размера, чтобы хранить 10 элементов типа "integer"
  
  // Выполнение
  
  prog['TEST'].Exec1(10, // используем 10 ядер
    
-    A.NewQueue.AddFillValue(1) // заполняем весь буфер значениями (1), прямо перед выполнением. Тип значения integer, потому что это тип по-умолчанию для целых чисел
-    as CommandQueue<Buffer> //ToDo нужно только из за issue компилятора #1981, иначе получаем странную ошибку. Когда исправят - можно будет убрать
+    // Заполняем весь буфер значениями (1), прямо перед выполнением
+    // Тип значения integer, потому что это тип по-умолчанию для целых чисел
+    A.NewQueue.AddFillValue(1)
    
  );
  
  // Чтение и вывод результата
  
-  A.GetArray1&<integer>.Println; // читаем всё содержимое как одномерный массив с элементами типа "integer"
+  A.GetArray1&<integer>.Println; // Читаем всё содержимое как одномерный массив с элементами типа "integer"
  
 end.
--- a/InstallerSamples/OpenGL
+++ b/InstallerSamples/OpenGL
@ -4,24 +4,28 @@ uses NamedQData;
 begin
  // График выполнения очередей:
  //
-  // A1--B1--C1------D1
+  // A1--B1--C1------E1
  //       \        /
  //        \      /
-  // A2------B2--C2--D2
+  // A2------C2--D2--E2
  //
  
  var A1 := NamedQ('A1');
-  var B1 := NamedQ('B1');
-  var C1 := NamedQ('C1');
-  var D1 := NamedQ('D1');
  var A2 := NamedQ('A2');
-  var B2 := NamedQ('B2');
+  
+  var B1 := NamedQ('B1');
+  
+  var C1 := NamedQ('C1');
  var C2 := NamedQ('C2');
+  
  var D2 := NamedQ('D2');
  
+  var E1 := NamedQ('E1');
+  var E2 := NamedQ('E2');
+  
  Context.Default.SyncInvoke(
-    ( A1 + B1          + C1 + WaitFor(C2) + D1 ) *
-    ( A2 + WaitFor(B1) + B2 + C2          + D2 )
+    ( A1 + B1          + C1 + WaitFor(D2) + E1 ) *
+    ( A2 + WaitFor(B1) + C2 + D2          + E2 )
  );
  
 end.
--- a/OpenCL/OpenCLABC/Прекомпиляция
+++ b/OpenCL/OpenCLABC/Прекомпиляция
@ -0,0 +1,11 @@
+
+
+
+__kernel void TEST(__global int* message)
+{
+	int gid = get_global_id(0); // номер текущего вызова TEST
+	
+	message[gid] += gid;
+}
+
+
--- a/OpenCL/OpenCLABC/Прекомпиляция
+++ b/OpenCL/OpenCLABC/Прекомпиляция
@ -0,0 +1,24 @@
+uses OpenCLABC;
+
+begin
+  
+  // Запустите "Прекомпилировать.pas", чтобы создать этот файл
+  {$resource 0.cl.temp_bin}
+  
+  var prog := ProgramCode.DeserializeFrom(Context.Default,
+    System.Reflection.Assembly.GetExecutingAssembly.GetManifestResourceStream('0.cl.temp_bin')
+  );
+  
+  // Дальше всё так же как в "0Простейшие примеры\SimpleAddition"
+  
+  var A := new Buffer( 10 * sizeof(integer) );
+  
+  prog['TEST'].Exec1(10,
+    
+    A.NewQueue.AddFillValue(1)
+    
+  );
+  
+  A.GetArray1&<integer>.Println;
+  
+end.
--- a/ProgramCode/Прекомпилировать.pas
+++ b/ProgramCode/Прекомпилировать.pas
@ -0,0 +1,8 @@
+uses OpenCLABC;
+
+begin
+  var prog := new ProgramCode(Context.Default, ReadAllText('0.cl'));
+  var f := System.IO.File.Create('0.cl.temp_bin');
+  prog.SerializeTo(f);
+  f.Close;
+end.
--- a/InstallerSamples/OpenGL
+++ b/InstallerSamples/OpenGL
@ -71,13 +71,27 @@ body {
 }
 .ps-page-name {
 	cursor: pointer;
-    padding: 1;
+	padding: 1;
 }

 /* ============================== *\
 	       page display
 \* ============================== */

+blockquote > *:first-child {
+	margin-top: 0;
+}
+blockquote > *:last-child {
+	margin-bottom: 0;
+}
+blockquote {
+	margin-left: 0;
+	margin-right: 0;
+	padding: 12px;
+	border-left: 5px solid #d0d0d0;
+	background-color: #f7f7f7;
+}
+
 .smart-link {
 	color: blue;
 	cursor: pointer;
@ -99,7 +113,6 @@ body {
 .code-block {
 	display: inline-block;
 	cursor: text;
-	background: #F0F0F0;
 	border-radius: 3px;
 	
 	padding: 4px;
@ -121,9 +134,13 @@ body {
 	box-shadow: -2px 2px 3px #0000004a;
 	background: linear-gradient(0deg, #f0f0f0, white);
 }
-.inline-code {
+.inline-code, .code-block {
 	background: #F0F0F0;
 }
+blockquote .inline-code,
+blockquote .code-block {
+	background: #DEDEDE;
+}

 .code-keyword {
 	font-weight: bold;
@ -903,7 +920,8 @@ window.onload = ()=>{
 <p>Если в кратце:</p>
 <p>Есть платформа .Net, объединяющая много языков программирования. В том числе C# и PascalABC.NET.</p>
 <p>.exe и .dll созданные такими языками содержат код в особом виде, позволяющем легко подключать
-.dll (и теоретически .exe, но это плохо) созданные на 1 .Net языке к программе на любом другом .Net языке.</p>
+.dll (и теоретически .exe, но это плохо) созданные на 1 .Net языке к программе на любом другом .Net языке.<br />
+(в паскале это делается с помощью <code>$reference</code>)</p>
 <p>Такие .exe и .dll называются управляемыми. .exe и .dll созданные на любом другом (то есть не .Net) языке называются неуправляемыми.</p>
 <p>OpenCL.dll и OpenGL.dll созданы на- и для языков C/C++, поэтому являются неуправляемыми.</p>
 </div>
@ -911,7 +929,7 @@ window.onload = ()=>{
 <div id="page-3" page_name="Участие в разработке н.у. модулей" hidden=true>
 <p>Большинство функций из н.у. модулей требует какой-либо чистки. В основном потому что
 при переводе с C++ тип <code>int*</code> может оказаться и указателем, и <code>var</code>-параметром, и массивом.</p>
-<p>При автоматическом переводе кода с C++ на паскаль - создаются все возможные перегркзки.
+<p>При автоматическом переводе кода с C++ на паскаль - создаются все возможные перегрузки.
 А те перегрузки, которые не подходят конкретным подпрограммам - надо убирать вручную.
 В этом и состоит чистка подпрограмм.</p>
 <p>Все разом функции не почистишь, в одном только <code>OpenGL</code> их &gt;3000. Но сделать это, всё же, надо.
@ -920,9 +938,9 @@ window.onload = ()=>{
 <hr />
 <p>Чтоб чистить подпрограммы было проще - я написал несколько инструментов.
 Вся упаковка н.у. модулей находится <a href="https://github.com/SunSerega/POCGL/tree/master/Packing/Template">тут</a>.
-Откройте папку <code>CL</code> или <code>GL</code>. Они устроены одинаково, поэтому объяснять можно на любой из них.</p>
-<p>Все файлы (кроме .exe и .pdb, которые создаёт паскаль) в этих папках
-открываются любым текстовым редактором (но я рекомендую Notepad++).</p>
+Откройте папку <code>OpenCL</code> или <code>OpenGL</code>. Они устроены одинаково, поэтому объяснять можно на любой из них.</p>
+<p>Все файлы (кроме .exe и .pdb, которые создаёт паскаль) в
+этих папках открываются любым текстовым редактором.</p>
 <hr />
 <h3>Fixers</h3>
 <p>Контроль содержимого модуля находится в папке <code>Fixers</code>. В этой папке есть следующие подпапки:</p>
@ -985,6 +1003,7 @@ window.onload = ()=>{
    public constructor(val: UInt32) := self.val := val;
    public static property Zero: BufferName read default(BufferName);
    public static property Size: integer read Marshal.SizeOf&amp;&lt;UInt32&gt;;
+    public function ToString: string; override := $'BufferName[{val}]';
  end;
 </code></pre>
 <p>Такой подход не замедляет готовую программу, но позволяет отловить некоторые ошибки на этапе компиляции.</p>
@ -1020,7 +1039,7 @@ MyProc(default(BufferName));
 </div>
 <script>on_page_added(document.getElementById("page-5"))</script>
 <div id="page-6" page_name="Записи-перечисления" hidden=true>
-<p>Многие параметры подпрограмм в исходных библиотеках принимают перечисления (enum's).</p>
+<p>Многие параметры подпрограмм в исходных библиотеках принимают перечисления (enum'ы).</p>
 <p>Перечисления, как и имена, это числа. Но в отличии от имён - перечисления принимают заданные заранее константные значения.</p>
 <p>В качестве примера:</p>
 <pre><code>procedure gl.BeginQuery(target: QueryTarget; id: QueryName);
@ -1068,8 +1087,8 @@ gl.BeginQuery(QueryTarget.VERTICES_SUBMITTED, name);
 <pre><code>procedure gl.Clear(mask: ClearBufferMask);
 </code></pre>
 <pre><code>gl.Clear(
-	ClearBufferMask.COLOR_BUFFER_BIT or // очистка поверхности рисования одним цветом
-	ClearBufferMask.DEPTH_BUFFER_BIT    // очистка буфера глубины - обязательна при рисовании 3D
+  ClearBufferMask.COLOR_BUFFER_BIT or // очистка поверхности рисования одним цветом
+  ClearBufferMask.DEPTH_BUFFER_BIT    // очистка буфера глубины - обязательна при рисовании 3D
 );
 </code></pre>
 <p><code>or</code>, в данном случае, побитовое, а не логическое.</p>
@ -1155,7 +1174,7 @@ glDebugOutputAMD.DebugMessageEnableAMD(...);
 <div class="spoiler" summary="Как это читать?" hidden="true"><p>Каждый тип вектора берёт по 1 из значений, перечисленных в <code>[]</code>, через запятую.</p>
 <p>К примеру, есть типы <code>Vec2d</code> и <code>Vec4ui64</code>.</p>
 <p>Число в первых скобках - значит кол-во измерений вектора.</p>
-<p>Буква (буквы) в следующих скобках - значит тип координат вектора:</p>
+<p>Буква (буквы) в следующих скобках - значат тип координат вектора:</p>
 <ul>
 <li><p><code>b=shortint</code>, <code>s=smallint</code>, <code>i=integer</code>, <code>i64=int64</code>: Все 4 типа - целые числа, имеющие бит знака (±) и занимающие 1, 2, 4 и 8 байт соответственно;</p>
 </li>
@ -1171,9 +1190,9 @@ glDebugOutputAMD.DebugMessageEnableAMD(...);
 и возвращает или задаёт значение вектора для соответствующего измерения.</p>
 <p>К примеру:</p>
 <pre><code>var v: Vec4d;
-v[0] := 123.456; // записываем 123.456 по индексу 0
-v[1].Println; // читаем и выводим значение по индексу 1
-v.val[2] := 1; // можно так же писать и имя свойства
+v[0] := 123.456; // Записываем 123.456 по индексу 0
+v[1].Println; // Читаем и выводим значение по индексу 1
+v.val[2] := 1; // Можно так же писать и имя свойства
 </code></pre>
 <p>Но использование этого свойства не рекомендуется. Прямое обращение
 к полю всегда будет быстрее. То есть аналогично предыдущему коду:</p>
@ -1189,7 +1208,7 @@ v.val2 := 1;
 // v1 будет иметь ту же длину, но
 // противоположное v0 направление
 var v1 := -v0;
-// а унарный + не делает ничего, он только
+// А унарный + не делает ничего, он только
 // для красоты. То есть v2 будет =v0
 var v2 := +v0;
 </code></pre>
@ -1197,29 +1216,29 @@ var v2 := +v0;
 <pre><code>var v1: Vec3d;
 var v2: Vec3i;
 ...
-// выведет вектор, имеющий то же
+// Выведет вектор, имеющий то же
 // направление что v1, но в 2 раза длиннее
 (v1*2).Println;

-// выведет вектор, имеющий то же
+// Выведет вектор, имеющий то же
 // направление что v1, но в 2 раза короче
 (v1/2).Println;

-// к целочисленным векторам вместо
+// К целочисленным векторам вместо
 // обычного деления надо применять div
 (v2 div 2).Println;
 </code></pre>
 <h3>Операции с 2 векторами</h3>
 <pre><code>var v1, v2: Vec3d;
 ...
-// скалярное произведение векторов
+// Скалярное произведение векторов
 (v1*v2).Println;

-// сумма векторов, складывает
+// Сумма векторов, складывает
 // отдельно каждый элемент вектора
 (v1+v2).Println;

-// разность векторов, тоже работает
+// Разность векторов, тоже работает
 // отдельно на каждый элемент вектора
 (v1-v2).Println;
 </code></pre>
@ -1236,8 +1255,8 @@ var v2: Vec2i;
 Каким образом находить корень полученного значения - дело программиста.</p>
 <pre><code>var v1: Vec3d;
 ...
-v1.SqrLength.Println; // квадрат длины
-v1.SqrLength.Sqrt.Println; // сама длина
+v1.SqrLength.Println; // Квадрат длины
+v1.SqrLength.Sqrt.Println; // Сама длина
 </code></pre>
 <h3>Normalized</h3>
 <p>Метод <code>.Normalized</code> возвращает нормализированную (с длиной =1) версию вектора.<br />
@ -1248,7 +1267,7 @@ v1.Println;
 v1.SqrLength.Sqrt.Println;
 var v2 := v1.Normalized;
 v2.Println;
-v2.SqrLength.Sqrt.Println; // обязательно будет 1
+v2.SqrLength.Sqrt.Println; // Обязательно будет 1
 </code></pre>
 <h3>Cross</h3>
 <p>Статичный метод <code>.Cross</code> возвращает произведение векторов (Cross product, не путать со скалярным произведением).
@ -1273,7 +1292,7 @@ v2.SqrLength.Sqrt.Println; // обязательно будет 1
 <h3>Превращение в строку</h3>
 <pre><code>var v1: Vec4d;
 ...
-v1.Println; // вывод вектора
+v1.Println; // Вывод вектора
 // s присвоит ту же строку, что выводит .Println
 var s := v1.ToString;
 </code></pre>
@ -1311,7 +1330,7 @@ m.val31 := 1;
 <p>Но у матриц так же есть свойства для столбцов и строк:</p>
 <pre><code>var m: Mtr3d;
 ...
-m.Row0.Println; // вывод нулевой строчки в виде вектора
+m.Row0.Println; // Вывод нулевой строчки в виде вектора
 m.Row1 := new Vec3d(1,2,3);
 m.Col2.Println;
 </code></pre>
@ -1319,7 +1338,7 @@ m.Col2.Println;
 получать по динамическому индексу (но, опять же, это медленнее):</p>
 <pre><code>var m: Mtr3d;
 ...
-m.Row[0].Println; // вывод нулевой строчки в виде вектора
+m.Row[0].Println;
 m.Row[1] := new Vec3d(1,2,3);
 m.Col[2].Println;
 </code></pre>
@ -1340,7 +1359,7 @@ var ptr2 := m.ColPtr[3];
 <p><code>Identity</code> создаёт новую единичную матрицу. То есть матрицу, у
 которой главная диагональ заполнена 1, а всё останое заполнено 0.</p>
 <pre><code>Mtr3d.Identity.Println;
-// работает и для не_квадратных матриц
+// Работает и для не_квадратных матриц
 Mtr2x3d.Identity.Println;
 </code></pre>
 <h3>UseColPtr*</h3>
@ -1378,8 +1397,8 @@ var v := new Vec3d(1,2,3);
 умножении на которую к вектору добавляется заданное значение.</p>
 <pre><code>var m := Mtr4d.Translate(1,2,3);

-// последний элемент должен быть 1,
-// чтоб матрица из .Translate правильно работала
+// Последний элемент должен быть 1,
+// чтобы матрица из .Translate правильно работала
 var v := new Vec4d(0,0,0,1);

 (m*v).Println; // (1;2;3)
@ -1407,7 +1426,7 @@ var v := new Vec4d(0,0,0,1);
  1,2,3,
  4,5,6
 );
-m.Transpose.Println; // выводит:
+m.Transpose.Println; // Выводит:
 // 1 4
 // 2 5
 // 3 6
@ -1423,7 +1442,7 @@ m.Transpose.Println; // выводит:
 <p>Как и у векторов - матрицы можно выводить и превращать в строку</p>
 <pre><code>var m: Mtr4d;
 ...
-m.Println; // вывод матрицы
+m.Println; // Вывод матрицы
 // s присвоит ту же строку, что выводит .Println
 var s := m.ToString;
 </code></pre>
@ -1828,7 +1847,7 @@ begin
  cl.SetEventCallback(ev, CommandExecutionStatus.COMPLETE, cb,IntPtr.Zero).RaiseIfError;
 end.
 </code></pre>
-<p>Этот код будет может время от времени вылетать, потому что:</p>
+<p>Этот код может время от времени вылетать, потому что:</p>
 <ol>
 <li><p><code>cl.SetEventCallback</code> вызывает свой коллбек тогда, когда посчитает нужным (но обычно после того как вызов <code>cl.SetEventCallback</code> завершился);</p>
 </li>
--- a/InstallerSamples/OpenGL
+++ b/InstallerSamples/OpenGL
--- a/InstallerSamples/OpenGL
+++ b/InstallerSamples/OpenGL
@ -71,13 +71,27 @@ body {
 }
 .ps-page-name {
 	cursor: pointer;
-    padding: 1;
+	padding: 1;
 }

 /* ============================== *\
 	       page display
 \* ============================== */

+blockquote > *:first-child {
+	margin-top: 0;
+}
+blockquote > *:last-child {
+	margin-bottom: 0;
+}
+blockquote {
+	margin-left: 0;
+	margin-right: 0;
+	padding: 12px;
+	border-left: 5px solid #d0d0d0;
+	background-color: #f7f7f7;
+}
+
 .smart-link {
 	color: blue;
 	cursor: pointer;
@ -99,7 +113,6 @@ body {
 .code-block {
 	display: inline-block;
 	cursor: text;
-	background: #F0F0F0;
 	border-radius: 3px;
 	
 	padding: 4px;
@ -121,9 +134,13 @@ body {
 	box-shadow: -2px 2px 3px #0000004a;
 	background: linear-gradient(0deg, #f0f0f0, white);
 }
-.inline-code {
+.inline-code, .code-block {
 	background: #F0F0F0;
 }
+blockquote .inline-code,
+blockquote .code-block {
+	background: #DEDEDE;
+}

 .code-keyword {
 	font-weight: bold;
@ -859,9 +876,9 @@ window.onload = ()=>{
 }
 </script>
 <div id="page-1" page_name="" hidden=true>
-<p><code>ToDo</code> модуля <code>OpenGLABC</code> пока что не существует. Когда то он появится, но пока что есть только это сообщение.</p>
+<p><code>ToDo</code> модуль <code>OpenGLABC</code> пока что не существует. Когда то он появится, но пока что есть только это сообщение.</p>
 <p>Если хотите понять, было ли это сообщение оставлено тут 10 лет назад и забыто - вот дата, когда автор сообщения последний раз прочитал его:<br />
-24.02.2020</p>
+03.08.2020</p>
 </div>
 <script>on_start_folder("GL ABC", document.getElementById("page-1"))</script>
 <script>on_end_folder()</script>
--- a/ReleaseGenerators/RebuildStandartModules.pas
+++ b/ReleaseGenerators/RebuildStandartModules.pas
@ -13,8 +13,8 @@
   RobotTaskMaker, RobotZadan, Sockets, Utils, VCL, Timers, PT4Exam, PT4TaskMakerNET, RBDMUtils,
   Collections, Arrays, Core, FormsABC, MPI, ClientServer, Speech, Sounds, Countries,
   ABCDatabases,BlockFileOfT,Controls,
-   OpenCL, OpenCLABC,
-   OpenGL, OpenGLABC,
+   OpenCL, OpenCLABCBase, OpenCLABC,
+   OpenGL, OpenGLABCBase, OpenGLABC,
   BBCMicroBit, School, SF, TwoPanelsWindow
   ;
   
--- a/ReleaseGenerators/sect_Core.nsh
+++ b/ReleaseGenerators/sect_Core.nsh
@ -129,8 +129,10 @@
    File ..\bin\Lib\Sounds.pcu
    File ..\bin\Lib\BlockFileOfT.pcu
    File ..\bin\Lib\OpenCL.pcu
+    File ..\bin\Lib\OpenCLABCBase.pcu
    File ..\bin\Lib\OpenCLABC.pcu
    File ..\bin\Lib\OpenGL.pcu
+    File ..\bin\Lib\OpenGLABCBase.pcu
    File ..\bin\Lib\OpenGLABC.pcu
    File ..\bin\Lib\School.pcu
    File ..\bin\Lib\SF.pcu
@ -199,8 +201,10 @@
    ${AddFile} "Sounds.pcu"
    ${AddFile} "BlockFileOfT.pcu"
    ${AddFile} "OpenCL.pcu"
+    ${AddFile} "OpenCLABCBase.pas"
    ${AddFile} "OpenCLABC.pcu"
    ${AddFile} "OpenGL.pcu"
+    ${AddFile} "OpenGLABCBase.pas"
    ${AddFile} "OpenGLABC.pcu"
    ${AddFile} "School.pcu"
    ${AddFile} "SF.pcu"
@ -281,8 +285,10 @@
    File ..\bin\Lib\Sounds.pas
    File ..\bin\Lib\BlockFileOfT.pas
    File ..\bin\Lib\OpenCL.pas
+    File ..\bin\Lib\OpenCLABCBase.pas
    File ..\bin\Lib\OpenCLABC.pas
    File ..\bin\Lib\OpenGL.pas
+    File ..\bin\Lib\OpenGLABCBase.pas
    File ..\bin\Lib\OpenGLABC.pas
    File ..\bin\Lib\School.pas
    File ..\bin\Lib\SF.pas
@ -342,14 +348,15 @@
    ${AddFile} "Core.pas"
    ${AddFile} "MPI.pas"
    ${AddFile} "ClientServer.pas"
-    ${AddFile} "OpenGL.pas"
    ${AddFile} "PT4MakerNetX.pas"
    ${AddFile} "Speech.pas"
    ${AddFile} "Sounds.pas"
    ${AddFile} "BlockFileOfT.pas"
    ${AddFile} "OpenCL.pas"
+    ${AddFile} "OpenCLABCBase.pas"
    ${AddFile} "OpenCLABC.pas"
    ${AddFile} "OpenGL.pas"
+    ${AddFile} "OpenGLABCBase.pas"
    ${AddFile} "OpenGLABC.pas"
    ${AddFile} "School.pas"
    ${AddFile} "SF.pas"
--- a/bin/Lib/OpenCL.pas
+++ b/bin/Lib/OpenCL.pas
--- a/bin/Lib/OpenCLABC.pas
+++ b/bin/Lib/OpenCLABC.pas
--- a/bin/Lib/OpenCLABCBase.pas
+++ b/bin/Lib/OpenCLABCBase.pas
--- a/bin/Lib/OpenGL.pas
+++ b/bin/Lib/OpenGL.pas
--- a/bin/Lib/OpenGLABC.pas
+++ b/bin/Lib/OpenGLABC.pas
@ -1,5 +1,4 @@
-
-///
+///
 ///Модуль, зарезервированный для высокоуровневой оболочки модуля OpenGL
 ///
 unit OpenGLABC;
@ -7,6 +6,7 @@ unit OpenGLABC;
 interface

 uses OpenGL;
+uses OpenGLABCBase;

 implementation

--- a/bin/Lib/OpenGLABCBase.pas
+++ b/bin/Lib/OpenGLABCBase.pas
@ -0,0 +1,4 @@
+/// Модуль для внутренних типов модуля OpenGLABC
+unit OpenGLABCBase;
+
+end.