From c3f149c91033b05733cec4e447b5ceeb0172397d Mon Sep 17 00:00:00 2001 From: Martin Fouilleul Date: Mon, 15 Aug 2022 17:20:44 +0200 Subject: [PATCH] Fixing metal compute shader coordinate bug that prevented rendering the last row/column of tiles --- src/graphics.c | 45 ++++++++++++++-- src/graphics.h | 5 +- src/metal_shader.metal | 37 +++++++++++-- src/osx_app.mm | 7 ++- src/ui.c | 120 +++++++++++++++++++++++++++++------------ src/ui.h | 18 ++++++- 6 files changed, 184 insertions(+), 48 deletions(-) diff --git a/src/graphics.c b/src/graphics.c index 1f70635..edb8273 100644 --- a/src/graphics.c +++ b/src/graphics.c @@ -3153,17 +3153,37 @@ mg_glyph_info* mg_font_get_glyph_info(mg_font_info* fontInfo, u32 glyphIndex) return(&(fontInfo->glyphs[glyphIndex-1])); } -int mg_font_get_extents(mg_font font, mg_font_extents* outExtents) +mg_font_extents mg_font_get_extents(mg_font font) { mg_font_info* fontInfo = mg_get_font_info(font); if(!fontInfo) { - return(-1); + return((mg_font_extents){}); } - *outExtents = fontInfo->extents; - return(0); + return(fontInfo->extents); } +mg_font_extents mg_font_get_scaled_extents(mg_font font, f32 emSize) +{ + mg_font_info* fontInfo = mg_get_font_info(font); + if(!fontInfo) + { + return((mg_font_extents){}); + } + f32 scale = emSize/fontInfo->unitsPerEm; + mg_font_extents extents = fontInfo->extents; + + extents.ascent *= scale; + extents.descent *= scale; + extents.leading *= scale; + extents.xHeight *= scale; + extents.capHeight *= scale; + extents.width *= scale; + + return(extents); +} + + f32 mg_font_get_scale_for_em_pixels(mg_font font, f32 emSize) { mg_font_info* fontInfo = mg_get_font_info(font); @@ -3763,6 +3783,23 @@ mp_rect mg_glyph_outlines(mg_canvas handle, str32 glyphIndices) return(mg_glyph_outlines_from_font_info(context, fontInfo, glyphIndices)); } +void mg_codepoints_outlines(mg_canvas handle, str32 codePoints) +{ + mg_canvas_data* context = mg_canvas_ptr_from_handle(handle); + if(!context) + { + return; + } + mg_font_info* fontInfo = mg_get_font_info(context->attributes.font); + if(!fontInfo) + { + return; + } + + str32 glyphIndices = mg_font_push_glyph_indices(context->attributes.font, mem_scratch(), codePoints); + mg_glyph_outlines_from_font_info(context, fontInfo, glyphIndices); +} + void mg_text_outlines(mg_canvas handle, str8 text) { mg_canvas_data* context = mg_canvas_ptr_from_handle(handle); diff --git a/src/graphics.h b/src/graphics.h index f677b97..6c83a65 100644 --- a/src/graphics.h +++ b/src/graphics.h @@ -115,13 +115,15 @@ void mg_stream_append(mg_canvas context, mg_stream stream); //------------------------------------------------------------------------------------------ //NOTE(martin): fonts management //------------------------------------------------------------------------------------------ +mg_font mg_font_nil(); mg_font mg_font_create_from_memory(u32 size, byte* buffer, u32 rangeCount, unicode_range* ranges); void mg_font_destroy(mg_font font); //NOTE(martin): the following int valued functions return -1 if font is invalid or codepoint is not present in font// //TODO(martin): add enum error codes -int mg_font_get_extents(mg_font font, mg_font_extents* outExtents); +mg_font_extents mg_font_get_extents(mg_font font); +mg_font_extents mg_font_get_scaled_extents(mg_font font, f32 emSize); f32 mg_font_get_scale_for_em_pixels(mg_font font, f32 emSize); //NOTE(martin): if you need to process more than one codepoint, first convert your codepoints to glyph indices, then use the @@ -187,6 +189,7 @@ void mg_cubic_to(mg_canvas context, f32 x1, f32 y1, f32 x2, f32 y2, f32 x3, f32 void mg_close_path(mg_canvas context); mp_rect mg_glyph_outlines(mg_canvas context, str32 glyphIndices); +void mg_codepoints_outlines(mg_canvas context, str32 string); void mg_text_outlines(mg_canvas context, str8 string); //------------------------------------------------------------------------------------------ diff --git a/src/metal_shader.metal b/src/metal_shader.metal index 86a9b27..d0faf02 100644 --- a/src/metal_shader.metal +++ b/src/metal_shader.metal @@ -101,12 +101,12 @@ kernel void TileKernel(const device float4* boxArray [[buffer(0)]], constant vector_uint2* viewport [[buffer(3)]], uint gid [[thread_position_in_grid]]) { - uint2 tilesMatrixDim = (*viewport + RENDERER_TILE_SIZE - 1) / RENDERER_TILE_SIZE; + uint2 tilesMatrixDim = (*viewport - 1) / RENDERER_TILE_SIZE + 1; uint nTilesX = tilesMatrixDim.x; uint nTilesY = tilesMatrixDim.y; uint triangleIndex = gid; - uint4 box = uint4(floor(boxArray[triangleIndex]/RENDERER_TILE_SIZE)); + uint4 box = uint4(floor(boxArray[triangleIndex]))/RENDERER_TILE_SIZE; uint xMin = max((uint)0, box.x); uint yMin = max((uint)0, box.y); uint xMax = min(box.z, nTilesX-1); @@ -185,15 +185,39 @@ kernel void RenderKernel(texture2d outTexture [[texture(0) uint2 gridSize [[threads_per_grid]]) { //TODO: guard against thread group size not equal to tile size? - const uint2 tilesMatrixDim = (gridSize + RENDERER_TILE_SIZE - 1) / RENDERER_TILE_SIZE; - const uint2 tilePos = tgid * threadsPerThreadgroup / RENDERER_TILE_SIZE; + const uint2 tilesMatrixDim = (gridSize - 1) / RENDERER_TILE_SIZE + 1; +// const uint2 tilePos = tgid * threadsPerThreadgroup / RENDERER_TILE_SIZE; + const uint2 tilePos = gid/RENDERER_TILE_SIZE; const uint tileIndex = tilePos.y * tilesMatrixDim.x + tilePos.x; const device uint* tileBuffer = tilesArray + tileIndex * RENDERER_TILE_BUFFER_SIZE; const uint tileBufferSize = tileCounters[tileIndex]; + +//#define RENDERER_DEBUG_TILES #ifdef RENDERER_DEBUG_TILES //NOTE(martin): color code debug values and show the tile grid + uint nTileX = tilesMatrixDim.x; + uint nTileY = tilesMatrixDim.y; + + if(tilePos.x == 2 && tilePos.y == 12) + { + outTexture.write(float4(1, 0.5, 1, 1), gid); + return; + } + + if(nTileY != 13 || nTileX != 13) + { + outTexture.write(float4(1, 1, 0, 1), gid); + return; + } + + if(tilePos.x > nTileX || tilePos.y > nTileY) + { + outTexture.write(float4(0, 1, 1, 1), gid); + return; + } + if((gid.x % RENDERER_TILE_SIZE == 0) || (gid.y % RENDERER_TILE_SIZE == 0)) { outTexture.write(float4(0, 0, 0, 1), gid); @@ -204,6 +228,11 @@ kernel void RenderKernel(texture2d outTexture [[texture(0) outTexture.write(float4(0, 1, 0, 1), gid); return; } + else + { + outTexture.write(float4(1, 0, 0, 1), gid); + return; + } #endif const float2 sampleOffsets[6] = { float2(5./12, 5./12), float2(-3./12, 3./12), diff --git a/src/osx_app.mm b/src/osx_app.mm index a9d1037..036d569 100644 --- a/src/osx_app.mm +++ b/src/osx_app.mm @@ -592,10 +592,7 @@ static void mp_update_key_state(mp_key_state* key, bool down) key->lastUpdate = frameCounter; } - if(key->down != down) - { - key->transitionCounter++; - } + key->transitionCounter++; key->down = down; } @@ -630,6 +627,8 @@ static void mp_update_mouse_wheel(f32 deltaX, f32 deltaY) static void mp_update_text(utf32 codepoint) { + + printf("update text\n"); u64 frameCounter = __mpAppData.inputState.frameCounter; mp_text_state* text = &__mpAppData.inputState.text; diff --git a/src/ui.c b/src/ui.c index af3e244..1db1153 100644 --- a/src/ui.c +++ b/src/ui.c @@ -564,6 +564,12 @@ ui_box* ui_box_end() return(box); } +void ui_box_set_render_proc(ui_box* box, ui_box_render_proc proc, void* data) +{ + box->renderProc = proc; + box->renderData = data; +} + void ui_box_set_layout(ui_box* box, ui_axis axis, ui_align alignX, ui_align alignY) { box->layout = (ui_layout){axis, {alignX, alignY}}; @@ -716,7 +722,7 @@ void ui_layout_prepass(ui_context* ui, ui_box* box) if(size.kind == UI_SIZE_TEXT) { - box->rect.c[2+i] = textBox.c[2+i]; + box->rect.c[2+i] = textBox.c[2+i] + desiredSize[i].value*2; } else if(size.kind == UI_SIZE_PIXELS) { @@ -793,7 +799,7 @@ void ui_layout_downward_dependent_size(ui_context* ui, ui_box* box, int axis) ui_size* size = &box->computedStyle.size[axis]; if(size->kind == UI_SIZE_CHILDREN) { - box->rect.c[2+axis] = sum; + box->rect.c[2+axis] = sum + size->value*2; } } @@ -980,6 +986,11 @@ void ui_draw_box(mg_canvas canvas, ui_box* box) mg_fill(canvas); } + if((box->flags & UI_FLAG_DRAW_RENDER_PROC) && box->renderProc) + { + box->renderProc(canvas, box, box->renderData); + } + if(box->flags & UI_FLAG_CLIP) { mg_clip_pop(canvas); @@ -1284,7 +1295,7 @@ ui_box* ui_scrollbar(const char* label, f32 thumbRatio, f32* scrollValue) return(frame); } -void ui_panel_begin(const char* name) +ui_box* ui_panel_begin(const char* name) { ui_flags panelFlags = UI_FLAG_DRAW_BACKGROUND | UI_FLAG_DRAW_BORDER @@ -1293,6 +1304,8 @@ void ui_panel_begin(const char* name) ui_box* panel = ui_box_begin(name, panelFlags); ui_box* innerView = ui_box_begin(name, 0); + + return(panel); } void ui_panel_end() @@ -1794,11 +1807,56 @@ str32 ui_edit_perform_operation(ui_context* ui, ui_edit_op operation, ui_edit_mo ui->editMark = codepoints.len; } break; } - ui->editCursorBlinkStart = ui->frameTime; + ui->editCursorBlinkStart = ui->frameTime; return(codepoints); } +void ui_text_box_render(mg_canvas canvas, ui_box* box, void* data) +{ + str32 codepoints = *(str32*)data; + ui_context* ui = ui_get_context(); + + u32 firstDisplayedChar = 0; + if(ui_box_active(box)) + { + firstDisplayedChar = ui->editFirstDisplayedChar; + } + + ui_style* style = &box->computedStyle; + mg_font_extents extents = mg_font_get_scaled_extents(style->font, style->fontSize); + f32 lineHeight = extents.ascent + extents.descent; + + str32 before = str32_slice(codepoints, 0, firstDisplayedChar); + mp_rect beforeBox = mg_text_bounding_box_utf32(style->font, style->fontSize, before); + + f32 textMargin = 5; + + f32 textX = textMargin + box->rect.x - beforeBox.w; + f32 textTop = box->rect.y + 0.5*(box->rect.h - lineHeight); + f32 textY = textTop + extents.ascent ; + + if(box->active && !((u64)(2*(ui->frameTime - ui->editCursorBlinkStart)) & 1)) + { + str32 beforeCaret = str32_slice(codepoints, 0, ui->editCursor); + mp_rect beforeCaretBox = mg_text_bounding_box_utf32(style->font, style->fontSize, beforeCaret); + + f32 caretX = box->rect.x + textMargin - beforeBox.w + beforeCaretBox.w; + f32 caretY = textTop; + + mg_set_color_rgba(canvas, 0, 0, 0, 1); + mg_rectangle_fill(canvas, caretX, caretY, 2, lineHeight); + } + + mg_set_font(canvas, style->font); + mg_set_font_size(canvas, style->fontSize); + mg_set_color(canvas, style->fontColor); + + mg_move_to(canvas, textX, textY); + mg_codepoints_outlines(canvas, codepoints); + mg_fill(canvas); +} + ui_text_box_result ui_text_box(const char* name, mem_arena* arena, str8 text) { ui_context* ui = ui_get_context(); @@ -1809,18 +1867,13 @@ ui_text_box_result ui_text_box(const char* name, mem_arena* arena, str8 text) | UI_FLAG_DRAW_BACKGROUND | UI_FLAG_DRAW_BORDER | UI_FLAG_CLIP + | UI_FLAG_DRAW_RENDER_PROC | UI_FLAG_SCROLLABLE; ui_box* frame = ui_box_make(name, frameFlags); - ui_box_set_layout(frame, UI_AXIS_X, UI_ALIGN_START, UI_ALIGN_CENTER); - - ui_box_push(frame); - - ui_flags contentsFlags = UI_FLAG_DRAW_TEXT; - - ui_box* contents = ui_box_make_str8(text, contentsFlags); - ui_box_set_size(contents, UI_AXIS_X, UI_SIZE_TEXT, 0, 1); - ui_box_set_size(contents, UI_AXIS_Y, UI_SIZE_PARENT_RATIO, 1, 1); + ui_style* style = &frame->computedStyle; + mg_font_extents extents = mg_font_get_scaled_extents(style->font, style->fontSize); + ui_box_set_size(frame, UI_AXIS_Y, UI_SIZE_PIXELS, extents.ascent+extents.descent+10, 1); ui_sig sig = ui_box_sig(frame); @@ -1865,7 +1918,6 @@ ui_text_box_result ui_text_box(const char* name, mem_arena* arena, str8 text) if(ui_box_active(frame)) { str32 codepoints = utf8_push_to_codepoints(&ui->frameArena, text); - ui->editCursor = Clamp(ui->editCursor, 0, codepoints.len); ui->editMark = Clamp(ui->editMark, 0, codepoints.len); @@ -1873,9 +1925,8 @@ ui_text_box_result ui_text_box(const char* name, mem_arena* arena, str8 text) str32 input = mp_input_text_utf32(&ui->frameArena); if(input.len) { - printf("text edit got input\n"); - codepoints = ui_edit_replace_selection_with_codepoints(ui, codepoints, input); + ui->editCursorBlinkStart = ui->frameTime; } //NOTE handle shortcuts @@ -1891,27 +1942,30 @@ ui_text_box_result ui_text_box(const char* name, mem_arena* arena, str8 text) } } + //NOTE(martin): text box focus shortcuts + if(mp_input_key_pressed(MP_KEY_ENTER)) + { + //TODO(martin): extract in gui_edit_complete() (and use below) + ui_box_deactivate(frame); + ui->focus = 0; + } + result.text = utf8_push_from_codepoints(arena, codepoints); - contents->string = str8_push_copy(&ui->frameArena, result.text); - //NOTE slide contents + //TODO slide contents - //NOTE position caret - - str32 displayedBeforeCaret = str32_slice(codepoints, ui->editFirstDisplayedChar, ui->editCursor); - ui_style* style = &contents->computedStyle; - mp_rect extents = mg_text_bounding_box_utf32(style->font, style->fontSize, displayedBeforeCaret); - - ui_push_fg_color((mg_color){0, 0, 0, 1}); - ui_box* caret = ui_box_make("caret", UI_FLAG_DRAW_FOREGROUND); - ui_box_set_size(caret, UI_AXIS_X, UI_SIZE_PIXELS, 2, 1); - ui_box_set_size(caret, UI_AXIS_Y, UI_SIZE_PARENT_RATIO, 0.9, 1); - ui_box_set_floating(caret, UI_AXIS_X, extents.x + extents.w); - - ui_pop_fg_color(); + //NOTE: set renderer + str32* renderCodepoints = mem_arena_alloc_type(&ui->frameArena, str32); + *renderCodepoints = str32_push_copy(&ui->frameArena, codepoints); + ui_box_set_render_proc(frame, ui_text_box_render, renderCodepoints); + } + else + { + //NOTE: set renderer + str32* renderCodepoints = mem_arena_alloc_type(&ui->frameArena, str32); + *renderCodepoints = utf8_push_to_codepoints(&ui->frameArena, text); + ui_box_set_render_proc(frame, ui_text_box_render, renderCodepoints); } - - ui_box_pop(); return(result); } diff --git a/src/ui.h b/src/ui.h index 4034ee0..fcf316e 100644 --- a/src/ui.h +++ b/src/ui.h @@ -29,6 +29,7 @@ typedef enum UI_FLAG_DRAW_FOREGROUND = (1<<7), UI_FLAG_DRAW_BORDER = (1<<8), UI_FLAG_DRAW_TEXT = (1<<9), + UI_FLAG_DRAW_RENDER_PROC = (1<<10), } ui_flags; @@ -118,6 +119,8 @@ typedef struct ui_sig } ui_sig; +typedef void(*ui_box_render_proc)(mg_canvas canvas, ui_box* box, void* data); + struct ui_box { // hierarchy @@ -135,6 +138,9 @@ struct ui_box str8 string; // styling and layout + ui_box_render_proc renderProc; + void* renderData; + ui_style_tag tag; ui_style* targetStyle; ui_style computedStyle; @@ -173,6 +179,12 @@ ui_box* ui_box_begin_str8(str8 string, ui_flags flags); ui_box* ui_box_end(); #define ui_container(name, flags) defer_loop(ui_box_begin(name, flags), ui_box_end()) +void ui_box_push(ui_box* box); +void ui_box_pop(); +ui_box* ui_box_top(); + +void ui_box_set_render_proc(ui_box* box, ui_box_render_proc proc, void* data); + void ui_box_set_layout(ui_box* box, ui_axis axis, ui_align alignX, ui_align alignY); void ui_box_set_size(ui_box* box, ui_axis axis, ui_size_kind kind, f32 value, f32 strictness); void ui_box_set_floating(ui_box* box, ui_axis axis, f32 pos); @@ -193,6 +205,7 @@ void ui_push_font_color(mg_color color); void ui_push_border_size(f32 size); void ui_push_border_color(mg_color color); void ui_push_roundness(f32 roundness); +void ui_push_animation_time(f32 time); void ui_push_bg_color_ext(ui_style_tag tag, ui_style_selector selector, mg_color color); void ui_push_fg_color_ext(ui_style_tag tag, ui_style_selector selector, mg_color color); @@ -202,6 +215,7 @@ void ui_push_font_color_ext(ui_style_tag tag, ui_style_selector selector, mg_col void ui_push_border_size_ext(ui_style_tag tag, ui_style_selector selector, f32 size); void ui_push_border_color_ext(ui_style_tag tag, ui_style_selector selector, mg_color color); void ui_push_roundness_ext(ui_style_tag tag, ui_style_selector selector, f32 roundness); +void ui_push_animation_time_ext(ui_style_tag tag, ui_style_selector selector, f32 time); void ui_pop_bg_color(); void ui_pop_fg_color(); @@ -211,7 +225,7 @@ void ui_pop_font_color(); void ui_pop_border_size(); void ui_pop_border_color(); void ui_pop_roundness(); - +void ui_pop_animation_time(); // Basic helpers enum { @@ -228,7 +242,7 @@ ui_sig ui_label(const char* label); ui_sig ui_button(const char* label); ui_box* ui_scrollbar(const char* label, f32 thumbRatio, f32* scrollValue); -void ui_panel_begin(const char* name); +ui_box* ui_panel_begin(const char* name); void ui_panel_end(); #define ui_panel(name) defer_loop(ui_panel_begin(name), ui_panel_end())