diff --git a/LastPack.log b/LastPack.log index 62d4e19f..ebb12d34 100644 --- a/LastPack.log +++ b/LastPack.log @@ -621,6 +621,12 @@ Tester: Compiling "Samples/OpenCLABC/Кеширование CLProgramCode/Simple Tester: Compiling: OK Tester: Compiling "Samples/OpenGLABC/!Крутящийся треугольник/Крутящийся треугольник.pas" Tester: Compiling: OK +Tester: Compiling "Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas" +Tester: Compiling: OK +Tester: Compiling "Samples/OpenGLABC/Mandelbrot/FieldTest.pas" +Tester: Compiling: OK +Tester: Compiling "Samples/OpenGLABC/Mandelbrot/SamplingTests.pas" +Tester: Compiling: OK Tester: Compiling "Samples/OpenGLABC/Точки на поле/Точки.pas" Tester: Compiling: OK Tester: Switched to platform "NVIDIA CUDA" and using 1 devices @@ -832,6 +838,10 @@ Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/Si Tester: Done executing Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/SimpleAddition но с кешем] Tester: Done executing +Tester: Executing Test[Samples/OpenGLABC/Mandelbrot/FieldTest] +Tester: Done executing +Tester: Executing Test[Samples/OpenGLABC/Mandelbrot/SamplingTests] +Tester: Done executing Tester: Switched to platform "Intel(R) OpenCL" and using 1 devices Tester: Executing Test[Tests/Exec/CL/ToString] Tester: Done executing @@ -1041,6 +1051,10 @@ Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/Si Tester: Done executing Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/SimpleAddition но с кешем] Tester: Done executing +Tester: Executing Test[Samples/OpenGLABC/Mandelbrot/FieldTest] +Tester: Done executing +Tester: Executing Test[Samples/OpenGLABC/Mandelbrot/SamplingTests] +Tester: Done executing Tester: Cleanup Tester: Done testing Finished runing Tester @@ -1070,6 +1084,22 @@ Packing sample file "Samples/OpenCLABC/Кеширование CLProgramCode/Simp Packing sample file "Samples/OpenGLABC/Common.pas" Packing sample file "Samples/OpenGLABC/!Крутящийся треугольник/Крутящийся треугольник.pas" Packing sample file "Samples/OpenGLABC/!Крутящийся треугольник/Крутящийся треугольник.vert" +Packing sample file "Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Blocks.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/CameraDef.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/FieldTest.bmp" +Packing sample file "Samples/OpenGLABC/Mandelbrot/FieldTest.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/GL_CL_Context.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.cl" +Packing sample file "Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/MemoryLayering.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/PointComponents.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/SamplingTests.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Settings.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/SheetTransfer.cl" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Shaders/Box.geom" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Shaders/Empty.vert" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Shaders/Rainbow.frag" Packing sample file "Samples/OpenGLABC/Точки на поле/Empty.vert" Packing sample file "Samples/OpenGLABC/Точки на поле/Mandelbrot.frag" Packing sample file "Samples/OpenGLABC/Точки на поле/SinglePointToScreen.geom" @@ -1079,5 +1109,5 @@ Packing sample file "Samples/OpenGLABC/Точки на поле/Минимум Packing sample file "Samples/OpenGLABC/Точки на поле/Спирали.frag" Packing sample file "Samples/OpenGLABC/Точки на поле/Сумма расстояний.frag" Packing sample file "Samples/OpenGLABC/Точки на поле/Точки.pas" -Packed 27 sample files +Packed 43 sample files Done packing diff --git a/Log/Release.log b/Log/Release.log index f1b714f3..dd18bd5f 100644 --- a/Log/Release.log +++ b/Log/Release.log @@ -23,6 +23,22 @@ Packing sample file "Samples/OpenCLABC/Кеширование CLProgramCode/Simp Packing sample file "Samples/OpenGLABC/Common.pas" Packing sample file "Samples/OpenGLABC/!Крутящийся треугольник/Крутящийся треугольник.pas" Packing sample file "Samples/OpenGLABC/!Крутящийся треугольник/Крутящийся треугольник.vert" +Packing sample file "Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Blocks.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/CameraDef.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/FieldTest.bmp" +Packing sample file "Samples/OpenGLABC/Mandelbrot/FieldTest.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/GL_CL_Context.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.cl" +Packing sample file "Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/MemoryLayering.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/PointComponents.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/SamplingTests.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Settings.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/SheetTransfer.cl" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Shaders/Box.geom" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Shaders/Empty.vert" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Shaders/Rainbow.frag" Packing sample file "Samples/OpenGLABC/Точки на поле/Empty.vert" Packing sample file "Samples/OpenGLABC/Точки на поле/Mandelbrot.frag" Packing sample file "Samples/OpenGLABC/Точки на поле/SinglePointToScreen.geom" @@ -32,4 +48,4 @@ Packing sample file "Samples/OpenGLABC/Точки на поле/Минимум Packing sample file "Samples/OpenGLABC/Точки на поле/Спирали.frag" Packing sample file "Samples/OpenGLABC/Точки на поле/Сумма расстояний.frag" Packing sample file "Samples/OpenGLABC/Точки на поле/Точки.pas" -Packed 27 sample files +Packed 43 sample files diff --git a/Log/Test.log b/Log/Test.log index 7e6a268d..52a0afdd 100644 --- a/Log/Test.log +++ b/Log/Test.log @@ -229,6 +229,12 @@ Tester: Compiling "Samples/OpenCLABC/Кеширование CLProgramCode/Simple Tester: Compiling: OK Tester: Compiling "Samples/OpenGLABC/!Крутящийся треугольник/Крутящийся треугольник.pas" Tester: Compiling: OK +Tester: Compiling "Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas" +Tester: Compiling: OK +Tester: Compiling "Samples/OpenGLABC/Mandelbrot/FieldTest.pas" +Tester: Compiling: OK +Tester: Compiling "Samples/OpenGLABC/Mandelbrot/SamplingTests.pas" +Tester: Compiling: OK Tester: Compiling "Samples/OpenGLABC/Точки на поле/Точки.pas" Tester: Compiling: OK Tester: Executing Test[Tests/Exec/CL/ToString] @@ -439,6 +445,10 @@ Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/Si Tester: Done executing Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/SimpleAddition но с кешем] Tester: Done executing +Tester: Executing Test[Samples/OpenGLABC/Mandelbrot/FieldTest] +Tester: Done executing +Tester: Executing Test[Samples/OpenGLABC/Mandelbrot/SamplingTests] +Tester: Done executing Tester: Cleanup Tester: Done testing Finished runing Tester diff --git a/Samples/OpenGLABC/Mandelbrot/.gitignore b/Samples/OpenGLABC/Mandelbrot/.gitignore new file mode 100644 index 00000000..ee813f11 --- /dev/null +++ b/Samples/OpenGLABC/Mandelbrot/.gitignore @@ -0,0 +1,6 @@ + + + +Cache/block_w_pow=*/*.point_block + + diff --git a/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas b/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas index 01319ff4..0f7b8637 100644 --- a/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas +++ b/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas @@ -4,22 +4,49 @@ {$apptype windows} // Управление: +// - Escape: Завершение программы (дважды или +Ctrl чтобы отменить сохранение) // - Space: Сбросить положение камеры -// - Ctrl+C: Скопировать положение камеры -// - Ctrl+V: Вставить положение камеры // - Mouse Drag: Быстрое движение камеры // - Arrows: Гладкое движение камеры // - Scroll: Быстрое изменение масштаба // - "+" и "-": Гладкое изменение масштаба +//TODO: +// - Ctrl+C: Скопировать положение камеры (+Shift чтобы добавить комментарий) +// - Ctrl+V: Вставить положение камеры +// - Alt: Вид без копирования информации предыдущих кадров // - Alt+Enter: Полноэкранный режим +// - B: Телепортировать камеру к курсору (Blink) +// --- Пока держат - выводить точку в начале заголовка, а телепортировать когда отпускают -//TODO Доделать ограничение view_bound -// - Сейчас его не возвращает из функции, создающей блоки - -// Константы, которые можно быстро менять +// В модуле Settings находятся все основных константы +// + объяснение логики программы, чтобы понимать зачем эти константы +// Ctrl+тыкните на название модуля в uses чтобы открыть его uses Settings; -uses GL_CL_Context; +//TODO Вывод шагов под курсором чтобы норм дебажить +// - in[123] +// - out[456] +//TODO Выводить отдельно для sheet и для блоков +// - Для этого надо находить номер блока и точки в нём и кидать (x;y) точки в CQ_GetData + +//TODO Alt-режим почему то не работает... +// - Точнее, если масштабировать с ним - начинаются глюки +// - Вообще, он уже и не полезен. Зато полезна была бы возможность сбрасывать sheet вручную + +//TODO Отдельная программа для полной прорисовки кардров с движением камеры от 1 точки (и масштаба) к другой + +//TODO При очень большом приближении край рисунка ведёт себя криво +// - Потому что FirstWordToReal +// - Надо в виде PointComponent считать разницу сначала + +//TODO mouse_grab_move ведёт себя не стабильно (точка которую держат может потихоньку сдвигаться) +// - Надо запоминать camera.pos в начале движения мышкой +// - И затем пересчитывать на каждом кадре относительно него + +//TODO Отдельное окно по нажатию какой-то клавиши с кучей инфы +// - Загруженность памяти (VRAM,RAM,Drive) +// - Скорость обработки блоков (ну и текущее кол-во слов там же) +// - Для начала выводить сколько памяти тратится на sheet-ы uses System; uses System.Windows.Forms; @@ -27,53 +54,16 @@ uses OpenGL; uses OpenGLABC; -//uses OpenCL; //TODO Remove uses OpenCLABC; -uses Common in '../Common'; +uses GL_CL_Context; + +uses Common;// in '../Common'; //TODO Merge uses PointComponents; uses CameraDef; uses Blocks; -//var extra_debug_output := false; - -//TODO Система отдельных блоков, как на гугл картах -// - Запоминать номер кадра последнего использования каждого блока -//TODO Что насчёт рисования кадра? -// - В идеале надо смотреть только на блок нужного масштаба -// - Но если он не полностью нарисован: -// --- Сначала пытаться нарисовать более мелкие блоки -// --- Затем наоборот, более большие (а знач неточные) -// - А что с точками, которые ещё считает? -// --- По сути можно иметь 3 глубины: -// --- 0=не нарисовано -// --- 1=нарисовано недосчитанным -// --- 2=нарисовано конечным -// - Конечный статус наверн будет у очень маленького кол-ва точек... -// - То есть, наверное, нет смысла давать блоку конечный статус - -//TODO + и - для стабильного зума -//TODO * и / для максимального кол-ва шагов (для цвета и расчётов) - -//TODO Множество разных фрагметных шейдеров, для разной раскраски в зависимости от глубины -// - По глубине: -// --- от depth/макс_глубину -// --- от depth/n%1 -// - По цвету (0 и MaxLongWord считать особенными) -// --- hsv2rgb -// --- чёрно-белое -//TODO Максимальную глубину передавать в уже сглаженном виде, чтобы цвета не менялись резко -//TODO Или может собирать фрагментные шейдеры на ходу, из комбинаций кусков кода? - -//TODO Значения z надо хранить с изменяемой точность, исходя из текущего масштаба - -//TODO Отдельное окно по нажатию какой-то клавиши с кучей инфы -// - Загруженность памяти (VRAM,RAM) -// - Скорость обработки блоков (ну и текущее кол-во слов там же) - -//TODO Попробовать вставить цикл в корень MandelbrotBlockStep, чтобы не запускать kernel кучу раз - type UpdateTimingQueue = sealed class private sw := Stopwatch.StartNew; @@ -115,7 +105,7 @@ BoundUniforms = record (xl, yl) := (gl.GetUniformLocation(shader_prog, $'{prefix}_skip_x_last'), gl.GetUniformLocation(shader_prog, $'{prefix}_skip_y_last')); end; - procedure Write(b: BoundDefs); + procedure Write(b: BoundDefs); begin gl.Uniform1f(xf, b.xf); gl.Uniform1f(yf, b.yf); @@ -125,24 +115,101 @@ BoundUniforms = record end; + CLGLArray = sealed class + where T: record; + + public b_gl: gl_buffer; + public b_cl: CLArray; + + public constructor(gl: OpenGL.gl); + begin + gl.CreateBuffers(1, self.b_gl); + b_cl := nil; + end; + + public function EnsureLen(len: integer): boolean; + begin + Result := false; + if (b_cl<>nil) and (b_cl.Length>=len) then exit; + + if b_cl<>nil then + CLMemoryObserver.Current.RemoveMemoryUse(b_cl.ByteSize, b_gl); + gl.NamedBufferData(b_gl, new UIntPtr(len*System.Runtime.InteropServices.Marshal.SizeOf&), IntPtr.Zero, glVertexBufferObjectUsage.STREAM_DRAW); + GL_CL_Context.WrapBuffer(b_gl, b_cl); + if b_cl.Length <> len then + raise new InvalidOperationException; + CLMemoryObserver.Current.AddMemoryUse(b_cl.ByteSize, b_gl); + + Result := true; + end; + + end; + begin var f := new Form; - f.WindowState := FormWindowState.Maximized; + CLMemoryObserver.Current := new TrackingMemoryObserver; // f.ControlBox := false; +// OpenCLABC.eh_debug_otp := nil; + + // Может понадобится увеличить max_VRAM в "Settings.pas", + // для раскрытия на полный экран с 1920х1080 (1080p) и больше пикселей +// f.WindowState := FormWindowState.Maximized; +// f.ClientSize := new System.Drawing.Size(1920,1080); // 1080p, 16:9 +// f.ClientSize := new System.Drawing.Size(1280,720); // 720p, 16:9 + f.ClientSize := new System.Drawing.Size(1440,720); // 720p, 2:1 +// f.ClientSize := new System.Drawing.Size(1072,603); // 603p, 16:9 +// f.ClientSize := new System.Drawing.Size(1206,603); // 603p, 2:1 + f.StartPosition := FormStartPosition.CenterScreen; + + {$region Закрытие} - // Моментальное закрытие при Alt+F4 и Esc - f.Closed += (o,e)->Halt(); f.KeyUp += (o,e)-> case e.KeyCode of - Keys.Escape: Halt; + Keys.Escape: f.Close; + end; + f.Closing += (o,e)-> + begin + if Control.ModifierKeys.HasFlag(Keys.Control) then Halt; + + var shutdown_progress_form := new Form; + shutdown_progress_form.StartPosition := FormStartPosition.CenterScreen; + shutdown_progress_form.FormBorderStyle := FormBorderStyle.None; + shutdown_progress_form.Closing += (o,e)->Halt(); + shutdown_progress_form.KeyUp += (o,e)-> + case e.KeyCode of + Keys.Escape: shutdown_progress_form.Close; + end; + + var progress_bar := new ProgressBar; + shutdown_progress_form.Controls.Add(progress_bar); + shutdown_progress_form.ClientSize := new System.Drawing.Size( + (Screen.PrimaryScreen.WorkingArea.Width * 0.7).Round, + progress_bar.Height + ); + progress_bar.Dock := DockStyle.Fill; + + var progress_t := new Timer; + progress_t.Interval := 10; + progress_t.Tick += (o,e)-> + begin + var (done,total) := BlockUpdater.ShutdownProgress; + progress_bar.Minimum := 0; + progress_bar.Maximum := total; + progress_bar.Value := done; + f.Text := $'Saving to drive: {done/total:P} ({done}/{total})'; + end; + + BlockUpdater.BeginShutdown(shutdown_progress_form.Close); + progress_t.Start; + + shutdown_progress_form.ShowDialog; end; - var camera: CameraPos; + {$endregion Закрытие} var need_resize := false; f.Shown += (o,e)-> begin - camera := new CameraPos(f.ClientSize.Width, f.ClientSize.Height); need_resize := true; f.Resize += (o,e)-> begin @@ -153,35 +220,84 @@ BoundUniforms = record end; end; - //TODO Использовать чтобы выдавать кол-во итераций под курсором - var mouse_pos: Vec2i; + var draw_alt_mode := false; + var mouse_pos := default(Vec2i); + var mouse_captured := true; + var mouse_grab_move := default(Vec2i); + var scale_speed_add := 0; + var camera_reset := false; + var slow_move_dir := default(Vec2i); + var slow_scale_dir1 := 0; + var slow_scale_dir2 := 0; {$region Управление} begin -// var CoordsFromScreen := function(X,Y: integer): Vec2d -> -// begin -// var logic_pos := new Vec2d(X/f.ClientSize.Width-0.5, 0.5-Y/f.ClientSize.Height)*2; -// var pos := new Vec2d(logic_pos.val0*camera.aspect, logic_pos.val1); -// Result := pos*camera.scale + camera.pos; -// end; + f.MouseEnter += (o,e)->(mouse_captured := true); + f.MouseLeave += (o,e)->(mouse_captured := false); - f.MouseWheel += (o,e)-> - begin - //TODO -// var pos := CoordsFromScreen(e.X, e.Y); -// -// var pow := 1 - Sign(e.Delta)*0.1; -// camera.scale := camera.scale * pow; -// -// camera.pos := camera.pos + (pos-CoordsFromScreen(e.X, e.Y)); - end; + f.KeyDown += (o,e)->(draw_alt_mode := e.Alt); + f.KeyUp += (o,e)->(draw_alt_mode := e.Alt); f.KeyDown += (o,e)-> case e.KeyCode of - Keys.Space: camera := new CameraPos(f.Width, f.Height); + Keys.Space: + begin + camera_reset := true; + scale_speed_add := 0; + end; end; - f.MouseMove += (o,e)->(mouse_pos := new Vec2i(e.X,e.Y)); + var mouse_grabbed := false; + f.MouseDown += (o,e)-> + case e.Button of + MouseButtons.Left: mouse_grabbed := true; + end; + f.MouseUp += (o,e)-> + case e.Button of + MouseButtons.Left: mouse_grabbed := false; + end; + + f.MouseWheel += (o,e)->System.Threading.Interlocked.Add(scale_speed_add, e.Delta); + f.MouseMove += (o,e)-> + begin + var n_mouse_pos := new Vec2i(e.X,e.Y); + if mouse_grabbed then + begin + var change := mouse_pos-n_mouse_pos; + System.Threading.Interlocked.Add(mouse_grab_move.val0, change.val0); + System.Threading.Interlocked.Add(mouse_grab_move.val1, change.val1); + end; + mouse_pos := n_mouse_pos; + end; + + var define_slow_control := procedure(key_low, key_high, modifiers: Keys; on_change: integer->())-> + begin + var low_pressed := false; + var high_pressed := false; + var update := procedure->on_change(Ord(high_pressed)-Ord(low_pressed)); + f.KeyDown += (o,e)-> + begin + if e.Modifiers and modifiers <> modifiers then exit; + if e.KeyCode=key_low then low_pressed := true else + if e.KeyCode=key_high then high_pressed := true else + exit; + update; + end; + f.KeyUp += (o,e)-> + begin + if e.Modifiers and modifiers <> modifiers then exit; + if e.KeyCode=key_low then low_pressed := false else + if e.KeyCode=key_high then high_pressed := false else + exit; + update; + end; + end; + + define_slow_control(Keys.Subtract, Keys.Add, Keys.None, x->(slow_scale_dir1:=x)); + define_slow_control(Keys.OemMinus, Keys.Oemplus, Keys.Shift, x->(slow_scale_dir2:=x)); + + define_slow_control(Keys.Left, Keys.Right, Keys.None, x->(slow_move_dir.val0:=x)); + define_slow_control(Keys.Down, Keys.Up, Keys.None, x->(slow_move_dir.val1:=x)); end; {$endregion Управление} @@ -192,19 +308,13 @@ BoundUniforms = record GL_CL_Context.Init(hdc); gl := new OpenGL.gl(pl); - begin - var org_swap_interval := wglSwapControlEXT.Instance.GetSwapIntervalEXT; - var new_swap_interval := Settings.frame_interval; - if extra_debug_output then - $'Swap interval: {org_swap_interval}=>{new_swap_interval}'.Println; - if not wglSwapControlEXT.Instance.SwapIntervalEXT(new_swap_interval) then - raise new System.InvalidOperationException; - end; + // Реализация от NVidia тратит 16мс на вызов clGLSharingKHR.EnqueueAcquireGLObjects (при чём синхронно), если не выключить vsync + // Точнее vsync применяется и к EndFrame, и затем ещё раз к .EnqueueAcquireGLObjects + // Ну, в этой программе vsync не сдался... + if not wglSwapControlEXT.Instance.SwapIntervalEXT(0) then + raise new System.InvalidOperationException; - var sw := Stopwatch.StartNew; - var timings_max_count := 30; - // sw.Elapsed для каждого из последних timings_max_count кадров - var timings := new Queue(timings_max_count); + {$region Общие данные для всех кадров} var s_vert_empty := InitShaderResource('Empty.vert', glShaderType.VERTEX_SHADER); {$resource Shaders/Empty.vert} @@ -216,7 +326,6 @@ BoundUniforms = record var uniform_view_bound, uniform_sheet_bound: BoundUniforms; var uniform_sheet_size: integer; var ssb_sheet: integer; - var uniform_max_steps: integer; var choose_frag_shader := procedure(s_frag: gl_shader)-> begin shader_prog := InitProgram(s_vert_empty, s_geom_box, s_frag); @@ -225,39 +334,16 @@ BoundUniforms = record uniform_sheet_bound := new BoundUniforms(shader_prog, 'sheet'); uniform_sheet_size := gl.GetUniformLocation(shader_prog, 'sheet_size'); ssb_sheet := gl.GetProgramResourceIndex(shader_prog, glProgramInterface.SHADER_STORAGE_BLOCK, 'sheet_block'); - uniform_max_steps := gl.GetUniformLocation(shader_prog, 'max_steps'); end; + //TODO Собственно использовать чтобы менять шейдеры choose_frag_shader(s_frag_rainbow); - var cl_err_buffer := new CLArray(3); - var cl_uc_buffer := new CLValue; + var sheet_draw := new CLGLArray(gl); + var sheet_back := new CLGLArray(gl); + var V_ExtractedCount := new CLValue; - var gl_sheet_buffer: gl_buffer; - gl.CreateBuffers(1, gl_sheet_buffer); - var cl_sheet_buffer: CLArray; - var cl_sheet_states: CLArray; - var curr_sheet_size := -1; - var Q_Init := CQNil; - var ensure_sheet_buffer_size := procedure(w, h: integer)-> - begin - var req_size := w*h; - if req_size<=0 then - raise new InvalidOperationException; - if curr_sheet_size>=req_size then exit; - curr_sheet_size := req_size; - gl.NamedBufferData(gl_sheet_buffer, new UIntPtr(req_size*sizeof(cardinal)), IntPtr.Zero, glVertexBufferObjectUsage.STREAM_DRAW); - GL_CL_Context.WrapBuffer(gl_sheet_buffer, cl_sheet_buffer); - if cl_sheet_buffer.Length <> req_size then - raise new InvalidOperationException; - cl_sheet_states := new CLArray(req_size); - Q_Init := CQNil - + cl_sheet_buffer.MakeCCQ.ThenFillValue(0) - + cl_sheet_states.MakeCCQ.ThenFillValue(0) - + cl_err_buffer.MakeCCQ.ThenFillValue(0) - + cl_uc_buffer.MakeCCQ.ThenWriteValue(0) - + CQNil - ; - end; + var last_render_info := default(BlockLayerRenderInfo?); + var last_render_sheet_w := 0; // Для дебага // var buffer_temp: gl_buffer; @@ -265,63 +351,122 @@ BoundUniforms = record // gl.NamedBufferData(buffer_temp, new IntPtr(3*sizeof(real)), IntPtr.Zero, VertexBufferObjectUsage.DYNAMIC_READ); // gl.BindBufferBase(BufferTarget.SHADER_STORAGE_BUFFER, 1, buffer_temp); - //TODO Calculate on-fly - var max_steps := 256; + var t_full := new UpdateTimingQueue(120); + var t_body := new UpdateTimingQueue(120); - while true do - begin + var camera := new CameraPos(f.ClientSize.Width, f.ClientSize.Height); + var scale_speed := 0.0; + + var frame_time_sw := Stopwatch.StartNew; + var last_frame_time := frame_time_sw.Elapsed; + + {$endregion Общие данные для всех кадров} + + while BlockUpdater.ShutdownProgress=nil do + try var curr_frame_resized := false; - if need_resize then begin var w_size := f.ClientSize; + if need_resize then + begin + gl.Viewport(0,0, w_size.Width,w_size.Height); + curr_frame_resized := true; + camera.Resize(w_size.Width, w_size.Height); + end; + if camera_reset then + begin + camera_reset := false; + camera := new CameraPos(w_size.Width, w_size.Height); + scale_speed := 0; + end; + end; + t_body.Start; + + begin + var next_frame_time := frame_time_sw.Elapsed; + var frame_len := (next_frame_time-last_frame_time).TotalSeconds; + last_frame_time := next_frame_time; - gl.Viewport(0,0, w_size.Width,w_size.Height); - curr_frame_resized := true; + scale_speed += System.Threading.Interlocked.Exchange(scale_speed_add,0) * 0.005; - camera.Resize(w_size.Width, w_size.Height); + camera.Move( + slow_move_dir.val0*10 * frame_len + System.Threading.Interlocked.Exchange(mouse_grab_move.val0, 0), + slow_move_dir.val1*10 * frame_len - System.Threading.Interlocked.Exchange(mouse_grab_move.val1, 0), + mouse_pos.val0, mouse_pos.val1, + scale_speed + frame_len * (slow_scale_dir1+slow_scale_dir2), + mouse_captured + ); + + scale_speed *= 0.5; end; - var render_info := BlockLayer.BlocksForCurrentScale(camera); - var b_cy := render_info.blocks.GetLength(0); - var b_cx := render_info.blocks.GetLength(1); - var render_block_size := Settings.block_w shr render_info.mipmap_lvl; - var render_sheet_w := b_cx * render_block_size; - var render_sheet_h := b_cy * render_block_size; - ensure_sheet_buffer_size(render_sheet_w, render_sheet_h); + camera.FixWordCount; + var render_info := BlockLayer.GetLayer(camera).GetRenderInfo(camera, last_render_info); + BlockUpdater.SetCurrent(render_info.block_area); + last_render_info := render_info; + + {$region Кадр} + gl.Clear(glClearBufferMask.COLOR_BUFFER_BIT); + + var blocks := render_info.block_area.MakeInitedBlocksMatr; + var b_cy := blocks.GetLength(0); + var b_cx := blocks.GetLength(1); + + var render_sheet_w := b_cx * block_w; + var render_sheet_h := b_cy * block_w; + + var need_back_sheet := not draw_alt_mode + and (render_info.last_sheet_diff<>nil) + and not render_info.last_sheet_diff.Value.IsNoChange; - var Q_Steps := CQNil; //TODO Calculate in separate thread - var Q_Extract := CQNil; + var Q_Acquire := CQNil; + var Q_Release := CQNil; + var Q_Init := CQNil; + if need_back_sheet then + begin + Q_Acquire += CQAcquireGL(sheet_draw.b_cl); + Q_Release += CQReleaseGL(sheet_draw.b_cl); + Swap(sheet_back, sheet_draw); + end; + var need_zero_out := sheet_draw.EnsureLen(render_sheet_w * render_sheet_h); + Q_Acquire += CQAcquireGL(sheet_draw.b_cl); + Q_Release += CQReleaseGL(sheet_draw.b_cl); + if need_back_sheet or need_zero_out then + Q_Init += sheet_draw.b_cl.MakeCCQ.ThenFillValue(0).DiscardResult; + if need_back_sheet then + Q_Init += render_info.last_sheet_diff.Value.CQ_CopySheet(sheet_back.b_cl, sheet_draw.b_cl, last_render_sheet_w, render_sheet_w, render_sheet_h); + last_render_sheet_w := render_sheet_w; + + //TODO С текущими очередями OpenCLABC не получится тут использовать имеющуюся очередь + // - Проблема в том что CQ_GetData может выдавать CQNil, если блок создан, но его vram_data не инициализировано + // - Чтобы это решить надо сначала доделать ветвление в OpenCLABC + var Q_Extract := V_ExtractedCount.MakeCCQ.ThenWriteValue(0).DiscardResult; for var b_y := 0 to b_cy-1 do for var b_x := 0 to b_cx-1 do begin - var b := render_info.blocks[b_y, b_x]; - - Q_Steps += b.CQ_MandelbrotBlockStep(max_steps, cl_uc_buffer, cl_err_buffer); - - var sheet_shift := render_block_size * (b_x + b_y*render_sheet_w); - Q_Extract += b.CQ_GetData(render_info.mipmap_lvl - , new ShiftedCLArray(cl_sheet_states, sheet_shift, render_sheet_w) - , new ShiftedCLArray(cl_sheet_buffer, sheet_shift, render_sheet_w) + var b := blocks[b_y, b_x]; + if b=nil then continue; + var sheet_shift := block_w * (b_x + b_y*render_sheet_w); + Q_Extract += b.CQ_GetData( + new ShiftedCLArray(sheet_draw.b_cl, sheet_shift, render_sheet_w), + V_ExtractedCount ); - end; - begin - var cl_err := CLContext.Default.SyncInvoke( - Q_Init + - Q_Steps + - Q_Extract + - cl_err_buffer.MakeCCQ.ThenGetArray - ); - if cl_err[0]<>0 then - $'OpenCL err at [{cl_err[1]},{cl_err[2]}]: {CLCodeExecutionError(cl_err[0])}'.Println; - end; +// var sw := Stopwatch.StartNew; + var extracted_count := CLContext.Default.SyncInvoke( + Q_Acquire + + Q_Init + + Q_Extract + + Q_Release + + V_ExtractedCount.MakeCCQ.ThenGetValue + ); +// Println(sw.Elapsed); uniform_view_bound.Write(render_info.view_bound); uniform_sheet_bound.Write(render_info.sheet_bound); gl.Uniform2i(uniform_sheet_size, render_sheet_w, render_sheet_h); - gl.BindBufferBase(glBufferTarget.SHADER_STORAGE_BUFFER, ssb_sheet, gl_sheet_buffer); - gl.Uniform1f(uniform_max_steps, max_steps); + gl.BindBufferBase(glBufferTarget.SHADER_STORAGE_BUFFER, ssb_sheet, sheet_draw.b_gl); // Для дебага // gl.NamedBufferSubData(buffer_temp, new IntPtr(0*sizeof(real)), new IntPtr(2*sizeof(real)), mouse_pos); @@ -332,37 +477,57 @@ BoundUniforms = record // var temp_data := new real[1]; // gl.GetNamedBufferSubData(buffer_temp, new IntPtr(2*sizeof(real)), new IntPtr(1*sizeof(real)), temp_data); +// gl.BindBufferBase(glBufferTarget.SHADER_STORAGE_BUFFER, ssb_sheet, gl_buffer.Zero); + + {$endregion Кадр} + gl.Finish; //TODO Использовать обмент ивентами OpenCL/OpenGL var err := gl.GetError; if err.IS_ERROR then MessageBox.Show(err.ToString); - gl.Finish; + t_body.Stop; if curr_frame_resized then need_resize := false; - var curr_time := sw.Elapsed; var title_parts := new List; // Для дебага // title_parts += $'temp_data={_ObjectToString(temp_data)}'; - //TODO Оттестировать и убрать - title_parts += $'sheet byte size={curr_sheet_size} (${curr_sheet_size:X})'; +// title_parts += $'mem={CLMemoryObserver.Current.CurrentlyUsedAmount}'; +// title_parts += $'rendering {b_cx} x {b_cy} blocks'; +// title_parts += $'sheet byte size={curr_sheet_size} (${curr_sheet_size:X})'; - if timings.Count=timings_max_count then - begin - var time_diff := curr_time - timings.Dequeue; - var mspf := time_diff.TotalMilliseconds / timings_max_count; - var fps := 1000/mspf; - title_parts += $'{fps:N2} fps'; - title_parts += $'{mspf:N2} mspf'; - end; - timings += curr_time; + t_full.Update; + t_body.Update; + title_parts += $'{t_full.UPS:N2} fps'; + title_parts += $'{t_full.MSPU:N2} full mspf'; +// title_parts += $'{t_body.UPS:N2} body fps'; + title_parts += $'{t_body.MSPU:N2} body mspf'; - title_parts += $'pos=({camera.pos.r}; {camera.pos.i})'; title_parts += $'scale={camera.scale_fine:N3}*2^{camera.scale_pow}'; +// title_parts += $'pos=({camera.pos.r}; {camera.pos.i})'; + + title_parts += $'{1-extracted_count/render_sheet_w/render_sheet_h:00.00%} old'; + + if BlockUpdater.StepInfoStr<>nil then + title_parts += BlockUpdater.StepInfoStr; + +// title_parts += $'RAM: {GC.GetTotalMemory(true)/1024/1024/1024:N5} GC vs {System.Diagnostics.Process.GetCurrentProcess.WorkingSet64/1024/1024/1024:N5} Process'; + + if BlockUpdater.LackingVRAM then + title_parts += $'LACKING VRAM!!!'; - f.Text := title_parts.JoinToString(', '); + f.BeginInvoke(()-> + try + f.Text := title_parts.JoinToString(', '); + except + on e: Exception do + MessageBox.Show(e.ToString); + end); EndFrame; + except + on e: Exception do + Println(e); end; except diff --git a/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.td b/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.td index d84ad68f..a8ed93d6 100644 --- a/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.td +++ b/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.td @@ -1,9 +1,72 @@  -#ExpErr -Compile errors: -[225,27] Blocks.pas: The type 'PointComponent' does not contain a definition for 'FirstWordToReal' +#Delegates +0Mandelbrot.$delegate? = procedure +0Mandelbrot.$delegate? = procedure +0Mandelbrot.$delegate? = procedure +0Mandelbrot.$delegate? = procedure +0Mandelbrot.$delegate? = procedure +0Mandelbrot.$delegate? = procedure +0Mandelbrot.$delegate? = procedure +0Mandelbrot.$delegate? = procedure +0Mandelbrot.$delegate? = procedure +0Mandelbrot.$delegate? = procedure(<>key_low: Keys; <>key_high: Keys; <>modifiers: Keys; <>on_change: integer -> ()) +0Mandelbrot.$delegate? = procedure(<>key_low: Keys; <>key_high: Keys; <>modifiers: Keys; <>on_change: integer -> ()) +0Mandelbrot.$delegate? = procedure(<>s_frag: gl_shader) +0Mandelbrot.$delegate? = procedure(<>s_frag: gl_shader) +0Mandelbrot.$delegate? = procedure(key_low: Keys; key_high: Keys; modifiers: Keys; on_change: integer -> ()) +0Mandelbrot.$delegate? = procedure(s_frag: gl_shader) +Blocks.$delegate? = procedure(<>ml: MemoryLayer) +Blocks.$delegate? = procedure(<>ml: MemoryLayer) +Blocks.$delegate? = procedure(e: Exception) +Blocks.$delegate? = procedure(ml: MemoryLayer) +OpenCL.$delegate? = function(command_queue: cl_command_queue; num_objects: longword; var mem_objects: cl_mem; num_events_in_wait_list: longword; var event_wait_list: cl_event; var event: cl_event): clErrorCode +OpenCL.$delegate? = function(command_queue: cl_command_queue; num_objects: longword; var mem_objects: cl_mem; num_events_in_wait_list: longword; var event_wait_list: cl_event; var event: cl_event): clErrorCode +OpenCL.$delegate? = function(context: cl_context; flags: clMemFlags; bufobj: longword; var errcode_ret: clErrorCode): cl_mem +OpenCL.clCreateContextCallback = procedure(errinfo: string; private_info: System.IntPtr; cb: System.UIntPtr; user_data: System.IntPtr) +OpenCL.clEventCallback = procedure(event: cl_event; event_command_status: clCommandExecutionStatus; user_data: System.IntPtr) +OpenCL.clProgramCallback = procedure(program: cl_program; user_data: System.IntPtr) +OpenCLABC.$delegate? = function(ntv: cl_program; var data: clBool; validate: boolean): clErrorCode +OpenCLABC._GetPropValueFunc = function(ntv: cl_program; var data: T): clErrorCode +OpenCLABC_implementation______.EnqFunc = function(prev_res: T; cq: cl_command_queue; ev_l2: EventList): ValueTuple ()> +OpenCLABC_implementation______.InvokeParamsFunc = function(enq_c: integer; o_const: boolean; g: CLTaskGlobalData; enq_evs: DoubleList; par_err_handlers: DoubleList): ValueTuple> +OpenGL.$delegate? = function(interval: integer): glBool32 +OpenGL.$delegate? = function(program: gl_program; name: System.IntPtr): integer +OpenGL.$delegate? = function(program: gl_program; programInterface: glProgramInterface; name: System.IntPtr): longword +OpenGL.$delegate? = function(type: glShaderType): gl_shader +OpenGL.$delegate? = function: gl_program +OpenGL.$delegate? = function: glErrorCode +OpenGL.$delegate? = procedure +OpenGL.$delegate? = procedure(buffer: gl_buffer; size: System.UIntPtr; data: System.IntPtr; usage: glVertexBufferObjectUsage) +OpenGL.$delegate? = procedure(buffer: gl_buffer; size: System.UIntPtr; var data: byte; usage: glVertexBufferObjectUsage) +OpenGL.$delegate? = procedure(location: integer; v0: integer; v1: integer) +OpenGL.$delegate? = procedure(location: integer; v0: single) +OpenGL.$delegate? = procedure(mask: glClearBufferMask) +OpenGL.$delegate? = procedure(mode: glPrimitiveType; first: integer; count: integer) +OpenGL.$delegate? = procedure(n: integer; buffers: System.IntPtr) +OpenGL.$delegate? = procedure(n: integer; var buffers: gl_buffer) +OpenGL.$delegate? = procedure(program: gl_program) +OpenGL.$delegate? = procedure(program: gl_program) +OpenGL.$delegate? = procedure(program: gl_program) +OpenGL.$delegate? = procedure(program: gl_program; bufSize: integer; length: System.IntPtr; infoLog: System.IntPtr) +OpenGL.$delegate? = procedure(program: gl_program; bufSize: integer; var length: integer; infoLog: System.IntPtr) +OpenGL.$delegate? = procedure(program: gl_program; pname: glProgramProperty; params: System.IntPtr) +OpenGL.$delegate? = procedure(program: gl_program; pname: glProgramProperty; var params: integer) +OpenGL.$delegate? = procedure(program: gl_program; shader: gl_shader) +OpenGL.$delegate? = procedure(shader: gl_shader) +OpenGL.$delegate? = procedure(shader: gl_shader) +OpenGL.$delegate? = procedure(shader: gl_shader; bufSize: integer; length: System.IntPtr; infoLog: System.IntPtr) +OpenGL.$delegate? = procedure(shader: gl_shader; bufSize: integer; var length: integer; infoLog: System.IntPtr) +OpenGL.$delegate? = procedure(shader: gl_shader; count: integer; string: Void*; length: System.IntPtr) +OpenGL.$delegate? = procedure(shader: gl_shader; count: integer; string: Void*; var length: integer) +OpenGL.$delegate? = procedure(shader: gl_shader; count: integer; var string: System.IntPtr; length: System.IntPtr) +OpenGL.$delegate? = procedure(shader: gl_shader; count: integer; var string: System.IntPtr; var length: integer) +OpenGL.$delegate? = procedure(shader: gl_shader; pname: glShaderParameterName; params: System.IntPtr) +OpenGL.$delegate? = procedure(shader: gl_shader; pname: glShaderParameterName; var params: integer) +OpenGL.$delegate? = procedure(target: glBufferTarget; index: longword; buffer: gl_buffer) +OpenGL.$delegate? = procedure(x: integer; y: integer; width: integer; height: integer) +OpenGLABC.RedrawThreadProc = procedure(pl: IGLPlatformLoader; EndFrame: procedure) #ReqModules OpenGL+OpenGLABC diff --git a/Samples/OpenGLABC/Mandelbrot/Blocks.pas b/Samples/OpenGLABC/Mandelbrot/Blocks.pas index e161cd4d..1047f2fc 100644 --- a/Samples/OpenGLABC/Mandelbrot/Blocks.pas +++ b/Samples/OpenGLABC/Mandelbrot/Blocks.pas @@ -8,6 +8,7 @@ uses PointComponents; uses CameraDef; uses MandelbrotSampling; +uses MemoryLayering; type CLCodeExecutionError = MandelbrotSampling.CLCodeExecutionError; @@ -26,272 +27,672 @@ ShiftedCLArray = record end; + BlockLayer = class; // Блок из block_w*block_w точек - PointBlock = sealed class + PointBlock = sealed class(IMemoryLayerData) + private layer: BlockLayer; // Принимает значения -∞..1 - // Длина стороны блока в логическом пространстве = 2**block_scale - private block_scale: integer; - private component_word_count: integer; - private pos00: CLArray; + // Длина стороны блока в логическом пространстве = 2**block_scale_pow + private block_scale_pow: integer; + private pos00: PointPos; - private gpu_data: CLArray; - private gpu_mipmaps_state: CLArray; - private gpu_mipmaps_steps: CLArray; - private gpu_mipmaps_need_update: CLArray; + private vram_pos00: CLArray; + private vram_data: CLArray := nil; - private ram_data: array of byte; + private ram_data: array of cardinal := nil; - public constructor(block_scale: integer; pos00: PointPos); + private drive_cache_file: string := nil; + // Current block format: + // + // version: int32 := cache_format_version + // scale: int32 + // word_count: int32 + // pos00: PointPos + // data: array of uint32 + // + // Also used in Drive_MemoryLayer + private const drive_cache_format_version = 1; + private static last_drive_cache_id := 0; + private static drive_cache_dir_name := $'Cache/block_w_pow={Settings.block_w_pow}'; + private const drive_cache_ext_name = 'point_block'; + private static function AllocDriveCacheFile: string; begin - self.block_scale := block_scale; - if block_scale>=2 then + while true do + begin + var id := System.Threading.Interlocked.Increment(last_drive_cache_id); + Result := $'{drive_cache_dir_name}/{id}.{drive_cache_ext_name}'; + if not FileExists(Result) then break; + end; + end; + + public constructor(layer: BlockLayer; block_scale_pow: integer; pos00: PointPos); + begin + self.layer := layer; + + self.block_scale_pow := block_scale_pow; + if block_scale_pow>Settings.max_block_scale_pow then raise new System.ArgumentOutOfRangeException; - self.component_word_count := pos00.Size; - //TODO Track GPU memory used amount - self.pos00 := new CLArray(pos00.r.Words+pos00.i.Words, CLMemoryUsage.ReadOnly, CLMemoryUsage.None); - - self.gpu_data := new CLArray( block_w*block_w * (2 + component_word_count*2) ); - self.gpu_mipmaps_state := new CLArray(mipmap_total_size); - self.gpu_mipmaps_steps := new CLArray(mipmap_total_size); - self.gpu_mipmaps_need_update := new CLArray(mipmap_total_size); - - //TODO Снять с этого потока выполнения... - CLContext.Default.SyncInvoke( - self.gpu_data.MakeCCQ.ThenFillValue(0) + - self.gpu_mipmaps_state.MakeCCQ.ThenFillValue(0) + - self.gpu_mipmaps_steps.MakeCCQ.ThenFillValue(0) + - self.gpu_mipmaps_need_update.MakeCCQ.ThenFillValue(0) - ); + self.pos00 := pos00; end; private constructor := raise new System.InvalidOperationException; + public static function GetMetaWordCount(component_word_count: integer) := component_word_count*2; + public static function GetDataWordCount(component_word_count: integer) := block_w.Sqr * (2 + component_word_count*2); + public static function GetWordCount(component_word_count: integer) := GetMetaWordCount(component_word_count) + GetDataWordCount(component_word_count); + public static function GetByteSize(component_word_count: integer) := GetWordCount(component_word_count) * sizeof(cardinal); + public function GetByteSize := GetByteSize(self.pos00.Size); + private static CLCodeCache := new System.Collections.Concurrent.ConcurrentDictionary; private static function CLCodeFor(word_c: cardinal) := CLCodeCache.GetOrAdd(word_c, word_c->MandelbrotSampling.CompiledCode(word_c)); - public function CQ_MandelbrotBlockStep(step_repeat_count: CommandQueue; V_UpdateCount: CLValue; A_Err: CLArray): CommandQueueNil := - CLCodeFor(self.component_word_count)['MandelbrotBlockSteps'] - .MakeCCQ.ThenExec2(block_w,block_w - , self.gpu_data - , self.pos00 - , Settings.z_int_bits-1 + -(self.block_scale-Settings.block_w_pow) - , self.gpu_mipmaps_need_update - , step_repeat_count - , V_UpdateCount - , A_Err - ).DiscardResult; - - public function CQ_GetData(target_mipmap_lvl: integer; A_State: ShiftedCLArray; A_Steps: ShiftedCLArray): CommandQueueNil; + + public function CQ_DownGradeToRAM: CommandQueueNil; begin + Result := CQNil; + var l_vram_data := System.Threading.Interlocked.Exchange(self.vram_data, nil); + if l_vram_data=nil then exit; - {$ifdef DEBUG} - if target_mipmap_lvl < -scale_shift then - raise new System.InvalidOperationException; - if target_mipmap_lvl > block_w_pow then - raise new System.InvalidOperationException; - {$endif DEBUG} + vram_pos00.Dispose; + vram_pos00 := nil; + + if ram_data<>nil then raise new System.InvalidOperationException; + var l_ram_data := new cardinal[l_vram_data.Length]; + ram_data := l_ram_data; + + //TODO l_vram_data may still be used by the render thread + // - Maybe, instead of disposing, add to some list that render thread can dispose at the end of frame + //TODO #????: & + Result += l_vram_data.MakeCCQ.ThenReadArray(HFQ&(()->l_ram_data, need_own_thread := false)) + + HPQ(()-> + begin + l_vram_data.Dispose; + //TODO Протестить как теперь тратится память. Должно убрать скачёк при сохранении + // - Нет, проблема ж вообще не в этом + // - На самом деле все l_ram_data выделяются вместе, перед началом выполнения очереди + // - Это очень фиговый дизайн получился... + l_ram_data := nil; + end, need_own_thread := false); + + end; + public function CQ_DownGradeToDrive: CommandQueueNil; + begin + Result := CQ_DownGradeToRAM; + + var l_ram_data := System.Threading.Interlocked.Exchange(self.ram_data, nil); + if l_ram_data=nil then exit; - var code := CLCodeFor(self.component_word_count); - var w := block_w; + if drive_cache_file<>nil then raise new System.InvalidOperationException; + var l_drive_cache_file := AllocDriveCacheFile; + drive_cache_file := l_drive_cache_file; - if target_mipmap_lvl=0 then + Result += HPQ(()-> begin - Result := code['ExtractRawSteps'].MakeCCQ.ThenExec2(w,w - , self.gpu_data - , A_State.a, A_State.shift, A_State.row_len - , A_Steps.a, A_Steps.shift, A_Steps.row_len - ).DiscardResult; - exit; - end; + var need_del := true; + var temp_fname := l_drive_cache_file + '.temp'; + var str := System.IO.File.Create(temp_fname); + try + var bw := new System.IO.BinaryWriter(str); + + bw.Write(drive_cache_format_version); + bw.Write(self.block_scale_pow); + bw.Write(self.pos00.Size); + self.pos00.Save(bw); + + // Much faster than reading/writing one word at a time + // But could be even faster, if stream directly accessed l_ram_data + var bytes := new byte[l_ram_data.Length * sizeof(cardinal)]; + System.Buffer.BlockCopy(l_ram_data,0, bytes,0, bytes.Length); + l_ram_data := nil; + str.Write(bytes, 0, bytes.Length); + + need_del := false; + finally + str.Close; + if need_del then + DeleteFile(temp_fname); + end; + System.IO.File.Move(temp_fname, l_drive_cache_file); + end, need_own_thread := false); + + end; + + public function CQ_UpGradeToRAM: CommandQueueNil; + begin + Result := CQNil; + var l_drive_cache_file := System.Threading.Interlocked.Exchange(drive_cache_file, nil); + if l_drive_cache_file=nil then exit; - w := w shr 1; - Result := code['FixFirstMipMap'].MakeCCQ.ThenExec2(w,w - , self.gpu_data - , self.gpu_mipmaps_state - , self.gpu_mipmaps_steps - , self.gpu_mipmaps_need_update - ).DiscardResult; - var mipmap_shift := w*w; + if ram_data<>nil then raise new System.InvalidOperationException; + var l_ram_data := new cardinal[GetDataWordCount(self.pos00.Size)]; + ram_data := l_ram_data; - for var mipmap_lvl := 2 to target_mipmap_lvl do + Result += HPQ(()-> begin - w := w shr 1; - Result += code['FixMipMap'].MakeCCQ.ThenExec2(w,w - , self.gpu_mipmaps_state - , self.gpu_mipmaps_steps - , self.gpu_mipmaps_need_update - , cardinal(mipmap_shift) - , cardinal(mipmap_lvl) + var str := System.IO.File.OpenRead(l_drive_cache_file); + try + str.Position := 12 + GetMetaWordCount(self.pos00.Size)*sizeof(cardinal); + + var bytes := new byte[l_ram_data.Length * sizeof(cardinal)]; + var read_len := str.Read(bytes, 0, bytes.Length); + if read_len<>bytes.Length then raise new System.InvalidOperationException; + + System.Buffer.BlockCopy(bytes,0, l_ram_data,0, bytes.Length); + finally + str.Close; + DeleteFile(l_drive_cache_file); + end; + end, need_own_thread := false); + + end; + public function CQ_UpGradeToVRAM: CommandQueueNil; + begin + Result := CQ_UpGradeToRAM; + var l_ram_data := System.Threading.Interlocked.Exchange(self.ram_data, nil); + // Init vram buffer even if there is no existing data +// if l_ram_data=nil then exit; +// if vram_data<>nil then raise new System.InvalidOperationException; + // Instead, exit if vram is already inited + if vram_data<>nil then exit; + + if vram_pos00<>nil then raise new System.InvalidOperationException; + vram_pos00 := new CLArray(pos00.r.Words + pos00.i.Words, CLMemoryUsage.ReadOnly, CLMemoryUsage.None); + + var l_vram_data := new CLArray(GetDataWordCount(self.pos00.Size), CLMemoryUsage.ReadWrite, CLMemoryUsage.ReadWrite); + vram_data := l_vram_data; + + Result += if l_ram_data=nil then + l_vram_data.MakeCCQ.ThenFillValue(0).DiscardResult else + l_vram_data.MakeCCQ.ThenWriteArray(l_ram_data).DiscardResult; + + end; + + public function CQ_MandelbrotBlockStep(Q_GetStepRepeatCount: CommandQueue; V_UpdateCount: CLValue; A_Err: CLArray): CommandQueueNil; + begin + Result := CQNil; + if vram_data=nil then exit; + + Result := CLCodeFor(self.pos00.Size)['MandelbrotBlockSteps'] + .MakeCCQ.ThenExec2(block_w,block_w + , vram_data + , vram_pos00 + , Settings.z_int_bits-1 + -(self.block_scale_pow-Settings.block_w_pow) + , Q_GetStepRepeatCount + , V_UpdateCount + , A_Err ).DiscardResult; - mipmap_shift += w*w; - end; - Result += code['ExtractMipMapSteps'].MakeCCQ.ThenExec2(w,w - , self.gpu_mipmaps_state, self.gpu_mipmaps_steps, cardinal(mipmap_shift - w*w) - , A_State.a, A_State.shift, A_State.row_len - , A_Steps.a, A_Steps.shift, A_Steps.row_len - , cardinal(target_mipmap_lvl) - ).DiscardResult; + end; + + public function CQ_GetData(A_Result: ShiftedCLArray; V_ExtractedCount: CLValue): CommandQueueNil; + begin + Result := CQNil; + + var l_vram_data := self.vram_data; + if l_vram_data=nil then exit; + + Result += CLCodeFor(self.pos00.Size)['ExtractSteps'] + .MakeCCQ.ThenExec2(block_w,block_w + , l_vram_data + , A_Result.a, A_Result.shift, A_Result.row_len + , V_ExtractedCount + ).DiscardResult; + + end; + + public function ToString: string; override := + $'block at scale_pow={self.block_scale_pow}, pos={self.pos00}'; + + public procedure Dispose; + begin + pos00 := default(PointPos); + + var l_vram_pos00 := System.Threading.Interlocked.Exchange(self.vram_pos00, nil); + if l_vram_pos00<>nil then l_vram_pos00.Dispose; + + var l_vram_data := System.Threading.Interlocked.Exchange(self.vram_data, nil); + if l_vram_data<>nil then l_vram_data.Dispose; + ram_data := nil; + + var l_drive_cache_file := System.Threading.Interlocked.Exchange(self.drive_cache_file, nil); + if l_drive_cache_file<>nil then DeleteFile(l_drive_cache_file); + + GC.SuppressFinalize(self); end; + protected procedure Finalize; override := Dispose; end; - BoundDefs = record - public xf, yf: single; - public xl, yl: single; + BoundDefs = record + public xf, yf: T; + public xl, yl: T; + + public function Convert(dx, dy: T->T2): BoundDefs; + begin + Result.xf := dx(self.xf); Result.yf := dy(self.yf); + Result.xl := dx(self.xl); Result.yl := dy(self.yl); + end; + public function Convert(d: T->T2): BoundDefs := Convert(d,d); + + public function ToString: string; override := + $'{xf}<=x=>{xl} | {yf}<=y=>{yl}'; + + end; + BlockLayerSubArea = record + private c_min, c_max: PointPos; + private r_block_poss, i_block_poss: array of PointComponent; + private layer: BlockLayer; + + /// Ordered by distance from the center + public function MakeOrderedArr(sel: PointPos->T): array of T; + begin + var rc := r_block_poss.Length; + var ic := i_block_poss.Length; + Result := new T[rc*ic]; + var keys := new int64[Result.Length]; + + var ind := 0; + for var i_ind := 0 to ic-1 do + for var r_ind := 0 to rc-1 do + begin + Result[ind] := sel(new PointPos(r_block_poss[r_ind], i_block_poss[i_ind])); + // Doesn't account for sheet_bound + // But in practice it would hardly matter + keys[ind] := Sqr(r_ind*2-rc) + Sqr(i_ind*2-ic); + ind += 1; + end; + {$ifdef DEBUG} + if ind<>Result.Length then + raise new System.InvalidOperationException; + {$endif DEBUG} + + System.Array.Sort(keys, Result); + end; + + end; + SheetDiff = record + + // All positive (or zero) if new sheet is fully inside old sheet + // All negative (or zero) if old sheet is fully inside new sheet + // Counted in terms of points of new sheet + private bounds_diff: BoundDefs; - public static procedure operator*=(var b: BoundDefs; k: System.ValueTuple); + // Positive if new sheet covers smaller scale than old sheet + private scale_diff: integer; + + {$resource SheetTransfer.cl} + private static sheet_transfer_code_text := System.IO.StreamReader.Create( + System.Reflection.Assembly.GetCallingAssembly.GetManifestResourceStream('SheetTransfer.cl') + ).ReadToEnd; + private static sheet_transfer_code := new CLProgramCode(sheet_transfer_code_text); + + public function IsNoChange := (bounds_diff = default(BoundDefs)) and (scale_diff=0); + + public function CQ_CopySheet(old_sheet, new_sheet: CLArray; old_row_len, new_row_len, new_col_len: integer): CommandQueueNil; begin - var (kx,ky) := k; - b.xf *= kx; b.yf *= ky; - b.xl *= kx; b.yl *= ky; + + // Both count in points of new scale + var old_bounds := self.bounds_diff.Convert(x->(+x).ClampBottom(0)); + var new_bounds := self.bounds_diff.Convert(x->(-x).ClampBottom(0)); + + var scale_k := 1 shl Abs(scale_diff); + + // To continue counting everything in points of new sheet, + // the scale_diff<0 should cause "old_row_len/scale_k", converting it to new points + // But that will either loose precision, or require floating point (or both) + // Same problem is old_shift is always in points of old sheet + // So instead, if scale_diff<0, old_bounds is coverted to points of old sheet +// var old_shift := old_row_len*old_bounds.yf*scale_k + old_bounds.xf * if scale_diff<0 then 1 else scale_k; + var new_shift := new_row_len*new_bounds.yf + new_bounds.xf; + + var w := new_row_len - (new_bounds.xf+new_bounds.xl); + var h := new_col_len - (new_bounds.yf+new_bounds.yl); + + if scale_diff<0 then + begin + var old_shift := (old_row_len*old_bounds.yf + old_bounds.xf)*scale_k; + + Result := sheet_transfer_code['DownScaleSheet'].MakeCCQ + .ThenExec2(w,h + , old_sheet, old_shift, old_row_len + , new_sheet, new_shift, new_row_len + , -scale_diff + ).DiscardResult; + + end else + if scale_diff=0 then + begin + var old_shift := old_row_len*old_bounds.yf + old_bounds.xf; + + Result := sheet_transfer_code['CopySheetRect'].MakeCCQ + .ThenExec2(w,h + , old_sheet, old_shift, old_row_len + , new_sheet, new_shift, new_row_len + ).DiscardResult; + + end else + if scale_diff>0 then + begin + + Result := sheet_transfer_code['UpScaleSheet'].MakeCCQ + .ThenExec2(w,h + , old_sheet, old_bounds.xf, old_bounds.yf, old_row_len + , new_sheet, new_shift, new_row_len + , +scale_diff + ).DiscardResult; + + end else + raise new System.InvalidOperationException; + end; - public static procedure operator/=(var b: BoundDefs; k: System.ValueTuple) := - b *= System.ValueTuple.Create(1/k.Item1, 1/k.Item2); end; BlockLayerRenderInfo = record - public blocks: array[,] of PointBlock; // [y,x] - public mipmap_lvl: integer; // How much of viewport is empty + // (between screen edge and the sheet) // [0;2) and first+last<=2 - public view_bound: BoundDefs; + public view_bound: BoundDefs; - // How much of edge blocks is hidden + // How much of edge blocks is hidden by window edge // [0;1) and first+last<=1 - public sheet_bound: BoundDefs; + public sheet_bound: BoundDefs; + // All coordinates of (partially or fully) visible blocks + public block_area: BlockLayerSubArea; + + // Info needed to render + public last_sheet_diff: SheetDiff?; + + private block_sz_bit_ind: integer; end; // Слой, содержащий кэш уже просчитанных блоков BlockLayer = sealed class - private scale: integer; + private scale_pow: integer; private blocks := new Dictionary; - public constructor(scale: integer); + public constructor(scale_pow: integer) := self.scale_pow := scale_pow; + private constructor := raise new System.InvalidOperationException; + + private static all_layers := new List; + public static function GetLayer(block_scale_pow: integer): BlockLayer; begin - self.scale := scale; + var layer_ind := Settings.max_block_scale_pow - block_scale_pow; + while all_layers.Count<=layer_ind do + all_layers += default(BlockLayer); + Result := all_layers[layer_ind]; + if Result<>nil then exit; + Result := new BlockLayer(block_scale_pow); + all_layers[layer_ind] := Result; end; - private constructor := raise new System.InvalidOperationException; - -// private static all_layers := new BlockLayer[1+max_z_scale_bits_rounded]; -// private static function scale_to_layer_ind(scale: integer) := 1-scale; -// -// public static function TakeBlocks(scale: integer; dx, dy: cardinal): List; -// begin -// var layer_ind := scale_to_layer_ind(scale).Clamp(0, all_layers.Length-1); -// scale := 1-layer_ind; -// -// -// end; -// -// public static procedure Cleanup; -// begin -// var TODO := 0; // Вообще вместо этого надо чистить когда заканчивается место -// end; + public static function GetLayer(camera_pos: CameraPos) := + GetLayer(camera_pos.GetPointScalePow + Settings.block_w_pow); - //TODO Сейчас блоки создаёт с 0 при каждом вызове. Использовать экземпляр типа BlockLayer - public static function BlocksForCurrentScale(camera_pos: CameraPos): BlockLayerRenderInfo; + private function GetBlockAt(pos00: PointPos; can_create: boolean): PointBlock; + begin + if self.blocks.TryGetValue(pos00, Result) then exit; + if not can_create then exit; + Result := new PointBlock(self, self.scale_pow, pos00); + self.blocks.Add(pos00, Result); + end; + public function GetRenderInfo(camera_pos: CameraPos; last_ri: BlockLayerRenderInfo?): BlockLayerRenderInfo; begin Result := default(BlockLayerRenderInfo); + {$ifdef DEBUG} + if self.scale_pow <> camera_pos.GetPointScalePow + Settings.block_w_pow then + raise new System.InvalidOperationException; + {$endif DEBUG} + var word_count := camera_pos.pos.Size; var c_ctr := camera_pos.pos; - //TODO Таким макаром view_size_bit_ind может быть <0, не говоря уже о переполнении сложения .AddLowestBits - // Если сильно отдалить камеру - значения границ не поместятся в c_min/c_max - // - Можно сразу складывать, округлять и ограничивать, всё в 1 операцию - // --- Но это будет 1 очень большая и сложная подпрограмма - // - Можно ограничивать отдельно вывод RoundToLowestBits и т.п. - // --- Но тогда границы блоков неправильно поставит относительно границ экрана - // - Можно ограничивать камеру - // --- Но обработать PointComponent границы вместе с fine+pow масштабом будет сложно... - // --- И очень широкое окно всё равно переполнит PointComponent - // - Можно использовать отдельный алгоритм для границ, в зависимости от того - нужна ли точность PointComponent - // --- Но это дубли кода, сложная проверка какой алгоритм выбрать и возможность дёрганья камеры при переходе от 1 алгоритма к другому - // - Можно менять область рендеринга если камера слишком отдалена, а внутри неё всё равно использовать обычный алгоритм с PointComponent - // --- Но тогда на каждый пиксель придётся очень много точек - //TODO В итоге решил ограничить .pos, чтобы оно за [-2;+2] не выходило - // - Тогда достаточно таки view_skip при большом отдалении делать, и больше ничего - if camera_pos.scale_pow>=1 then begin - if c_ctr.Size<>1 then raise new System.NotImplementedException; var cx := c_ctr.r.FirstWordToReal; var cy := c_ctr.i.FirstWordToReal; + var ar := camera_pos.AspectRatio; + + var visible_space_dy := camera_pos.scale_fine * 2.0**camera_pos.scale_pow; + var visible_space_dx := visible_space_dy * ar; + + // Whole logical space is -2 .. +2 + // Visible logical space is c_ctr-visible_space_d .. c_ctr+visible_space_d - //TODO Масштабировать cx и cy, привести их к диапазону -1..+1 + Result.view_bound.xf := (-2-(cx-visible_space_dx)).ClampBottom(0) / visible_space_dx; + Result.view_bound.xl := ((cx+visible_space_dx)-2).ClampBottom(0) / visible_space_dx; + + Result.view_bound.yf := (-2-(cy-visible_space_dy)).ClampBottom(0) / visible_space_dy; + Result.view_bound.yl := ((cy+visible_space_dy)-2).ClampBottom(0) / visible_space_dy; end; - var view_size_bit_ind := Settings.z_int_bits-1 - camera_pos.scale_pow; - var di := PointComponent.RoundToLowestBits(word_count, view_size_bit_ind, camera_pos.scale_fine); - var dr := PointComponent.RoundToLowestBits(word_count, view_size_bit_ind, camera_pos.scale_fine * camera_pos.AspectRatio); - var c_min := c_ctr.AddLowestBits(-dr,-di); - var c_max := c_ctr.AddLowestBits(+dr,+di); - c_max.SelfFlipIfMinusZero; + var dr := new PointComponentShift(word_count, camera_pos.scale_pow, camera_pos.scale_fine * camera_pos.AspectRatio); + var di := new PointComponentShift(word_count, camera_pos.scale_pow, camera_pos.scale_fine); + Result.block_area.c_min := c_ctr.WithShiftClamp2(-dr,-di, false); + Result.block_area.c_max := c_ctr.WithShiftClamp2(+dr,+di, true); - var (block_scale, main_mipmap_lvl) := camera_pos.GetPointScaleAndMainMipMapLvl; - block_scale += Settings.block_w_pow; - Result.mipmap_lvl := main_mipmap_lvl; - var block_sz_bit_ind := Settings.z_int_bits-1 + -block_scale; - c_min.SelfBlockRound(block_sz_bit_ind, false, Result.sheet_bound.xf, Result.sheet_bound.yf); - c_max.SelfBlockRound(block_sz_bit_ind, true, Result.sheet_bound.xl, Result.sheet_bound.yl); - //TODO Debug error for when bounds are >2 + Result.block_sz_bit_ind := Settings.z_int_bits-1 + -self.scale_pow; + Result.block_area.c_min.SelfBlockRound(Result.block_sz_bit_ind, false, Result.sheet_bound.xf, Result.sheet_bound.yf); + Result.block_area.c_max.SelfBlockRound(Result.block_sz_bit_ind, true, Result.sheet_bound.xl, Result.sheet_bound.yl); - var r_blocks_count := PointComponent.BlocksCount(c_min.r, c_max.r, block_sz_bit_ind); - var i_blocks_count := PointComponent.BlocksCount(c_min.i, c_max.i, block_sz_bit_ind); + Result.block_area.r_block_poss := PointComponent.Range(Result.block_area.c_min.r, Result.block_area.c_max.r, Result.block_sz_bit_ind); + Result.block_area.i_block_poss := PointComponent.Range(Result.block_area.c_min.i, Result.block_area.c_max.i, Result.block_sz_bit_ind); + Result.block_area.layer := self; - Result.sheet_bound /= new System.ValueTuple(r_blocks_count, i_blocks_count); + var kx: single := 1/Result.block_area.r_block_poss.Length; + var ky: single := 1/Result.block_area.i_block_poss.Length; + Result.sheet_bound := Result.sheet_bound.Convert(x->x*kx, y->y*ky); - var pc_rs := new PointComponent[r_blocks_count]; - begin - var pc_r := c_min.r; - for var ri := 0 to r_blocks_count-1 do - begin - pc_rs[ri] := pc_r; - pc_r := pc_r.MakeNextBlockBound(block_sz_bit_ind); - end; - {$ifdef DEBUG} - if pc_r<>c_max.r then - raise new System.InvalidOperationException; - {$endif DEBUG} + if last_ri<>nil then + try + var last_ri_v := last_ri.Value; + var last_sheet_diff: SheetDiff; + + last_sheet_diff.scale_diff := Result.block_sz_bit_ind - last_ri_v.block_sz_bit_ind; + if Abs(last_sheet_diff.scale_diff)>16 then + // Handling such jump in scale would require + // increasing precision in SheetDiff.CQ_CopySheet + // But it's not worth it, just recalculate the sheet + raise new System.OverflowException; + + var point_sz_bit_ind := Result.block_sz_bit_ind + Settings.block_w_pow; + last_sheet_diff.bounds_diff.xf := +PointComponent.BlocksCount(last_ri_v.block_area.c_min.r.WithBlockRound(point_sz_bit_ind, false), Result.block_area.c_min.r, point_sz_bit_ind); + last_sheet_diff.bounds_diff.yf := +PointComponent.BlocksCount(last_ri_v.block_area.c_min.i.WithBlockRound(point_sz_bit_ind, false), Result.block_area.c_min.i, point_sz_bit_ind); + last_sheet_diff.bounds_diff.xl := -PointComponent.BlocksCount(last_ri_v.block_area.c_max.r.WithBlockRound(point_sz_bit_ind, false), Result.block_area.c_max.r, point_sz_bit_ind); + last_sheet_diff.bounds_diff.yl := -PointComponent.BlocksCount(last_ri_v.block_area.c_max.i.WithBlockRound(point_sz_bit_ind, false), Result.block_area.c_max.i, point_sz_bit_ind); + + Result.last_sheet_diff := last_sheet_diff; +// if not last_sheet_diff.IsNoChange then +// begin +// Println($'{c_min} .. {c_max}'); +// Println($'{Result.block_area.r_block_poss.Length} x {Result.block_area.i_block_poss.Length}'); +// Println(last_sheet_diff.bounds_diff); +// end; + except + on e: System.OverflowException do ; end; - Result.blocks := new PointBlock[i_blocks_count, r_blocks_count]; -// $'Need {i_blocks_count} x {r_blocks_count} = {Result.blocks.Length} blocks'.Println; Halt; + end; + + end; + + //TODO Вытащить код для вывода KB и т.п. +// // VRAM/RAM/Drive +// MemoryLayer = sealed class +// private blocks := new List; +// private new_blocks := new List; +// private layer_name: string; +// private max_size: int64; +// private get_curr_size: ()->int64; +// private next := default(MemoryLayer); +// +// public constructor(layer_name: string; max_size: int64; get_curr_size: ()->int64); +// begin +// self.layer_name := layer_name; +// self.max_size := max_size; +// self.get_curr_size := get_curr_size; +// end; +// private constructor := raise new System.InvalidOperationException; +// +//// public procedure Flush; +//// begin +//// blocks.AddRange(new_blocks); +//// new_blocks.Clear; +//// end; +//// public procedure FlushAll := Enmr.ForEach(l->l.Flush); +// +// public function MemoryInfoStr: string; +// begin +// var c1 := real(get_curr_size()); +// var c2 := real(max_size); +// +// var pow := 0; +// var pow_step := 1024; +// var pow_names := |'KB','MB','GB'|; +// while (c1>=pow_step) or (c2>=pow_step) do +// begin +// if pow=pow_names.Length then break; +// pow += 1; +// c1 /= pow_step; +// c2 /= pow_step; +// end; +// +// var pow_name := if pow=0 then nil else ' '+pow_names[pow]; +// Result := $'{layer_name}: {c1}/{c2}{pow_name} ({c1/c2:000.00%})'; +// end; +// public function Enmr: sequence of MemoryLayer; +// begin +// var curr := self; +// repeat +// yield curr; +// curr := curr.next; +// until curr=nil; +// end; +// +// +// +// end; + + VRAM_MemoryLayer = sealed class(MemoryLayer) + + public constructor := inherited Create('VRAM', Settings.max_VRAM); + + protected function GetRealFilledSize: int64; override := CLMemoryObserver.Current.CurrentlyUsedAmount; + + end; + + RAM_MemoryLayer = sealed class(MemoryLayer) + + public constructor := inherited Create('RAM', Settings.max_RAM); + + protected function GetRealFilledSize: int64; override := GC.GetTotalMemory(true); + + end; + + Drive_MemoryLayer = sealed class(MemoryLayer) + + public constructor; + begin + inherited Create('Drive', Settings.max_drive_space); + + var cache_dir := System.IO.Directory.CreateDirectory(PointBlock.drive_cache_dir_name); + + var filled_size := int64(0); + foreach var fi in cache_dir.EnumerateFiles($'*.{PointBlock.drive_cache_ext_name}').OrderByDescending(fi->fi.LastWriteTimeUtc) do begin - var pc_i := c_min.i; - for var ii := 0 to i_blocks_count-1 do - begin - for var ri := 0 to r_blocks_count-1 do - Result.blocks[ii, ri] := new PointBlock(block_scale, new PointPos(pc_rs[ri], pc_i)); - pc_i := pc_i.MakeNextBlockBound(block_sz_bit_ind); + try + var str := fi.OpenRead; + try + var br := new System.IO.BinaryReader(str); + var version := br.ReadInt32; + var scale := br.ReadInt32; + var word_count := br.ReadInt32; + var pos00 := PointPos.Load(br, word_count); + if version<>PointBlock.drive_cache_format_version then + raise new Exception('Unsupported cache format'); + var expected_len := str.Position + PointBlock.GetDataWordCount(word_count)*sizeof(cardinal); + var actual_len := str.Length; + if expected_len <> actual_len then + raise new Exception($'Bad cache file size: {expected_len} vs {actual_len}'); + var layer := BlockLayer.GetLayer(scale); + var bl := new PointBlock(layer, scale, pos00); + bl.drive_cache_file := fi.FullName; + self.TryAdd(bl, bl->raise new Exception($'Not enough allocated disk space to load all the cache')); + layer.blocks.Add(pos00, bl); + filled_size += actual_len; + finally + str.Close; + end; + except + on e: Exception do + begin + $'Failed to load cached block: {fi.FullName}'.Println; + Println(e); + fi.Delete; + end; end; - {$ifdef DEBUG} - if pc_i<>c_max.i then - raise new System.InvalidOperationException; - {$endif DEBUG} end; + self.EndUpdate; // Flush newly added blocks + if filled_size + System.IO.DriveInfo.Create(GetCurrentDir).TotalFreeSpace < Settings.max_drive_space then + raise new System.InvalidOperationException('No enough disk space'); + end; + protected function GetRealFilledSize: int64; override := + System.IO.Directory.CreateDirectory(PointBlock.drive_cache_dir_name) + .EnumerateFiles($'*.{PointBlock.drive_cache_ext_name}').Select(fi->fi.Length).Sum; + end; + // Цикл обработки видимых блоков BlockUpdater = static class - private static output_update_info := false; private static procedure ExceptionToConsole(e: Exception) := Println(e); private static current_ex_handler := ExceptionToConsole; public static procedure SetExHandler(h: Exception->()) := current_ex_handler := h; - private static current_blocks: array[,] of PointBlock; - public static procedure SetCurrent(a: array[,] of PointBlock) := current_blocks := a; + private static current_area := default(System.Tuple); + public static procedure SetCurrent(area: BlockLayerSubArea) := current_area := Tuple.Create(area); + + private static lacking_vram := true; + public static property LackingVRAM: boolean read lacking_vram; + + private static step_info_str := default(string); + public static property StepInfoStr: string read step_info_str; + + private static shutdown_progress: (integer,integer) := nil; + private static on_shutdown_done: Action0 := nil; + public static property ShutdownProgress: (integer,integer) read shutdown_progress; + public static procedure BeginShutdown(on_shutdown_done: Action0); + begin + shutdown_progress := (0,1); + BlockUpdater.on_shutdown_done += on_shutdown_done; + end; + + private static function MatrItemsSortedFromCenter(m: array[,] of T): array of T; + begin + Result := new T[m.Length]; + var keys := new integer[Result.Length]; + + var res_i := 0; + var (c1,c2) := m.Size; + for var i1 := 0 to c1-1 do + for var i2 := 0 to c2-1 do + begin + Result[res_i] := m[i1,i2]; + keys[res_i] := Sqr(i1*2 - c1) + Sqr(i2*2 - c2); + res_i += 1; + end; + + System.Array.Sort(keys, Result); + end; static constructor := System.Threading.Thread.Create(()-> begin @@ -301,57 +702,158 @@ BlockLayerRenderInfo = record var A_Err := new CLArray(3); var err := new cardinal[3]; - var last_blocks: array[,] of PointBlock := nil; + var last_blocks := new HashSet; var Q_StepAll: CommandQueue; - var sw := new Stopwatch; + var step_sw := new Stopwatch; var step_count := 1; + var ml_vram := new VRAM_MemoryLayer; + var ml_ram := new RAM_MemoryLayer; + var ml_drive := new Drive_MemoryLayer; + + var all_mem_layers := new MemoryLayer[]( + ml_vram, ml_ram, ml_drive + ); + while true do try + if shutdown_progress<>nil then + begin + + var unloadable_blocks := new List; + var add_to_unload := procedure(ml: MemoryLayer)-> + begin + foreach var bl in ml.Enmr.Reverse do + if ml_drive.TryAdd(bl, bl->bl.Dispose()) then + unloadable_blocks += bl; + ml_drive.EndUpdate; + end; + add_to_unload( ml_ram); + add_to_unload(ml_vram); + + (ml_ram.Enmr+ml_vram.Enmr).ToArray; + foreach var bl in unloadable_blocks do + ml_drive.TryAdd(bl, bl->bl.Dispose()); + + var Q_ShutDown := CQNil; + foreach var bl in unloadable_blocks index bl_ind do + begin + Q_ShutDown += bl.CQ_DownGradeToDrive; + var shutdown_ind := bl_ind+1; + Q_ShutDown += HPQ(()->(shutdown_progress := (shutdown_ind,unloadable_blocks.Count)), need_own_thread := false); + end; + CLContext.Default.SyncInvoke(Q_ShutDown); + + var on_shutdown_done := on_shutdown_done; + if on_shutdown_done<>nil then on_shutdown_done(); + break; + end; - var blocks := current_blocks; - if blocks=nil then + var area: BlockLayerSubArea; begin - Sleep(1); - continue; + var area_t := current_area; + if area_t=nil then + begin + Sleep(1); + continue; + end; + area := area_t.Item1; end; - if blocks<>last_blocks then + var req_blocks := area.MakeOrderedArr(pos00->area.layer.GetBlockAt(pos00, true)); + if req_blocks.Length=0 then raise new System.NotImplementedException($'0 block area'); + + var block_new_table := new Dictionary(area.r_block_poss.Length * area.i_block_poss.Length); + foreach var bl in req_blocks do block_new_table.Add(bl, true); + foreach var ml in all_mem_layers do + ml.BeginUpdate(block_new_table); + + foreach var bl in req_blocks do + ml_vram.TryAdd(bl, bl-> + ml_ram.TryAdd(bl, bl-> + ml_drive.TryAdd(bl, bl-> + begin + if not bl.layer.blocks.Remove(bl.pos00) then + begin + $'Could not delete {bl}'.Println; + $'Current blocks:'.Println; + foreach var curr_bl in req_blocks do + $'{curr_bl}'.Println; + Println('='*50); + end; + bl.Dispose; + end) + ) + ); + + foreach var ml in all_mem_layers do + ml.EndUpdate; + + lacking_vram := ml_vram.Enmr.Take(req_blocks.Length).Count <> req_blocks.Length; + + begin + var Q_Init := CQNil; + + foreach var bl in ml_drive.Enmr do Q_Init += bl.CQ_DownGradeToDrive; + foreach var bl in ml_ram.Enmr do Q_Init += bl.CQ_DownGradeToRAM; + + foreach var bl in ml_ram.Enmr do Q_Init += bl.CQ_UpGradeToRAM; + foreach var bl in ml_vram.Enmr do Q_Init += bl.CQ_UpGradeToVRAM; + +// var init_sw := Stopwatch.StartNew; + CLContext.Default.SyncInvoke(Q_Init); +// $'Inited in {init_sw.Elapsed}'.Println; + end; + + var update_blocks := req_blocks.Intersect(ml_vram.Enmr).ToArray; + if not last_blocks.SetEquals(update_blocks) then begin + last_blocks.Clear; + last_blocks.UnionWith(update_blocks); + + // After blocks have changed, the update might take far longer + // Reset step count to avoid sudden horrible lag + step_count := 1; + var branches := ArrFill(Settings.max_parallel_blocks, CQNil); - foreach var b: PointBlock in blocks index b_i do - branches[b_i mod branches.Length] += b.CQ_MandelbrotBlockStep(P_StepCount, V_UpdateCount, A_Err); - last_blocks := blocks; + foreach var bl in update_blocks index bl_i do + branches[bl_i mod branches.Length] += bl.CQ_MandelbrotBlockStep(P_StepCount, V_UpdateCount, A_Err); + Q_StepAll := V_UpdateCount.MakeCCQ.ThenWriteValue(0) + A_Err.MakeCCQ.ThenFillValue(0) + CombineAsyncQueue(branches) + A_Err.MakeCCQ.ThenReadArray(err) + - V_UpdateCount.MakeCCQ.ThenGetValue + V_UpdateCount.MakeCCQ.ThenGetValue; + end; - sw.Restart; + {$ifdef DEBUG} + //TODO BlockLayer.all_layers использует из нескольких потоков, надо его блокировать... + var blocks_by_scale := BlockLayer.all_layers.Where(l->l<>nil).SelectMany(l->l.blocks.Values).ToArray; + var blocks_by_memory := all_mem_layers.SelectMany(l->l.Enmr).ToArray; + if not blocks_by_scale.ToHashSet.SetEquals(blocks_by_memory) then + raise new System.InvalidOperationException; + {$endif DEBUG} + + step_sw.Restart; var update_count := CLContext.Default.SyncInvoke(Q_StepAll , P_StepCount.NewSetter(step_count) ); - sw.Stop; + step_sw.Stop; - if output_update_info then - begin - $'Updated {update_count} points'.Println; - $'Updated current {blocks.Length} blocks {step_count} times in {sw.Elapsed}'.Println; - Println('='*30); - end; + step_info_str := $'u={update_count/step_count:N0}/step, {step_count} steps in {step_sw.Elapsed.TotalSeconds:N3}s'; // Тут бы PID контроллер реализовать вообще, потому что // зависимость времени от кол-ва шагов не линейная - // Но пока и так работает... - step_count := (step_count * Settings.target_step_time_seconds/sw.Elapsed.TotalSeconds).Clamp(1,max_steps_at_once).Round; + // Но на практике этого достаточно... + step_count := (step_count * Settings.target_step_time_seconds/step_sw.Elapsed.TotalSeconds).Clamp(1,max_steps_at_once).Round; if err[0]<>0 then - //TODO Надо бы выводить какой-то id блока тоже... + // Надо бы выводить какой-то id блока тоже... // - Достаточно [x,y] индекс. Тут из него можно получить все данные блока + // - Но пока что ошибок на стороне .cl кода (после FieldTest) вообще не видел. Слишком хорошо написал? raise new Exception($'Step err {CLCodeExecutionError(err[0])} at [{err[1]},{err[2]}]'); except @@ -363,4 +865,31 @@ BlockLayerRenderInfo = record end; +// For FieldTest, to go around the inconsistent BlockUpdater +procedure InitAllBlocks(self: BlockLayerSubArea); extensionmethod; +begin + var rc := self.r_block_poss.Length; + var ic := self.i_block_poss.Length; + for var i_ind := 0 to ic-1 do + for var r_ind := 0 to rc-1 do + self.layer.GetBlockAt(new PointPos( + self.r_block_poss[r_ind], self.i_block_poss[i_ind] + ), true); +end; + +function MakeInitedBlocksMatr(self: BlockLayerSubArea): array[,] of PointBlock; extensionmethod; +begin + var rc := self.r_block_poss.Length; + var ic := self.i_block_poss.Length; + Result := new PointBlock[ic, rc]; + for var i_ind := 0 to ic-1 do + for var r_ind := 0 to rc-1 do + // Cannot create blocks here, because they + // would not be marked by memory layers + // Instead BlockUpdater manages all creation and initing + Result[i_ind, r_ind] := self.layer.GetBlockAt(new PointPos( + self.r_block_poss[r_ind], self.i_block_poss[i_ind] + ), false); +end; + end. \ No newline at end of file diff --git a/Samples/OpenGLABC/Mandelbrot/CameraDef.pas b/Samples/OpenGLABC/Mandelbrot/CameraDef.pas index e57c5a7a..a21c6e68 100644 --- a/Samples/OpenGLABC/Mandelbrot/CameraDef.pas +++ b/Samples/OpenGLABC/Mandelbrot/CameraDef.pas @@ -31,38 +31,31 @@ CameraPos = record // By default fit 4x4 of logical space around (0;0) inside the window // Logical space Y is -2..+2, scale=2 self.scale_fine := 2; - // But if dw; - const max_point_scale = Settings.max_block_scale-Settings.block_w_pow; - begin - var point_scale := self.scale_pow + Settings.scale_shift + Floor(Log2(self.scale_fine/dh)); - var main_mipmap_lvl := -scale_shift; - if point_scale > max_point_scale then - begin - main_mipmap_lvl += point_scale-max_point_scale; - point_scale := max_point_scale; - end; - Result := System.ValueTuple.Create(point_scale, main_mipmap_lvl.Clamp(0,block_w)); - end; - private function GetBitCount := Settings.z_int_bits + -GetPointScaleAndMainMipMapLvl.Item1 + Settings.z_extra_precision_bits; + public function GetPointScalePow := + (self.scale_pow + Settings.scale_pow_shift + Floor(Log2(self.scale_fine/dh))) + .ClampTop(Settings.max_block_scale_pow - Settings.block_w_pow); + private function GetBitCount := Settings.z_int_bits + -GetPointScalePow + Settings.z_extra_precision_bits; private function GetWordCount := Ceil(GetBitCount/32).ClampBottom(1); // public function GetPosBitCount := Settings.z_int_bits - (self.scale_pow + Floor(Log2(self.scale_fine/dh))); // public function GetBlockBitCount := Settings.z_int_bits - (self.scale_pow + Settings.scale_shift - Settings.block_w_pow) + Settings.z_extra_precision_bits; - public procedure FixScalePow; + private procedure FixScalePow; begin if (scale_fine>=1) and (scale_fine<2) then exit; @@ -75,10 +68,33 @@ CameraPos = record raise new System.InvalidOperationException; end; + end; + + public procedure FixWordCount; + begin var pos_word_count := GetWordCount; if pos.Size = pos_word_count then exit; pos := pos.WithSize(pos_word_count); + end; + + public procedure Move(move_x,move_y, mouse_x,mouse_y, scale_speed: real; mouse_captured: boolean); + begin + var scale_mlt := 0.9 ** scale_speed; + begin + var shift_mlt := scale_fine * (1 - scale_mlt) * Ord(mouse_captured); + var dx := mouse_x-dw; + var dy := mouse_y-dh; + dx *= +shift_mlt/dw * AspectRatio; + dy *= -shift_mlt/dh; + var word_count := self.pos.Size; + var dr := new PointComponentShift(word_count, self.scale_pow, dx + move_x*scale_fine/dh); + var di := new PointComponentShift(word_count, self.scale_pow, dy + move_y*scale_fine/dh); + self.pos := self.pos.WithShiftClamp2(dr,di, true); + end; + + self.scale_fine *= scale_mlt; + FixScalePow; end; end; diff --git a/Samples/OpenGLABC/Mandelbrot/FieldTest.bmp b/Samples/OpenGLABC/Mandelbrot/FieldTest.bmp index 553d5008..dabc2687 100644 Binary files a/Samples/OpenGLABC/Mandelbrot/FieldTest.bmp and b/Samples/OpenGLABC/Mandelbrot/FieldTest.bmp differ diff --git a/Samples/OpenGLABC/Mandelbrot/FieldTest.pas b/Samples/OpenGLABC/Mandelbrot/FieldTest.pas index fd54b719..5ba0bad4 100644 --- a/Samples/OpenGLABC/Mandelbrot/FieldTest.pas +++ b/Samples/OpenGLABC/Mandelbrot/FieldTest.pas @@ -1,9 +1,11 @@ -## uses OpenCLABC, Settings, Blocks, CameraDef; +## uses OpenCLABC, Blocks, CameraDef, Settings; var max_iterations := 256; -//TODO Сейчас только степени двойки -// - Пожалуй частичный рендеринг тут тестить уже после того как напишу основной рендеринг +// Только степени двойки +// Если нет - округлит наверх до следующей степени двойки +// Потому что из GetRenderInfo берёт только .block_area, +// выбрасывая всю информацию о границах экрана var view_w := 1024; var view_h := 1024; @@ -12,23 +14,28 @@ more_output := false; {$endif ForceMaxDebug} -var render_info := BlockLayer.BlocksForCurrentScale(new CameraPos(view_w,view_h)); -var b_cy := render_info.blocks.GetLength(0); -var b_cx := render_info.blocks.GetLength(1); +CLMemoryObserver.Current := new TrackingMemoryObserver; -var render_block_size := Settings.block_w shr render_info.mipmap_lvl; +var camera := new CameraPos(view_w,view_h); +var layer := BlockLayer.GetLayer(camera); +var block_area: BlockLayerSubArea := layer.GetRenderInfo(camera, nil).block_area; +block_area.InitAllBlocks; +var blocks := block_area.MakeInitedBlocksMatr; +var b_cy := blocks.GetLength(0); +var b_cx := blocks.GetLength(1); + +var render_block_size := Settings.block_w; var render_sheet_w := b_cx * render_block_size; var render_sheet_h := b_cy * render_block_size; -var A_State := new CLArray(render_sheet_w*render_sheet_h); -var A_Steps := new CLArray(A_State.Length); +var A_Sheet := new CLArray(render_sheet_w*render_sheet_h); var V_UpdateCount := new CLValue(0); +var V_ExtractedCount := new CLValue(0); var A_Err := new CLArray(3); var Q_Init := CQNil - + A_State.MakeCCQ.ThenFillValue(0) - + A_Steps.MakeCCQ.ThenFillValue(0) - + A_Err.MakeCCQ.ThenFillValue(0) + + A_Sheet.MakeCCQ.ThenFillValue(0).DiscardResult + + A_Err.MakeCCQ.ThenFillValue(0).DiscardResult ; var Q_Steps := CQNil; @@ -36,19 +43,18 @@ for var b_y := 0 to b_cy-1 do for var b_x := 0 to b_cx-1 do begin - var b := render_info.blocks[b_y, b_x]; + var b := blocks[b_y, b_x]; // if (b_x, b_y) <> (1,1) then continue; + Q_Init += b.CQ_UpGradeToVRAM; + // Можно распараллелить, заменив += на *= (и в случае Q_Extract) // Но из тестирования - затраты на синхронизацию тут того не стоят //TODO Раз всё последовательно считает - можно по 1 блоку выделять и освобождать сразу... Q_Steps += b.CQ_MandelbrotBlockStep(max_iterations, V_UpdateCount, A_Err); var sheet_shift := render_block_size * (b_x + b_y*render_sheet_w); - Q_Extract += b.CQ_GetData(render_info.mipmap_lvl - , new ShiftedCLArray(A_State, sheet_shift, render_sheet_w) - , new ShiftedCLArray(A_Steps, sheet_shift, render_sheet_w) - ); + Q_Extract += b.CQ_GetData(new ShiftedCLArray(A_Sheet, sheet_shift, render_sheet_w), V_ExtractedCount); end; @@ -59,21 +65,32 @@ Q_Steps + Q_Extract + A_Err.MakeCCQ.ThenReadArray(err_data) + - A_Steps.MakeCCQ.ThenGetArray2(render_sheet_h, render_sheet_w) + V_UpdateCount.MakeCCQ.ThenGetValue.ThenUse(v->Println($'Updates: {v}')) + + V_ExtractedCount.MakeCCQ.ThenGetValue.ThenUse(v->Println($'Extracted: {v}')) + + A_Sheet.MakeCCQ.ThenGetArray2(render_sheet_h, render_sheet_w) ); if more_output then Println(sw.Elapsed); if err_data[0]<>0 then $'Err at [{err_data[1]},{err_data[2]}]: {CLCodeExecutionError(err_data[0])}'.Println; -$'Updates: {V_UpdateCount.GetValue}'.Println; +var sps := -Settings.scale_pow_shift; {$reference System.Drawing.dll} -var bmp := new System.Drawing.Bitmap(render_sheet_w, render_sheet_h); -for var y := 0 to render_sheet_h-1 do - for var x := 0 to render_sheet_w-1 do +var bmp := new System.Drawing.Bitmap(render_sheet_w shr sps, render_sheet_h shr sps); +for var y := 0 to bmp.Height-1 do + for var x := 0 to bmp.Width-1 do begin - var v := Round( (sheet_data[y,x]/max_iterations)**0.5 * 255); + var v_r := 0.0; + for var dy := 0 to 1 shl sps - 1 do + for var dx := 0 to 1 shl sps - 1 do + begin + var sy := y shl sps + dy; + var sx := x shl sps + dx; + v_r += ((sheet_data[sy,sx] and integer.MaxValue)/max_iterations) ** 0.5; + end; + v_r /= 1 shl (sps*2); + var v := Round(v_r * 255); bmp.SetPixel(x, y, System.Drawing.Color.FromArgb(v,v,v)); end; @@ -107,6 +124,14 @@ System.IO.File.Delete(res_fname); System.IO.File.Move(tmp_fname, res_fname); +$'Used a total of {CLMemoryObserver.Current.CurrentlyUsedAmount} VRAM bytes'.Println; +foreach var bl in blocks do + bl.Dispose; +A_Sheet.Dispose; +V_UpdateCount.Dispose; +V_ExtractedCount.Dispose; +A_Err.Dispose; + if more_output and ('[REDIRECTIOMODE]' not in System.Environment.CommandLine) then begin $'Press enter to exit'.Println; diff --git a/Samples/OpenGLABC/Mandelbrot/FieldTest.td b/Samples/OpenGLABC/Mandelbrot/FieldTest.td index b4f528df..dcc57cf8 100644 --- a/Samples/OpenGLABC/Mandelbrot/FieldTest.td +++ b/Samples/OpenGLABC/Mandelbrot/FieldTest.td @@ -1,24 +1,39 @@  -#ExpErr -Compile errors: -[225,27] Blocks.pas: The type 'PointComponent' does not contain a definition for 'FirstWordToReal' +#Delegates +Blocks.$delegate? = procedure(<>ml: MemoryLayer) +Blocks.$delegate? = procedure(<>ml: MemoryLayer) +Blocks.$delegate? = procedure(e: Exception) +Blocks.$delegate? = procedure(ml: MemoryLayer) +FieldTest.$delegate? = function: boolean +FieldTest.$delegate? = function: boolean +FieldTest.$delegate? = function: boolean +OpenCL.clCreateContextCallback = procedure(errinfo: string; private_info: System.IntPtr; cb: System.UIntPtr; user_data: System.IntPtr) +OpenCL.clEventCallback = procedure(event: cl_event; event_command_status: clCommandExecutionStatus; user_data: System.IntPtr) +OpenCL.clProgramCallback = procedure(program: cl_program; user_data: System.IntPtr) +OpenCLABC.$delegate? = function(ntv: cl_program; var data: clBool; validate: boolean): clErrorCode +OpenCLABC._GetPropValueFunc = function(ntv: cl_program; var data: T): clErrorCode +OpenCLABC_implementation______.EnqFunc = function(prev_res: T; cq: cl_command_queue; ev_l2: EventList): ValueTuple ()> +OpenCLABC_implementation______.InvokeParamsFunc = function(enq_c: integer; o_const: boolean; g: CLTaskGlobalData; enq_evs: DoubleList; par_err_handlers: DoubleList): ValueTuple> #ExpExecOtp Updates: 115036204 -[EventDebug]: 118 event's created -[QueueDebug]: 18 queue's created -[ExecDebug]: 48 cache entries created -[QueueResNil]: 334 -[QueueRes]: 119 - QueueResPtr: 1 - QueueResValDirect: 1 - QueueResValDirect: 1 - QueueResValDirect>: 16 - QueueResValDirect>: 1 - QueueResValDirect>: 1 - QueueResValDirect: 98 +Extracted: 4194304 +final = ErrHandlerInitial#0: ExecuteEnqFunc container +(total 1 handlers) +Used a total of 83886228 VRAM bytes +[EventDebug]: 56 event's created +[QueueDebug]: 1 queue's created +[ExecDebug]: 32 cache entries created +[QueueResNil]: 158 +[QueueRes]: 73 + QueueResPtr: 2 + QueueResValDirect: 1 + QueueResValDirect: 1 + QueueResValDirect>: 1 + QueueResValDirect>: 2 + QueueResValDirect: 66 #ReqModules OpenCLABC diff --git a/Samples/OpenGLABC/Mandelbrot/GL_CL_Context.pas b/Samples/OpenGLABC/Mandelbrot/GL_CL_Context.pas index 3472df59..41c782c3 100644 --- a/Samples/OpenGLABC/Mandelbrot/GL_CL_Context.pas +++ b/Samples/OpenGLABC/Mandelbrot/GL_CL_Context.pas @@ -39,7 +39,7 @@ procedure Init(hdc: gdi_device_context); var cl_c := cl.CreateContext(cl_c_props, cl_dvcs.Length, cl_dvcs, nil, IntPtr.Zero, ec); ec.RaiseIfError; - OpenCLABC.CLContext.Default := new CLContext(cl_c); + OpenCLABC.CLContext.Default := new CLContext(cl_c, false); end; procedure WrapBuffer(gl_b: gl_buffer; var cl_a: CLArray); where T: record; @@ -52,7 +52,7 @@ procedure WrapBuffer(gl_b: gl_buffer; var cl_a: CLArray); where T: record; if cl_a<>nil then cl_a.Dispose; - cl_a := new CLArray(cl_b); + cl_a := new CLArray(cl_b, false, false); end; end. \ No newline at end of file diff --git a/Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.cl b/Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.cl index e58034e6..74e97ba5 100644 --- a/Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.cl +++ b/Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.cl @@ -3,6 +3,8 @@ +//TODO Попробовать менять x и y местами, для более эффективного доступа к памяти + //TODO Minimize code dupe, using this: // - https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_C.html#declaring-and-using-a-block @@ -252,7 +254,7 @@ void point_component_sqr(point_component* x, uint* err) { } -// add bit at "bit_pos" to "x" multiple ("mlt"+0.5) times +// Add bit at "bit_pos" to "x" multiple ("mlt"+0.5) times // x += (1/2)^(bit_pos-Z_INT_BITS+1) * (mlt+0.5) void point_component_add_bit_mlt(point_component* x, uint bit_pos, uint mlt, uint* err) { err_cond(mlt>=BLOCK_W, CCEE_OVERFLOW, err); @@ -287,7 +289,7 @@ void point_component_add_bit_mlt(point_component* x, uint bit_pos, uint mlt, uin typedef struct { - // real and imaginary components of complex number + // Real and imaginary components of complex number point_component r, i; } point_pos; @@ -303,6 +305,7 @@ void point_sqr(point_pos* x, uint* err) { // = (x.r + x.i*i)*(x.r + x.i*i) = // = (sqr(x.r) - sqr(x.i)) + 2*(x.r*x.i)*i + // x_i_sqr = sqr(x.i) point_component x_i_sqr = x->i; point_component_sqr(&x_i_sqr, err); @@ -312,7 +315,7 @@ void point_sqr(point_pos* x, uint* err) { // x.r = sqr(x.r) point_component_sqr(&x->r, err); - // x.r -= sqr(x.i) + // x.r -= x_i_sqr x_i_sqr.words[0] ^= SIGN_BIT_MASK; point_component_add(&x->r, x_i_sqr, err); @@ -339,7 +342,7 @@ bool point_too_big(const point_pos x, uint* err) { // Should be >4 instead of >=4 // But we don't account here for bits in lower words // And exactly =4 is a point, impossible to see in an image - // Also checking for =4 would require adding +1 to Z_INT_BITS + // While checking for =4 would require adding +1 to Z_INT_BITS if (r2.words[0] >= Z_V_4) return true; return false; @@ -356,7 +359,6 @@ typedef struct { // block - current block of BLOCK_W*BLOCK_W points // pos00 - the "c" value of block[0,0] point // point_size_bit_pos - index of bit, which, when set alone, results in size of 1 point -// mipmap_need_update - will write 1's where mipmap would need to get recalculated // step_count - number of steps to try at once // update_count - will add 1 every time next z value is computed // err - will set !0 if there was an error during calculation @@ -364,7 +366,6 @@ kernel void MandelbrotBlockSteps( global point_info* block, constant point_pos* pos00, int point_size_bit_pos, - global uchar* mipmap_need_update, int step_count, global volatile uint* update_count, global uint* err @@ -399,159 +400,35 @@ kernel void MandelbrotBlockSteps( } if (!any_step_done) return; point->last_z = z; - - uint dx = x; - uint dy = y; - uint w = BLOCK_W; - for (uint w = BLOCK_W>>1; w > 0; w >>= 1) { - dx >>= 1; - dy >>= 1; - mipmap_need_update[dx + dy*w] = 1; - mipmap_need_update += w*w; - } - } -kernel void ExtractHighScaleSteps( - global point_info* block, uint scale_pow, - global uchar* result_state, uint result_state_shift, uint result_state_row_len, - global uint* result_steps, uint result_steps_shift, uint result_steps_row_len -) { - //TODO -} - -void ExtractSteps( - global uchar* source_state, uint source_state_item_len, uint source_state_row_len, - global uint* source_steps, uint source_steps_item_len, uint source_steps_row_len, - global uchar* result_state, uint result_state_item_len, uint result_state_row_len, - global uint* result_steps, uint result_steps_item_len, uint result_steps_row_len -) { - uint x = get_global_id(0); - uint y = get_global_id(1); - - uchar state = source_state[x*source_state_item_len + y*source_state_row_len]; - global uchar* p_result_state = &result_state[x*result_state_item_len + y*result_state_row_len]; - if (state < *p_result_state) return; - *p_result_state = state; - - uint steps = source_steps[x*source_steps_item_len + y*source_steps_row_len]; - //TODO Некоторые state не установлены, хотя должны быть - // - Пока что закомментировал проверку в FixMipMapLvl - но это плохо конечно... - // - Впрочем не похоже чтобы влияло на производительность - // - Может тогда нафиг мипмапы вообще? - // - С мипмапами или без, это слишком медленно... - // - Надо таки использовать только блоки текущего уровня, но вместо этого ещё брать данные из предыдущего кадра - steps |= SIGN_BIT_MASK * (uint)state; - global uint* p_result_steps = &result_steps[x*result_steps_item_len + y*result_steps_row_len]; - if (steps < *p_result_steps) return; - *p_result_steps = steps; - -} - -kernel void ExtractRawSteps( +kernel void ExtractSteps( global point_info* block, - global uchar* result_state, uint result_state_shift, uint result_state_row_len, - global uint* result_steps, uint result_steps_shift, uint result_steps_row_len -) { - - global uint* p_i_state = &block->state; - global uchar* p_state = (global uchar*)p_i_state; - global uint* p_steps = &block->steps; - - ExtractSteps( - p_state, sizeof(point_info)/sizeof(uchar), sizeof(point_info)/sizeof(uchar) * BLOCK_W, - p_steps, sizeof(point_info)/sizeof( uint), sizeof(point_info)/sizeof( uint) * BLOCK_W, - &result_state[result_state_shift], 1, result_state_row_len, - &result_steps[result_steps_shift], 1, result_steps_row_len - ); - -} - -kernel void ExtractMipMapSteps( - global uchar* mip_map_state, global uint* mip_map_steps, uint mip_map_shift, - global uchar* result_state, uint result_state_shift, uint result_state_row_len, - global uint* result_steps, uint result_steps_shift, uint result_steps_row_len, - uint mip_map_lvl -) { - - ExtractSteps( - &mip_map_state[mip_map_shift], 1, BLOCK_W>>mip_map_lvl, - &mip_map_steps[mip_map_shift], 1, BLOCK_W>>mip_map_lvl, - &result_state[result_state_shift], 1, result_state_row_len, - &result_steps[result_steps_shift], 1, result_steps_row_len - ); - -} - - - -void FixMipMapLvl( - global uchar* source_state, uint source_state_item_len, uint source_state_row_len, - global uint* source_steps, uint source_steps_item_len, uint source_steps_row_len, - global uchar* result_state, uint result_state_item_len, uint result_state_row_len, - global uint* result_steps, uint result_steps_item_len, uint result_steps_row_len, - global uchar* need_update, uint need_update_row_len + global uint* result_data, uint result_shift, uint result_row_len, + global volatile uint* extracted_count ) { uint x = get_global_id(0); uint y = get_global_id(1); - uchar* p_need_update = &need_update[x + y*need_update_row_len]; - if (!*p_need_update) return; - *p_need_update = false; + bool new_done = block[x + y*BLOCK_W].state != 0; + uint new_steps = block[x + y*BLOCK_W].steps & SIGN_BIT_ANTI_MASK; + uint new_data = (new_done ? SIGN_BIT_MASK : 0) ^ new_steps; - global uchar* p_source_state = &source_state[2*source_state_item_len*x + 2*source_state_row_len*y]; - global uchar* p_result_state = &result_state[1*result_state_item_len*x + 1*result_state_row_len*y]; - *p_result_state = p_source_state[0] & p_source_state[source_state_item_len] & p_source_state[source_state_row_len] & p_source_state[source_state_item_len+source_state_row_len]; + global uint* p_result_data = &result_data[result_shift + x + y*result_row_len]; + uint old_data = *p_result_data; + bool old_done = old_data>>31; + uint old_steps = old_data & SIGN_BIT_ANTI_MASK; - global uint* p_source_steps = &source_steps[2*source_steps_item_len*x + 2*source_steps_row_len*y]; - global uint* p_result_steps = &result_steps[1*result_steps_item_len*x + 1*result_steps_row_len*y]; - //*p_result_steps = max(max(p_source_steps[0], p_source_steps[source_steps_item_len]), max(p_source_steps[source_steps_row_len], p_source_steps[source_steps_item_len+source_steps_row_len])); - *p_result_steps = (p_source_steps[0] + p_source_steps[source_steps_item_len] + p_source_steps[source_steps_row_len] + p_source_steps[source_steps_item_len+source_steps_row_len] + 2) / 4; + // Попробовал быстрее избавится от артефактов на границе белой области, но pow значительно замедляет кадр + //old_steps *= pow(0.9, frames_since_sheet_transfer); -} - -kernel void FixFirstMipMap( - global point_info* block, - global uchar* mip_map_state, - global uint* mip_map_steps, - global uchar* need_update -) { - - global uint* p_i_state = &block->state; - global uchar* p_state = (global uchar*)p_i_state; - global uint* p_steps = &block->steps; - - FixMipMapLvl( - p_state, sizeof(point_info)/sizeof(uchar), sizeof(point_info)/sizeof(uchar) * BLOCK_W, - p_steps, sizeof(point_info)/sizeof( uint), sizeof(point_info)/sizeof( uint) * BLOCK_W, - mip_map_state, 1, BLOCK_W>>1, - mip_map_steps, 1, BLOCK_W>>1, - need_update, BLOCK_W>>1 - ); - -} - -kernel void FixMipMap( - global uchar* mip_map_state, - global uint* mip_map_steps, - global uchar* need_update, - uint next_mip_map_shift, - uint next_mip_map_lvl -) { - - uint prev_w = BLOCK_W >> (next_mip_map_lvl-1); - uint next_w = BLOCK_W >> next_mip_map_lvl; - uint prev_mip_map_shift = next_mip_map_shift - prev_w*prev_w; - - FixMipMapLvl( - &mip_map_state[prev_mip_map_shift], 1, prev_w, - &mip_map_steps[prev_mip_map_shift], 1, prev_w, - &mip_map_state[next_mip_map_shift], 1, next_w, - &mip_map_steps[next_mip_map_shift], 1, next_w, - &need_update[next_mip_map_shift], next_w - ); + //if (new_data > old_data) + if (new_done>=old_done || new_steps>=old_steps) { + atomic_inc(extracted_count); + *p_result_data = new_data; + } } diff --git a/Samples/OpenGLABC/Mandelbrot/MemoryLayering.pas b/Samples/OpenGLABC/Mandelbrot/MemoryLayering.pas new file mode 100644 index 00000000..fba03654 --- /dev/null +++ b/Samples/OpenGLABC/Mandelbrot/MemoryLayering.pas @@ -0,0 +1,195 @@ +unit MemoryLayering; + +{$savepcu false} //TODO + +type + MemoryLayerDataList = sealed class(IEnumerable) + private a := new T[16]; + private i1 := 0; + private i2 := 0; + {$ifdef DEBUG} + private list_v := int64(0); + {$endif DEBUG} + + private procedure IncInd(var ind: integer) := + ind := (ind+1) mod a.Length; + + public function IsEmpty := i1=i2; + + public function PeekOldest: T; + begin + {$ifdef DEBUG} + if IsEmpty then raise new System.InvalidOperationException; + {$endif DEBUG} + Result := a[i1]; + end; + public function RemoveOldest: T; + begin + {$ifdef DEBUG} + if IsEmpty then raise new System.InvalidOperationException; + list_v += 1; + {$endif DEBUG} + + Result := a[i1]; + a[i1] := default(T); + IncInd(i1); + + end; + + public procedure TryRemoveEachSingle(valid_remove_table: Dictionary; on_rem: T->()); + begin + {$ifdef DEBUG} + list_v += 1; + {$endif DEBUG} + + var look_i := i1; + var store_i := i1; + while look_i<>i2 do + begin + var o := a[look_i]; + + var valid_remove: boolean; + var need_remove := valid_remove_table.TryGetValue(o, valid_remove); + {$ifdef DEBUG} + if need_remove and not valid_remove then + $'Could not remove from memory: {o}'.Println; +// raise new System.InvalidOperationException; + {$endif DEBUG} + valid_remove_table[o] := false; + + if need_remove then + on_rem(o) else + begin + if look_i<>store_i then + a[store_i] := o; + IncInd(store_i); + end; + + IncInd(look_i); + end; + + self.i2 := store_i; + while look_i<>store_i do + begin + a[store_i] := default(T); + IncInd(store_i); + end; + + end; + + public procedure AddNewest(o: T); + begin + {$ifdef DEBUG} + if o in self then + raise new System.InvalidOperationException; + list_v += 1; + {$endif DEBUG} + a[i2] := o; + i2 := (i2+1) mod a.Length; + if i1<>i2 then exit; + var n_a := new T[a.Length*2]; + for var d := 0 to a.Length-1 - i1 do + n_a[d] := a[i1+d]; + for var d := 0 to i1-1 do + n_a[a.Length-i1+d] := a[d]; + i1 := 0; + i2 := a.Length; + a := n_a; + end; + + private function Enmr: sequence of T; + begin + var enmr_i := self.i1; + {$ifdef DEBUG} + var org_list_v := list_v; + {$endif DEBUG} + while enmr_i<>i2 do + begin + {$ifdef DEBUG} + if org_list_v <> self.list_v then + raise new System.InvalidOperationException; + {$endif DEBUG} + yield a[enmr_i]; + IncInd(enmr_i); + end; + {$ifdef DEBUG} + while enmr_i<>i1 do + begin + if a[enmr_i] <> default(T) then + raise new System.InvalidOperationException; + IncInd(enmr_i); + end; + {$endif DEBUG} + end; + public function GetEnumerator := Enmr.GetEnumerator; + public function System.Collections.IEnumerable.GetEnumerator: System.Collections.IEnumerator := GetEnumerator; + + end; + + IMemoryLayerData = interface + function GetByteSize: int64; + end; + MemoryLayer = abstract class + where TData: IMemoryLayerData; + private l := new MemoryLayerDataList; + private newly_added := new List; + + private _name: string; + + private allowed_size: int64; + private filled_size := int64(0); + protected function GetRealFilledSize: int64; abstract; + + public constructor(name: string; allowed_size: int64); + begin + self._name := name; + self.allowed_size := allowed_size; + end; + private constructor := raise new System.InvalidOperationException; + + public property Name: string read _name; + public function Enmr := l.Enmr; + + public procedure BeginUpdate(item_new_table: Dictionary); + begin + self.filled_size := GetRealFilledSize; + l.TryRemoveEachSingle(item_new_table, data->(self.filled_size -= data.GetByteSize)); + end; + public procedure EndUpdate; + begin + // First added is the most important + // Store it last it's the newest in l + for var i := newly_added.Count-1 downto 0 do + l.AddNewest(newly_added[i]); + newly_added.Clear; + end; + + public function TryAdd(new_data: TData; on_displaced: TData->()): boolean; + begin + Result := false; + var new_data_sz := new_data.GetByteSize; + var max_filled_size := allowed_size-new_data_sz; + + while filled_size > max_filled_size do + begin + if l.IsEmpty then + begin + on_displaced(new_data); + exit; + end; + var old_data := l.PeekOldest; // First only peek, in case on_displaced throws + var old_data_size := old_data.GetByteSize; + on_displaced(old_data); + if old_data <> l.RemoveOldest then + raise new System.InvalidOperationException; + filled_size -= old_data_size; + end; + + newly_added += new_data; + filled_size += new_data_sz; + Result := true; + end; + + end; + +end. \ No newline at end of file diff --git a/Samples/OpenGLABC/Mandelbrot/MemoryLayering.td b/Samples/OpenGLABC/Mandelbrot/MemoryLayering.td new file mode 100644 index 00000000..ab9afd1f --- /dev/null +++ b/Samples/OpenGLABC/Mandelbrot/MemoryLayering.td @@ -0,0 +1,6 @@ + + + +#SkipTest + + diff --git a/Samples/OpenGLABC/Mandelbrot/PointComponents.pas b/Samples/OpenGLABC/Mandelbrot/PointComponents.pas index 951f7ec2..b0ecfbce 100644 --- a/Samples/OpenGLABC/Mandelbrot/PointComponents.pas +++ b/Samples/OpenGLABC/Mandelbrot/PointComponents.pas @@ -6,6 +6,104 @@ type + PointComponentShift = record + private bits: int64; + private word_ind := 0; + private overflow := false; + + public constructor(word_c, pow: integer; shift: real); + // 64 bits of int64 + // -= 1 bit for sign + // -= 1 bit because shift is normalized to [1;2), the "1.xxx" form, with 1 leading int bit + const max_local_pow = 62; + begin + var sign := Sign(shift); + shift *= sign; + self.bits := sign; + if sign=0 then exit; + + if (shift<1) or (shift>=2) then + begin + var extra_pow_r := Log2(shift); + var extra_pow_i := Floor(extra_pow_r); + pow += extra_pow_i; + shift := 2 ** (extra_pow_r-extra_pow_i); + {$ifdef DEBUG} + if (shift<1) or (shift>=2) then + raise new System.InvalidOperationException; + {$endif DEBUG} + end; + + // With pow=0,shift=1 this is the only bit to set + var head_bit_ind := Settings.z_int_bits-1 - pow; + // In ideal case this is on the word bound or within 12 more bits, + // allowing to save all 52 bits of mantissa of "real" + // But if not, we round word_ind down (div) to prev word + // And then rounding to int64, loosing some precision (max 19/52 bits) + var after_tail_bit_ind := head_bit_ind + max_local_pow; + + var word_ind_raw := after_tail_bit_ind div 32; + var word_ind := word_ind_raw.Clamp(0, word_c-1); + self.word_ind := word_ind; + + // Bit shift by max_local_pow if there was no rounding (div) + // If there was, bit shift a bit less + shift *= 2 ** (max_local_pow-1 - (after_tail_bit_ind - (word_ind+1)*32)); + + {$ifdef DEBUG} + //TODO Слишном сонный чтобы доразобраться + // - Вылетает если pow>32 (очень далеко отдалить) +// if (after_tail_bit_ind>=0) and (shift>int64.MaxValue) then +// raise new System.InvalidOperationException; + if (after_tail_bit_indinteger.MaxValue) downto 0 do + begin + sb.Append(carry shr bit_i and 1); + if bit_i=32 then + sb += '|'; + end; + end; + Result := sb.ToString; + end; + + end; + PointComponent = record(System.IEquatable) private _words: array of cardinal; private const sign_bit_mask: cardinal = 1 shl 31; @@ -26,7 +124,7 @@ PointComponent = record(System.IEquatable) var size := c1.Words.Length; {$ifdef DEBUG} if size <> c2.Words.Length then - raise new System.InvalidOperationException; + raise new System.InvalidOperationException($'{c1} vs {c2}'); {$endif DEBUG} Result := false; for var i := size-1 downto 0 do @@ -46,7 +144,7 @@ PointComponent = record(System.IEquatable) Result := Words[l-1]; if l<>1 then Result := Result xor Words[l-2]; - // GetHashCode должно считать быстро + // GetHashCode должно считаться быстро // 64 нижних бита уже достаточно end; @@ -77,6 +175,25 @@ PointComponent = record(System.IEquatable) Result := sb.ToString; end; + public procedure Save(bw: System.IO.BinaryWriter) := + foreach var x in Words do bw.Write(x); + public static function Load(br: System.IO.BinaryReader; word_count: integer): PointComponent; + begin + Result := new PointComponent(word_count); + for var i := 0 to word_count-1 do + Result.WOrds[i] := br.ReadUInt32; + end; + + public function FirstWordToReal: real; + const body_d = 1 shl (32-z_int_bits); + const body_m = 1/body_d; + begin + Result := Words[0] and sign_bit_anti_mask; + Result *= body_m; + if Words[0] and sign_bit_mask <> 0 then + Result := -Result; + end; + public function WithSize(size: integer): PointComponent; begin {$ifdef DEBUG} @@ -88,50 +205,70 @@ PointComponent = record(System.IEquatable) Result.Words[i] := self.Words[i]; end; - public static function RoundToLowestBits(size, bit_ind: integer; x: real): int64; + private procedure HandleMinusZero(expect_minus_zero: boolean); begin - var d_shift := size*32-1-bit_ind; - {$ifdef DEBUG} - if d_shift>=63 then - raise new System.OverflowException; - {$endif DEBUG} - x *= int64(1) shl d_shift; - Result := Convert.ToInt64(x); + if Words[0] <> sign_bit_mask then exit; + for var i := 1 to Words.Length-1 do + if Words[i]<>0 then exit; + if not expect_minus_zero then + raise new System.InvalidOperationException; + Words[0] += sign_bit_mask; end; - public function AddLowestBits(d: int64): PointComponent; + + public function WithShiftClamp2(shift: PointComponentShift; expect_minus_zero: boolean): PointComponent; + const v2: cardinal = 1 shl (33-Settings.z_int_bits); begin + Result := self; + if shift.bits=0 then exit; var size := self.Words.Length; + Result := new PointComponent(size); + + if shift.overflow then + begin + Result.Words[0] := v2; + if shift.bits<0 then + Result.Words[0] += sign_bit_mask; + // All other words are already 0 on init + exit; + end; + + var d := shift.bits; var self_sign := self.Words[0] and sign_bit_mask; var same_sign := (self_sign<>0) = (d<0); if self_sign<>0 then d := -d; - Result := new PointComponent(size); + for var i := size-1 downto shift.word_ind+1 do + Result.Words[i] := self.Words[i]; + var carry := d; - for var i := size-1 downto 1 do + for var i := shift.word_ind downto 1 do begin carry += self.Words[i]; - {$ifdef DEBUG} - if (d<0) <> (carry<0) then - raise new System.OverflowException; - {$endif DEBUG} Result.Words[i] := carry; carry := carry shr 32; end; begin carry += self.Words[0] and sign_bit_anti_mask; - {$ifdef DEBUG} - if (d<0) <> (carry<0) then - raise new System.OverflowException; - {$endif DEBUG} Result.Words[0] := carry xor self_sign; end; + if Abs(carry) shr (33-Settings.z_int_bits) <> 0 then + begin + var res_sign := self_sign; + if not same_sign then + res_sign := sign_bit_mask - res_sign; + Result.Words[0] := v2 xor res_sign; + for var i := 1 to size-1 do + Result.Words[i] := 0; + exit; + end; + if self_sign <> (Result.Words[0] and sign_bit_mask) then begin - if same_sign or (Abs(carry) shr 32 <> 0) then - raise new System.OverflowException; + if same_sign then + raise new System.InvalidOperationException; var compliment := true; for var i := size-1 downto 1 do begin @@ -141,9 +278,10 @@ PointComponent = record(System.IEquatable) Result.Words[0] := sign_bit_mask xor not Result.Words[0] + Ord(compliment); end; + self.HandleMinusZero(expect_minus_zero); end; - // Round to the block bound + // Round self to the block boundry // PointComponent with only bit at block_sz_bit_ind set is the block side length // - Self is set to the bound between blocks // - block_len_outside is set to value in [0;1), indicating how much of the block was outside of initial bounds @@ -217,69 +355,102 @@ PointComponent = record(System.IEquatable) if Words[0] and sign_bit_mask <> self_sign then raise new System.OverflowException; + HandleMinusZero(round_up); + end; - - public procedure SelfFlipIfMinusZero; + public function WithBlockRound(block_sz_bit_ind: integer; round_up: boolean): PointComponent; begin - if Words[0] <> sign_bit_mask then exit; - for var i := 1 to Words.Length-1 do - if Words[i]<>0 then exit; - Words[0] += sign_bit_mask; + Result._words := self.Words.ToArray; + var block_len_outside: single; + Result.SelfBlockRound(block_sz_bit_ind, round_up, block_len_outside); end; - private function BodyWordAs64At(ind: integer): int64 := - if ind=0 then Words[0] and sign_bit_anti_mask else Words[ind]; + private function BodyWordAs64AtOr0(ind: integer): int64 := + if ind=0 then Words[0] and sign_bit_anti_mask else + if ind c2.Words.Length then - raise new System.InvalidOperationException; - {$endif DEBUG} - var word_inner_pos: integer; // 0..31 var word_ind := System.Math.DivRem(block_sz_bit_ind, 32, word_inner_pos); + var lower_bits_mask: int64 := (1 shl (31-word_inner_pos))-1; var c1_sign := c1.Words[0] and sign_bit_mask; var c2_sign := c2.Words[0] and sign_bit_mask; - {$ifdef DEBUG} - if (c1_sign=0) and (c2_sign<>0) then - raise new System.InvalidOperationException; - {$endif DEBUG} var same_sign := c1_sign=c2_sign; - var lower_bits_mask: cardinal := (1 shl (31-word_inner_pos))-1; + for var i := 0 to word_ind-2 do + begin + var word1 := c1.BodyWordAs64AtOr0(i); + var word2 := c2.BodyWordAs64AtOr0(i); + if word1<>word2 then + raise new System.OverflowException; + if not same_sign and (word1<>0) then + raise new System.OverflowException; + end; + + {$ifdef DEBUG} + for var i := word_ind+1 to c1.Words.Length-1 do + if c1.Words[i] <> 0 then raise new System.InvalidOperationException; + for var i := word_ind+1 to c2.Words.Length-1 do + if c2.Words[i] <> 0 then raise new System.InvalidOperationException; + {$endif DEBUG} - var diff := int64(0); + var total_diff := int64(0); if word_ind<>0 then begin - var prev_word1 := c1.BodyWordAs64At(word_ind-1); - var prev_word2 := c2.BodyWordAs64At(word_ind-1); - diff := lower_bits_mask and if same_sign then prev_word2-prev_word1 else prev_word2+prev_word1; + var word1 := c1.BodyWordAs64AtOr0(word_ind-1); + var word2 := c2.BodyWordAs64AtOr0(word_ind-1); + + var diff := if same_sign then word2-word1 else word2+word1; + var diff_sign := Sign(diff); + diff *= diff_sign; + diff := diff and lower_bits_mask; diff := diff shl (1+word_inner_pos); + diff *= diff_sign; + total_diff += diff; + + word1 := word1 and not lower_bits_mask; + word2 := word2 and not lower_bits_mask; + if word1<>word2 then + raise new System.OverflowException; + if not same_sign and (word1<>0) then + raise new System.OverflowException; + + end; + begin + var word1 := c1.BodyWordAs64AtOr0(word_ind); + var word2 := c2.BodyWordAs64AtOr0(word_ind); + + var diff := if same_sign then word2-word1 else word2+word1; + var diff_sign := Sign(diff); + diff *= diff_sign; + diff := diff shr (31-word_inner_pos); + diff *= diff_sign; + total_diff += diff; + {$ifdef DEBUG} - if (prev_word1 and not lower_bits_mask) <> (prev_word2 and not lower_bits_mask) then - raise new System.InvalidOperationException; + if word1 and lower_bits_mask <> 0 then raise new System.InvalidOperationException; + if word2 and lower_bits_mask <> 0 then raise new System.InvalidOperationException; {$endif DEBUG} + end; - var curr_word1 := c1.BodyWordAs64At(word_ind); - var curr_word2 := c2.BodyWordAs64At(word_ind); - {$ifdef DEBUG} - if curr_word1 and lower_bits_mask <> 0 then raise new System.InvalidOperationException; - if curr_word2 and lower_bits_mask <> 0 then raise new System.InvalidOperationException; - {$endif DEBUG} - diff += (if same_sign then curr_word2-curr_word1 else curr_word2+curr_word1) shr (31-word_inner_pos); - Result := diff; + total_diff := if c2_sign=0 then +total_diff else -total_diff; + Result := total_diff; {$ifdef DEBUG} - if int64(Result) <> diff then + if Result <> total_diff then + begin +// BlocksCount(c1, c2, block_sz_bit_ind); + //TODO "diff := lower_bits_mask and ..." даёт неправильное значение для отрицательных diff + // - shr/shl тоже не расчитаны на этот случай + // - Сейчас падает если приблизить на scale_pow=-30 + // - Тестил на мини-мандельбротах в левой, легко-просчитываемой части raise new System.OverflowException; - for var i := 0 to word_ind-2 do - if c1.BodyWordAs64At(i) <> c2.BodyWordAs64At(i) then - raise new System.OverflowException; + end; {$endif DEBUG} end; @@ -290,7 +461,6 @@ PointComponent = record(System.IEquatable) var word_inner_pos: integer; // 0..31 var word_ind := System.Math.DivRem(block_sz_bit_ind, 32, word_inner_pos); - var curr_block_sz_bit: cardinal := 1 shl (31-word_inner_pos); {$ifdef DEBUG} @@ -302,19 +472,19 @@ PointComponent = record(System.IEquatable) {$endif DEBUG} var self_sign := Words[0] and sign_bit_mask; + var i := word_ind; + if self_sign=0 then begin - var i := word_ind; Result.Words[i] := self.Words[i] + curr_block_sz_bit; while (Result.Words[i]=0) and (i<>0) do begin i -= 1; - Result.Words[word_ind] := self.Words[word_ind] + 1; + Result.Words[i] := self.Words[i] + 1; end; end else // Substract instead of adding begin - var i := word_ind; Result.Words[i] := self.Words[i] - curr_block_sz_bit; while (Result.Words[i]>self.Words[i]) and (i<>0) do begin @@ -323,28 +493,72 @@ PointComponent = record(System.IEquatable) end; end; + while i<>0 do + begin + i -= 1; + Result.Words[i] := self.Words[i]; + end; + if Result.Words[0] and sign_bit_mask <> self_sign then raise new System.OverflowException; - Result.SelfFlipIfMinusZero; + Result.HandleMinusZero(true); + end; + + // [c1,c2) range of points + public static function Range(c1, c2: PointComponent; block_sz_bit_ind: integer): array of PointComponent; + begin + Result := new PointComponent[BlocksCount(c1,c2,block_sz_bit_ind)]; + var x := c1; + Result[0] := x; + for var i := 1 to Result.Length-1 do + begin + x := x.MakeNextBlockBound(block_sz_bit_ind); + Result[i] := x; + end; + {$ifdef DEBUG} + if x.MakeNextBlockBound(block_sz_bit_ind)<>c2 then + raise new System.InvalidOperationException; + {$endif DEBUG} end; end; - PointPos = record + PointPos = record(System.IEquatable) + // Real and imaginary components of complex number public r,i: PointComponent; public constructor(r,i: PointComponent) := (self.r,self.i) := (r,i); + public static function operator=(p1, p2: PointPos) := (p1.r=p2.r) and (p1.i=p2.i); + public static function operator<>(p1, p2: PointPos) := not(p1=p2); + public function Equals(other: PointPos) := self=other; + public function Equals(o: object): boolean; override := + (o is PointPos(var other)) and Equals(other); + public function GetHashCode: integer; override := + r.GetHashCode xor i.GetHashCode*668265263; + + public function ToString: string; override := $'({r}; {i})'; + public property Size: integer read r.Words.Length; public function WithSize(size: integer) := new PointPos( r.WithSize(size), i.WithSize(size) ); - public function AddLowestBits(dr,di: int64) := new PointPos( - r.AddLowestBits(dr), - i.AddLowestBits(di) + public procedure Save(bw: System.IO.BinaryWriter); + begin + r.Save(bw); + i.Save(bw); + end; + public static function Load(br: System.IO.BinaryReader; word_count: integer) := new PointPos( + PointComponent.Load(br, word_count), + PointComponent.Load(br, word_count) + ); + + public function WithShiftClamp2(dr, di: PointComponentShift; expect_minus_zero: boolean) := new PointPos( + self.r.WithShiftClamp2(dr, expect_minus_zero), + self.i.WithShiftClamp2(di, expect_minus_zero) ); public procedure SelfBlockRound(block_sz_bit_ind: integer; round_up: boolean; var skip_r: single; var skip_i: single); @@ -353,12 +567,6 @@ PointPos = record self.i.SelfBlockRound(block_sz_bit_ind, round_up, skip_i); end; - public procedure SelfFlipIfMinusZero; - begin - self.r.SelfFlipIfMinusZero; - self.i.SelfFlipIfMinusZero; - end; - end; end. \ No newline at end of file diff --git a/Samples/OpenGLABC/Mandelbrot/Settings.pas b/Samples/OpenGLABC/Mandelbrot/Settings.pas index c1d1319c..0ec4154a 100644 --- a/Samples/OpenGLABC/Mandelbrot/Settings.pas +++ b/Samples/OpenGLABC/Mandelbrot/Settings.pas @@ -2,71 +2,92 @@ {$savepcu false} //TODO -// Всё пространство разделено на блоки этой ширины и высоты -// Если камеру приближают - просчитываются 4 более блока меньшего масштаба: -// Cтолько же точек, но эти точки упакованы ближе друг к другу -// Чем больше блоки - тем быстрее их считает и тем больше использует памяти +// Я разделяю в этой программе 3 вида пространства: +// - Графическое: 1 пиксель экрана это 1х1 графического пространства +// - Логическое: Весь рисунок Mandelbrot помещается в -2..+2 логического пространства (горизонтально и вертикально), независимо от графического масштаба +// - Пространство блока/точек: Каждый блок любого масштаба хранит block_w*block_w точек, каждая из которых просчитывается отдельно + +// Если в блоках слишком мало точек, получается много отдельных обращений к GPU +// А если слишном много - значительная часть просчитываемых точек +// (для которых надо VRAM и время GPU) может быть за пределами экрана const block_w_pow = 9; const block_w = 1 shl block_w_pow; // 512 -// Одновременно может просчитываться max_GPU_inst блоков -// Каждый из видимых блоков нужного масштаба будет по очереди -// пытаться сделать ещё 1 шаг рекуррентной функции "z_next(z) = z*z + c" для каждой точки -// Где "z" это предыдущее значение (изначально 0), а "c" это координата точки -const max_GPU_inst = 5; -// Сколько надо прибавить к масштабу камеры чтобы получить текущий масштаб просчитываемых блоков -// Масштаб это степень двойки -// -1 значит на каждый пиксель экрана придётся хотя бы 4 (2х2) просчитанные точки -const scale_shift = -1; // <=0 -// Блок максимального размера это 2х2 -// Таким образом вся просчитываемая область разбита на минимум 4 блока, по 1 на угол -// При приближении камеры каждый из этих угловых блоков будет далее разбивать на 4 меньших блока -const max_block_scale = 1; // Не менять - эта константа много где неявно задана -// Отрисованные блоки хранит в VRAM (памяти GPU) -const max_VRAM = 1610612736; // 1.5 GB +// Просчитываемые в любой момент блоки хранит в VRAM (личной памяти GPU) +// Для 1080p надо чуть больше 1.25 GB при поверхностном приближении +// (чем глубже - тем больше точность расчётов и тем больше надо памяти) +// Но если VRAM у системы мало (к примеру, браузер много скушал), +// драйвера могут крашнутся при слишком большом использовании VRAM данной программой +// Смотрите в диспетчер задач чтобы выставить сколько вы можете себе позволить +const max_VRAM = 838860800; // 800 MB // Если VRAM заканчивается - старые блоки отправляет в RAM (в обычную оперативную память) -const max_RAM = 4294967296; // 4 GB -// Если в RAM есть место, но отрисованный блок был там долго - тоже удалить -// (иначе его отправит в файл подкачки, а потом, при подгрузке назад - будет лагать) -const RAM_life_span_seconds = 15*60; // 15 минут +// Пока что 0, что значит старые блоки в VRAM уничтожаются как только перестают помещаться +const max_RAM = 0;// 4294967296; // 4 GB +// А если заканчивается и RAM - самые старые блоки отправляет на диск +// Если у вас медленный диск, есть смысл оставить тут 0 +// Но если включать - надо сразу установить max_RAM>=max_VRAM +// Потому что блоки не могут прыгнуть сразу из VRAM на диск или назад +const max_drive_space = 0;// 10737418240; // 10 GB + +// Сторона блока в логическом пространстве это всегда степень двойки (обычно отрицательная степень) +// Если scale_pow_shift=0, размер блоков в логическом пространстве будет выбирать так, +// чтобы на каждый пиксель экрана приходилась хотя бы область из 1x1 точек в соответствующем блоке +// Когда отдаляют достаточно чтобы в пиксель поместилась область из 2x2 точек, масштаб блоков переключается на следующий +// Если scale_pow_shift=-1, на каждый пиксель придётся хотя бы 2x2 точек, и всегда меньше 4x4 +// Значение -1 значительно улучшает визуал, но уже требует в 4 раза больше памяти и производительности +// Желательно дальше не понижать без очень мощного компьютера +// Иначе просчитывать будет только блоки в центре экрана (если закончится VRAM) +const scale_pow_shift = -1; // <=0 +// Блок максимального размера занимает в логическом пространстве область 2х2 +// Таким образом вся просчитываемая область (4х4 вокруг точки 0;0) разбита на минимум 4 блока, по 1 на угол +// При приближении камеры каждый из этих угловых блоков будет далее разбивать на 4 меньших блока +// Не менять, эта константа много где неявно задана, самим алгоритмом +const max_block_scale_pow = 1; -// Чтобы для каждого кадра не пересчитывать -// среднее арифметическое всех точек под каждым пикселем, -// Каждому блоку создаст мипмапы таких размеров: -// -// (block_w/2) * (block_w/2) -// (block_w/4) * (block_w/4) -// ... -// 2 * 2 -// 1 * 1 -const mipmap_total_size = block_w*block_w div 3; +// Каждая точка в каждом блоке хранит текущее состояние +// (z и номер шага) рекуррентной функции "z_next(z) = z*z + c" +// Где "z" это предыдущее значение (изначально 0), а "c" это "x+y*i" +// (комплексное число представляющее координаты точки в логическом пространстве) +// На экране рисует только кол-во шагов этой функции (ещё_считает=белый; 0 шагов=чёрный; остальное радугой) +// А в отдельном от графики потоке выполнения (класс Blocks.BlockUpdater) +// циклически берёт блоки текущего масштаба и просчитывает их на несколько шагов вперёд +// Максимум за 1 итерацию обработки может выполнить max_steps_at_once шагов для всех точек +// Это ограничение на всякий случай, чтобы Blocks.BlockUpdater не застряло на несколько секунд +// (от такого застревания может даже полностью перезапустится графический драйвер) +// Но это и так не должно происходить, потому что количество шагов и так сбрасывается к 1 при смене текущих блоков +const max_steps_at_once = 1024; +// Все блоки можно обрабатывать по-очереди или параллельно +// Чтобы достичь большОй параллельности - надо выделить много ресурсов в +// виде OpenCL.cl_command_queue (личные данные каждого потока выполнения GPU) +// Но если обрабатывать все блоки один за другим, GPU будет повторно оставаться +// без работы на короткие промежутки времени, между обработками блоков +// Поэтому одновременно обрабатывать будет не больше чем max_parallel_blocks блоков +const max_parallel_blocks = 2; +// Чем больше шагов - тем меньшая доля времени будет потрачена на синхронизацию и т.п. между обработками +// Но если обработка блоков использует GPU слишком эффективно, +// она заберёт все ресурсы у системы и графика начнёт лагать +// Поэтому кол-во шагов в 1 обработке меняет в диапазоне 1..max_steps_at_once, +// так чтобы 1 обработка заменяла примерно столько секунд: +const target_step_time_seconds = 0.050; +//TODO Может делать несколько запусков kernel-а подряд, чтобы у потока графики было больше шансов получить нужное ему время GPU... -// Последним шагом считаеться шаг где |z_next|>2 -// На шаге=0 z=0, а значит для точек |c|>2 на шаг 1 уже не переключится -// На шаге>0 |z|<=2, иначе на этот шаг не переключилось бы -// Тогда |z*z+c| максимум будет 2*2+2=6 +// Комплексное число z на каждом шаге представляет как 2 компоненты (PointComponents.pas и point_component в MandelbrotSampling.cl) +// Каждая компонента это число число с фиксированной точкой (fixed-point number) +// Это всё число разделено на несколько слов типа UInt32 (чтобы обрабатывать сразу кучу битов каждой операцией процессора) +// Кол-во слов выбирается в CameraDef.CameraPos.GetWordCount, исходя из нужного кол-ва бит для текущего масштаба +// +// Последним шагом считается шаг где |z_next|>2 +// Модуль комплексной точки "x+y*i" это модуль (длина) вектора (x;y), то есть "Sqrt(x.Sqr+y.Sqr)" +// На шаге=0: z=0, а значит для точек |c|>2 на шаг 1 уже не переключится +// На шаге>0: |z|<=2, иначе на этот шаг не переключилось бы +// Тогда |z*z+c| будет максимум 2*2+2=6 // Таким образом чтобы представить целую часть результата вычисления надо максимум 3 бита -// И затем ещё +1 бит для знака (+ или -) +// И затем ещё +1 бит в самом начале для знака (+ или -) const z_int_bits = 4; // 0..31 -// А точность (кол-во бит) после точки будет масштаб+block_w_pow+z_extra_precision_bits -// 16 доп. бит значит что для заметной ошибки надо минимум столько операций: +// А точность (кол-во бит) после точки будет -масштаб_точки+z_extra_precision_bits +// 16 дополнительных бит значит что для заметной ошибки надо минимум столько операций: // LogN(1+2**-16, 1.5) ~= 26572 +// Каждый шаг выполняет несколько операций, но +16 битов это всё равно очень много const z_extra_precision_bits = 16; // >=0 -// Таким образом кол-во байт на 1 точку: -// > 4 (состояние) + 4 (кол-во шагов) + z_component_size * 2 (действительная и мнимая компоненты "z") -// Где компонента z будет записью из нескольких cardinal (32-битных беззнаковых целых) -// > z_component_size = Trunc( (z_int_bits + -point_scale + z_extra_precision_bits) / 32 )*4 -// Значит объём VRAM на 1 блок: -// > VRAM = (block_w*block_w) * (8 + 2*z_component_size) -// > VRAM/(block_w*block_w) = 8 + 2*z_component_size -// > VRAM/(block_w*block_w) - 8 = 2*z_component_size -// > VRAM/(block_w*block_w)/2 - 4 = z_component_size -const max_z_component_size = max_VRAM/(block_w*block_w)/2 - 4; // 3068 -// Минимальный уровень масштаба, на котором максимум можно нарисовать 1 блок -// На самом деле, если у вас на экране одновременно показывает "n" блоков -// То максимальный масштаб без глюков будет max_z_scale_bits/n -// Но до этого масштаба всё равно будет сложно до-скролить -const max_z_scale_bits_raw = max_z_component_size/4*32 - z_int_bits - block_w_pow - z_extra_precision_bits; // 24515 -const max_z_scale_bits_rounded = Trunc(max_z_component_size/4)*32 - z_int_bits - block_w_pow - z_extra_precision_bits; // 24515 end. \ No newline at end of file diff --git a/Samples/OpenGLABC/Mandelbrot/Shaders/Box.geom b/Samples/OpenGLABC/Mandelbrot/Shaders/Box.geom index b6c61c0c..0d46320d 100644 --- a/Samples/OpenGLABC/Mandelbrot/Shaders/Box.geom +++ b/Samples/OpenGLABC/Mandelbrot/Shaders/Box.geom @@ -5,10 +5,10 @@ uniform float view_skip_x_last, view_skip_y_last; layout(points) in; layout(triangle_strip, max_vertices = 4) out; -out vec2 logic_pos; +out vec2 screen_pos; void SendVertex(float coord1, float dx, float coord2, float dy) { - logic_pos = vec2(coord1, coord2); + screen_pos = vec2(coord1, coord2); gl_Position = vec4(coord1+dx, coord2+dy, 0, 1); EmitVertex(); } diff --git a/Samples/OpenGLABC/Mandelbrot/Shaders/Rainbow.frag b/Samples/OpenGLABC/Mandelbrot/Shaders/Rainbow.frag index fbd2b1b7..e8964099 100644 --- a/Samples/OpenGLABC/Mandelbrot/Shaders/Rainbow.frag +++ b/Samples/OpenGLABC/Mandelbrot/Shaders/Rainbow.frag @@ -1,6 +1,6 @@ #version 460 core -noperspective in vec2 logic_pos; +noperspective in vec2 screen_pos; uniform float sheet_skip_x_frst, sheet_skip_y_frst; uniform float sheet_skip_x_last, sheet_skip_y_last; @@ -17,17 +17,19 @@ layout(binding = 1) buffer temp_otp { } temp; /**/ -out vec3 color; -// All components are in the range [0…1], including hue. -vec3 hsv2rgb(float hue) + +// All 3 components are in the [0…1) range +vec3 hsv2rgb(float hue, float saturation, float value) { - vec3 c = {hue, 0.75, 0.5}; + vec3 c = {hue, saturation, value}; vec4 K = vec4(1.0, 2.0 / 3.0, 1.0 / 3.0, 3.0); vec3 p = abs(fract(c.xxx + K.xyz) * 6.0 - K.www); return c.z * mix(K.xxx, clamp(p - K.xxx, 0.0, 1.0), c.y); } + + struct point_state { bool done; uint steps; @@ -42,11 +44,16 @@ point_state PointAt(ivec2 ind) { ); } + + +out vec3 color_out; + const int n_pts_avg = 5; +const int avg_pts_count = n_pts_avg*n_pts_avg; void main() { - // logic_pos is -1..+1 - vec2 sheet_pos_f = (logic_pos+1f)/2f; + // screen_pos is -1..+1 + vec2 sheet_pos_f = (screen_pos+1f)/2f; sheet_pos_f *= vec2(1f-sheet_skip_x_frst-sheet_skip_x_last, 1f-sheet_skip_y_frst-sheet_skip_y_last); sheet_pos_f += vec2(sheet_skip_x_frst, sheet_skip_y_frst); @@ -54,28 +61,49 @@ void main() { sheet_pos_f *= sheet_size-n_pts_avg; ivec2 sheet_pos_i = ivec2(round(sheet_pos_f)); - float done = 0; - float steps = 0; + //TODO More priority to points at the center? + //TODO That can help, but still doesn't fix cases where colors change too quickly and turn into a mess + // - First actually try how avg colors look, maybe it's not that bad... + // - Well, now I avg all 3 options. Now it's an extreamly expensive anti-aliasing + // - Much better to just draw +4 points in each coordinate and only then avg + // - Find how it's usually done in OpenGL + // - Or maybe make a pixel for every sheet point use that as a texture??? A waste otherwise + + //float done = 0; + //float steps = 0; + vec3 color = vec3(0); for (int dx = 0; dx>scale_change) + ((y+old_shift_y)>>scale_change)*old_row_len]; + global uint* new_ptr = &new_data[new_shift + x + y*new_row_len]; + + *new_ptr = *old_ptr; +} + +// map multiple points of new_data to single point of old_data +// old_shift and old_row_len count in old_data points +kernel void DownScaleSheet( + global uint* old_data, uint old_shift, uint old_row_len, + global uint* new_data, uint new_shift, uint new_row_len, + int scale_change +) { + uint x = get_global_id(0); + uint y = get_global_id(1); + + uint old_scale_shift = 1u << (scale_change-1); + + constant static uint SIGN_BIT_MASK = 1u << 31; + constant static uint SIGN_BIT_ANTI_MASK = ~SIGN_BIT_MASK; + constant static uint avg_c = 2; + constant static uint avg_area = avg_c*avg_c; + + uint done_count = avg_area/2; + uint steps_sum = avg_area/2; + + for (uint dy=0; dy> 31; + steps_sum += old_v & SIGN_BIT_ANTI_MASK; + } + + new_data[new_shift + x + y*new_row_len] = (done_count/avg_area*SIGN_BIT_MASK) ^ (steps_sum/avg_area); +} + +