diff --git a/LastPack.log b/LastPack.log index 92499fcf..a8c3f53f 100644 --- a/LastPack.log +++ b/LastPack.log @@ -621,6 +621,14 @@ 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: Compile errоrs: +Tester: Compiling: [42,48] GL_CL_Context.pas: Incompatible types: 'cl_context' and 'CLDevice' +Tester: Compiling "Samples/OpenGLABC/Mandelbrot/FieldTest.pas" +Tester: Compiling: Compile errоrs: +Tester: Compiling: [15,31] FieldTest.pas: Undefined name reference 'BlocksForCurrentScale' +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 +840,9 @@ Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/Si Tester: Done executing Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/SimpleAddition но с кешем] Tester: Done executing +Tester: WARNING: Test[Samples/OpenGLABC/Mandelbrot/FieldTest] wasn't executed because of prior errоrs or changes to .td +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,7 +1052,11 @@ Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/Si Tester: Done executing Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/SimpleAddition но с кешем] Tester: Done executing +Tester: WARNING: Test[Samples/OpenGLABC/Mandelbrot/FieldTest] wasn't executed because of prior errоrs or changes to .td +Tester: Executing Test[Samples/OpenGLABC/Mandelbrot/SamplingTests] +Tester: Done executing Tester: Cleanup +Tester: WARNING: Not all loaded tests were executed Tester: Done testing Finished runing Tester ===================================== Release ==================================== @@ -1070,6 +1085,23 @@ 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/FieldTest00.bmp" +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/PointComponents.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/SamplingTests.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Settings.pas" +WARNING: Sample file with unknown extension: +Packing sample file "Samples/OpenGLABC/Mandelbrot/0/Mandel_zoom_00_mandelbrot_set.jpg" +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 +1111,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..5119230f 100644 --- a/Log/Release.log +++ b/Log/Release.log @@ -23,6 +23,23 @@ 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/FieldTest00.bmp" +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/PointComponents.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/SamplingTests.pas" +Packing sample file "Samples/OpenGLABC/Mandelbrot/Settings.pas" +WARNING: Sample file with unknown extension: +Packing sample file "Samples/OpenGLABC/Mandelbrot/0/Mandel_zoom_00_mandelbrot_set.jpg" +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 +49,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..7e86634f 100644 --- a/Log/Test.log +++ b/Log/Test.log @@ -440,5 +440,6 @@ Tester: Done executing Tester: Executing Test[Samples/OpenCLABC/Кеширование CLProgramCode/SimpleAddition но с кешем] Tester: Done executing Tester: Cleanup +Tester: WARNING: Not all loaded tests were executed Tester: Done testing Finished runing Tester diff --git a/Samples/OpenGLABC/Mandelbrot/0/Mandel_zoom_00_mandelbrot_set.jpg b/Samples/OpenGLABC/Mandelbrot/0/Mandel_zoom_00_mandelbrot_set.jpg new file mode 100644 index 00000000..e35842e1 Binary files /dev/null and b/Samples/OpenGLABC/Mandelbrot/0/Mandel_zoom_00_mandelbrot_set.jpg differ diff --git a/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas b/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas index 01319ff4..30548b56 100644 --- a/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas +++ b/Samples/OpenGLABC/Mandelbrot/0Mandelbrot.pas @@ -13,8 +13,11 @@ // - "+" и "-": Гладкое изменение масштаба // - Alt+Enter: Полноэкранный режим -//TODO Доделать ограничение view_bound -// - Сейчас его не возвращает из функции, создающей блоки +//TODO Сделать слежение за выделением памяти GPU +// - Может в OpenCLABC, но тогда выделения на стороне OpenGL надо вручную всё равно обрабатывать + +//TODO Исследовать прямые концы внешнего круга - сколько там шагов и почему +// - Надо возможность приближать и возможность наводить мышку чтобы увидеть инфу // Константы, которые можно быстро менять uses Settings; @@ -155,6 +158,7 @@ BoundUniforms = record //TODO Использовать чтобы выдавать кол-во итераций под курсором var mouse_pos: Vec2i; + var scale_speed := 0.0; {$region Управление} begin @@ -165,22 +169,12 @@ BoundUniforms = record // Result := pos*camera.scale + camera.pos; // end; - 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)-> case e.KeyCode of Keys.Space: camera := new CameraPos(f.Width, f.Height); end; + f.MouseWheel += (o,e)->(scale_speed += e.Delta/1000); f.MouseMove += (o,e)->(mouse_pos := new Vec2i(e.X,e.Y)); end; @@ -192,19 +186,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 +204,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,19 +212,14 @@ 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; choose_frag_shader(s_frag_rainbow); - var cl_err_buffer := new CLArray(3); - var cl_uc_buffer := 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 Q_Init, Q_Acquire, Q_Release: CommandQueueNil; var ensure_sheet_buffer_size := procedure(w, h: integer)-> begin var req_size := w*h; @@ -249,14 +231,13 @@ BoundUniforms = record 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); + //TODO Вместо этого надо сделать 2 буфера и копировать данные из предыдущего когда изменяются границы sheet 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 ; + Q_Acquire := CQAcquireGL(cl_sheet_buffer); + Q_Release := CQReleaseGL(cl_sheet_buffer); end; // Для дебага @@ -265,11 +246,23 @@ 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); + + {$endregion Общие данные для всех кадров} while true do begin + if scale_speed<>0 then + begin + camera.ApplyScaleSpeed(scale_speed); + scale_speed /= 2; + end; + + var camera := camera; // Скопировать локально + var render_info := BlockLayer.GetLayer(camera).GetRenderInfo(camera); + BlockUpdater.SetCurrent(render_info.blocks); + var curr_frame_resized := false; if need_resize then begin @@ -280,48 +273,55 @@ BoundUniforms = record camera.Resize(w_size.Width, w_size.Height); end; - var render_info := BlockLayer.BlocksForCurrentScale(camera); + t_body.Start; + {$region Кадр} + gl.Clear(glClearBufferMask.COLOR_BUFFER_BIT); + +// clGLSharingKHR.PlatformLess.EnqueueAcquireGLObjects(q, 1, cl_temp0, 0,IntPtr.Zero,IntPtr.Zero).RaiseIfError; +// +// cl.SetKernelArg(k, 0, new UIntPtr(cl_mem.Size), cl_temp0).RaiseIfError; +// cl.EnqueueNDRangeKernel(q, k, 1, nil, |new UIntPtr(ram_temp0.Length)|, nil, 0,IntPtr.Zero,IntPtr.Zero).RaiseIfError; +// +// clGLSharingKHR.PlatformLess.EnqueueReleaseGLObjects(q, 1, cl_temp0, 0,IntPtr.Zero,IntPtr.Zero).RaiseIfError; + +// cl.Finish(q); +// +// gl.GetNamedBufferSubData(gl_temp0, IntPtr.Zero, new UIntPtr(ram_temp0.Length*sizeof(integer)), ram_temp0); +// ram_temp0.Println; + 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; + var render_sheet_w := b_cx * block_w; + var render_sheet_h := b_cy * block_w; ensure_sheet_buffer_size(render_sheet_w, render_sheet_h); - var Q_Steps := CQNil; //TODO Calculate in separate thread var Q_Extract := CQNil; 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 sheet_shift := block_w * (b_x + b_y*render_sheet_w); + Q_Extract += b.CQ_GetData( + new ShiftedCLArray(cl_sheet_buffer, sheet_shift, render_sheet_w) ); 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; + CLContext.Default.SyncInvoke( + Q_Acquire + + Q_Init + + Q_Extract + + Q_Release + ); +// 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.NamedBufferSubData(buffer_temp, new IntPtr(0*sizeof(real)), new IntPtr(2*sizeof(real)), mouse_pos); @@ -332,13 +332,14 @@ BoundUniforms = record // var temp_data := new real[1]; // gl.GetNamedBufferSubData(buffer_temp, new IntPtr(2*sizeof(real)), new IntPtr(1*sizeof(real)), temp_data); + {$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; @@ -346,22 +347,26 @@ BoundUniforms = record // title_parts += $'temp_data={_ObjectToString(temp_data)}'; //TODO Оттестировать и убрать + 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})'; - f.Text := title_parts.JoinToString(', '); + f.BeginInvoke(()-> + try + f.Text := title_parts.JoinToString(', '); + except + on e: Exception do + MessageBox.Show(e.ToString); + end); EndFrame; end; diff --git a/Samples/OpenGLABC/Mandelbrot/Blocks.pas b/Samples/OpenGLABC/Mandelbrot/Blocks.pas index e161cd4d..920c7ded 100644 --- a/Samples/OpenGLABC/Mandelbrot/Blocks.pas +++ b/Samples/OpenGLABC/Mandelbrot/Blocks.pas @@ -33,14 +33,11 @@ ShiftedCLArray = record // Длина стороны блока в логическом пространстве = 2**block_scale private block_scale: integer; private component_word_count: integer; - private pos00: CLArray; - - private gpu_data: CLArray; - private gpu_mipmaps_state: CLArray; - private gpu_mipmaps_steps: CLArray; - private gpu_mipmaps_need_update: CLArray; + private ram_pos00: PointPos; + private gpu_pos00: CLArray; private ram_data: array of byte; + private gpu_data: CLArray; public constructor(block_scale: integer; pos00: PointPos); begin @@ -50,20 +47,14 @@ ShiftedCLArray = record 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.ram_pos00 := pos00; + self.gpu_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.gpu_data := new CLArray( block_w*block_w * (2 + component_word_count*2) ); +// //TODO Снять с этого потока выполнения... +// // - Что то типа CQ_Init +// // - Но +// self.gpu_data.FillValue(0); end; private constructor := raise new System.InvalidOperationException; @@ -71,71 +62,48 @@ ShiftedCLArray = record 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_Init: CommandQueueNil; + begin + Result := CQNil; + if gpu_data<>nil then exit; + + if ram_data<>nil then raise new System.NotImplementedException; + + gpu_data := new CLArray( block_w*block_w * (2 + component_word_count*2) ); + Result += gpu_data.MakeCCQ.ThenFillValue(0).DiscardResult; + + end; 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 + , self.gpu_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_GetData(A_Result: ShiftedCLArray): CommandQueueNil; begin - - {$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} - - var code := CLCodeFor(self.component_word_count); - var w := block_w; - - if target_mipmap_lvl=0 then - 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; - - 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; - - for var mipmap_lvl := 2 to target_mipmap_lvl do - 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) + Result := CQNil; + var gpu_data := self.gpu_data; + if gpu_data=nil then exit; + Result := CLCodeFor(self.component_word_count)['ExtractSteps'] + .MakeCCQ.ThenExec2(block_w,block_w + , gpu_data + , A_Result.a, A_Result.shift, A_Result.row_len ).DiscardResult; - mipmap_shift += w*w; - end; + end; + + public procedure Dispose; + begin - 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; + var gpu_data := System.Threading.Interlocked.Exchange(self.gpu_data, nil); + if gpu_data<>nil then gpu_data.Dispose; end; + protected procedure Finalize; override := Dispose; end; @@ -154,8 +122,7 @@ BoundDefs = record end; BlockLayerRenderInfo = record - public blocks: array[,] of PointBlock; // [y,x] - public mipmap_lvl: integer; + public c_min, c_max: PointPos; // How much of viewport is empty // [0;2) and first+last<=2 @@ -165,6 +132,7 @@ BlockLayerRenderInfo = record // [0;1) and first+last<=1 public sheet_bound: BoundDefs; + public blocks: array[,] of PointBlock; // [y,x] end; // Слой, содержащий кэш уже просчитанных блоков @@ -180,85 +148,95 @@ BlockLayerRenderInfo = record 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; + private static all_layers := new List; + public static function GetLayer(camera_pos: CameraPos): BlockLayer; + begin + var block_scale := camera_pos.GetPointScale + Settings.block_w_pow; + var layer_ind := Settings.max_block_scale - block_scale; + while all_layers.Count<=layer_ind do + all_layers += default(BlockLayer); + Result := all_layers[layer_ind]; + if Result<>nil then exit; + + //TODO Remove + for var i := 0 to all_layers.Count-1 do + begin + if all_layers[i]=nil then continue; + foreach var bl in all_layers[i].blocks.Values do + bl.Dispose; + all_layers[i] := nil; + end; + + Result := new BlockLayer(block_scale); + all_layers[layer_ind] := Result; + end; - //TODO Сейчас блоки создаёт с 0 при каждом вызове. Использовать экземпляр типа BlockLayer - public static function BlocksForCurrentScale(camera_pos: CameraPos): BlockLayerRenderInfo; + private function GetBlockAt(r, i: PointComponent): PointBlock; + begin + var p := new PointPos(r, i); + if self.blocks.TryGetValue(p, Result) then exit; + Result := new PointBlock(self.scale, p); + self.blocks.Add(p, Result); + end; + public function GetRenderInfo(camera_pos: CameraPos): BlockLayerRenderInfo; begin Result := default(BlockLayerRenderInfo); + {$ifdef DEBUG} + if self.scale <> camera_pos.GetPointScale + 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; + + // Actual space is -2 .. +2 + // Visible space is c-visible_space_d .. c+visible_space_d + // View bound is - //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; + Result.c_min := c_ctr.AddLowestBitsC2(-dr,-di); + Result.c_max := c_ctr.AddLowestBitsC2(+dr,+di); + Result.c_max.SelfFlipIfMinusZero; - 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); + var block_sz_bit_ind := Settings.z_int_bits-1 + -self.scale; + Result.c_min.SelfBlockRound(block_sz_bit_ind, false, Result.sheet_bound.xf, Result.sheet_bound.yf); + Result.c_max.SelfBlockRound(block_sz_bit_ind, true, Result.sheet_bound.xl, Result.sheet_bound.yl); //TODO Debug error for when bounds are >2 - 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); + var r_blocks_count := PointComponent.BlocksCount(Result.c_min.r, Result.c_max.r, block_sz_bit_ind); + var i_blocks_count := PointComponent.BlocksCount(Result.c_min.i, Result.c_max.i, block_sz_bit_ind); Result.sheet_bound /= new System.ValueTuple(r_blocks_count, i_blocks_count); var pc_rs := new PointComponent[r_blocks_count]; begin - var pc_r := c_min.r; + var pc_r := Result.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 + if pc_r<>Result.c_max.r then raise new System.InvalidOperationException; {$endif DEBUG} end; @@ -266,15 +244,15 @@ BlockLayerRenderInfo = record Result.blocks := new PointBlock[i_blocks_count, r_blocks_count]; // $'Need {i_blocks_count} x {r_blocks_count} = {Result.blocks.Length} blocks'.Println; Halt; begin - var pc_i := c_min.i; + var pc_i := Result.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)); + Result.blocks[ii, ri] := GetBlockAt(pc_rs[ri], pc_i); pc_i := pc_i.MakeNextBlockBound(block_sz_bit_ind); end; {$ifdef DEBUG} - if pc_i<>c_max.i then + if pc_i<>Result.c_max.i then raise new System.InvalidOperationException; {$endif DEBUG} end; @@ -320,15 +298,23 @@ BlockLayerRenderInfo = record if blocks<>last_blocks then begin var branches := ArrFill(Settings.max_parallel_blocks, CQNil); + + var Q_Init := CQNil; foreach var b: PointBlock in blocks index b_i do + begin + Q_Init += b.CQ_Init; branches[b_i mod branches.Length] += b.CQ_MandelbrotBlockStep(P_StepCount, V_UpdateCount, A_Err); - last_blocks := blocks; + end; + CLContext.Default.SyncInvoke(Q_Init); + 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; + + last_blocks := blocks; end; sw.Restart; diff --git a/Samples/OpenGLABC/Mandelbrot/CameraDef.pas b/Samples/OpenGLABC/Mandelbrot/CameraDef.pas index e57c5a7a..82c2cca8 100644 --- a/Samples/OpenGLABC/Mandelbrot/CameraDef.pas +++ b/Samples/OpenGLABC/Mandelbrot/CameraDef.pas @@ -45,24 +45,15 @@ CameraPos = record public function AspectRatio := real(dw)/real(dh); - public function GetPointScaleAndMainMipMapLvl: System.ValueTuple; - 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 GetPointScale := + (self.scale_pow + Settings.scale_shift + Floor(Log2(self.scale_fine/dh))) + .ClampTop(Settings.max_block_scale - Settings.block_w_pow); + private function GetBitCount := Settings.z_int_bits + -GetPointScale + 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; @@ -81,6 +72,12 @@ CameraPos = record end; + public procedure ApplyScaleSpeed(scale_speed: real); + begin + self.scale_fine *= 0.9 ** scale_speed; + FixScalePow; + end; + end; end. \ No newline at end of file diff --git a/Samples/OpenGLABC/Mandelbrot/FieldTest.td b/Samples/OpenGLABC/Mandelbrot/FieldTest.td index b4f528df..f7503ea1 100644 --- a/Samples/OpenGLABC/Mandelbrot/FieldTest.td +++ b/Samples/OpenGLABC/Mandelbrot/FieldTest.td @@ -5,20 +5,13 @@ Compile errors: [225,27] Blocks.pas: The type 'PointComponent' does not contain a definition for 'FirstWordToReal' -#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 +#ExpExecErr +System.ArgumentException: Value of '738559' is not valid for 'red'. 'red' should be greater than or equal to 0 and less than or equal to 255. + at System.Drawing.Color.CheckByte(Int32 value, String name) + at System.Drawing.Color.FromArgb(Int32 alpha, Int32 red, Int32 green, Int32 blue) + at System.Drawing.Color.FromArgb(Int32 red, Int32 green, Int32 blue) + at FieldTest.Program.$Main() in *\POCGL\Samples\OpenGLABC\Mandelbrot\FieldTest.pas:line * + at FieldTest.Program.Main() #ReqModules OpenCLABC diff --git a/Samples/OpenGLABC/Mandelbrot/FieldTest00.bmp b/Samples/OpenGLABC/Mandelbrot/FieldTest00.bmp new file mode 100644 index 00000000..553d5008 Binary files /dev/null and b/Samples/OpenGLABC/Mandelbrot/FieldTest00.bmp differ diff --git a/Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.cl b/Samples/OpenGLABC/Mandelbrot/MandelbrotSampling.cl index e58034e6..01c5d248 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 @@ -356,7 +358,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 +365,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,160 +399,50 @@ 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 +kernel void ExtractSteps( + global point_info* block, + global uint* result_data, uint result_shift, uint result_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( - 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 -) { + bool state = block[x + y*BLOCK_W].state != 0; + uint steps = block[x + y*BLOCK_W].steps; + uint data = steps ^ (state ? SIGN_BIT_MASK : 0); - 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 - ); + global uint* p_result_data = &result_data[result_shift + x + y*result_row_len]; + if (data < *p_result_data) return; + *p_result_data = data; } -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 +kernel void CopySheetRect( + global uint* old_data, uint old_shift, uint old_row_len, + global uint* new_data, uint new_shift, uint new_row_len ) { - 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; - - 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_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; - + //TODO } -kernel void FixFirstMipMap( - global point_info* block, - global uchar* mip_map_state, - global uint* mip_map_steps, - global uchar* need_update +kernel void UpScaleSheet( + 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 ) { - - 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 - ); - + //TODO } -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 +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 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 - ); - + //TODO } diff --git a/Samples/OpenGLABC/Mandelbrot/PointComponents.pas b/Samples/OpenGLABC/Mandelbrot/PointComponents.pas index 951f7ec2..45d8bb17 100644 --- a/Samples/OpenGLABC/Mandelbrot/PointComponents.pas +++ b/Samples/OpenGLABC/Mandelbrot/PointComponents.pas @@ -77,6 +77,16 @@ PointComponent = record(System.IEquatable) Result := sb.ToString; 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} @@ -98,7 +108,9 @@ PointComponent = record(System.IEquatable) x *= int64(1) shl d_shift; Result := Convert.ToInt64(x); end; - public function AddLowestBits(d: int64): PointComponent; + /// Add d to the lowest words, while clamping the result to -2..+2 + public function AddLowestBitsC2(d: int64): PointComponent; + const v2: cardinal = 1 shl (33-Settings.z_int_bits); begin var size := self.Words.Length; var self_sign := self.Words[0] and sign_bit_mask; @@ -128,10 +140,21 @@ PointComponent = record(System.IEquatable) 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 @@ -330,21 +353,29 @@ PointComponent = record(System.IEquatable) end; - PointPos = record + PointPos = record(System.IEquatable) 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 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 function AddLowestBitsC2(dr,di: int64) := new PointPos( + r.AddLowestBitsC2(dr), + i.AddLowestBitsC2(di) ); public procedure SelfBlockRound(block_sz_bit_ind: integer; round_up: boolean; var skip_r: single; var skip_i: single); diff --git a/Samples/OpenGLABC/Mandelbrot/Settings.pas b/Samples/OpenGLABC/Mandelbrot/Settings.pas index c1d1319c..f4678250 100644 --- a/Samples/OpenGLABC/Mandelbrot/Settings.pas +++ b/Samples/OpenGLABC/Mandelbrot/Settings.pas @@ -6,22 +6,37 @@ // Если камеру приближают - просчитываются 4 более блока меньшего масштаба: // Cтолько же точек, но эти точки упакованы ближе друг к другу // Чем больше блоки - тем быстрее их считает и тем больше использует памяти -const block_w_pow = 9; +const block_w_pow = 11; //TODO 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х4 вокруг точки 0;0) разбита на минимум 4 блока, по 1 на угол // При приближении камеры каждый из этих угловых блоков будет далее разбивать на 4 меньших блока const max_block_scale = 1; // Не менять - эта константа много где неявно задана +// Данная программа хранит текущее состояние и номер шага +// рекуррентной функции "z_next(z) = z*z + c" для каждой точки в каждом блоке +// Где "z" это предыдущее значение (изначально 0), а "c" это координата точки +// В отдельном от графики потоке выполнения (класс Blocks.BlockUpdater) +// циклически берёт блоки текущего масштаба и делает ещё несколько шагов +// Максимум за 1 итерацию обработки может выполнить max_steps_at_once шагов для всех точек +const max_steps_at_once = 1024; +// Чем больше шагов - тем меньшая доля времени будет потрачена на синхронизацию и т.п. между обработками +// Но если обработка блоков использует GPU слишком эффективно - она заберёт все ресурсы у системы и графика начнёт лагать +// Поэтому кол-во шагов в 1 обработке меняет в диапазоне 1..max_steps_at_once, так чтобы 1 обработка заменяла примерно столько секунд: +const target_step_time_seconds = 0.01; //TODO 0.1 +//TODO Когда кадр закончило рисовать - можно всё остальное время использовать GPU на максимум +// - Но это не тривиально описать в программе... +// Все блоки можно обрабатывать по-очереди или параллельно +// Чтобы достичь большой параллельности - надо выделить много ресурсов в виде OpenCL.cl_command_queue (личные данные каждого потока выполнения GPU) +// Но если обрабатывать все блоки 1 потоком - GPU будет повторно оставаться без работы на короткие промежутки времени, между обработками блоков +// Поэтому для обработки блоков будет выделено не больше чем max_parallel_blocks потоков GPU: +const max_parallel_blocks = 2; + // Отрисованные блоки хранит в VRAM (памяти GPU) const max_VRAM = 1610612736; // 1.5 GB // Если VRAM заканчивается - старые блоки отправляет в RAM (в обычную оперативную память) @@ -30,17 +45,6 @@ // (иначе его отправит в файл подкачки, а потом, при подгрузке назад - будет лагать) const RAM_life_span_seconds = 15*60; // 15 минут -// Чтобы для каждого кадра не пересчитывать -// среднее арифметическое всех точек под каждым пикселем, -// Каждому блоку создаст мипмапы таких размеров: -// -// (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_next|>2 // На шаге=0 z=0, а значит для точек |c|>2 на шаг 1 уже не переключится // На шаге>0 |z|<=2, иначе на этот шаг не переключилось бы