Fixing metal compute shader coordinate bug that prevented rendering the last row/column of tiles

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(fontInfo->extents);
-	return(0);
 }
 
+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);

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);
 
 //------------------------------------------------------------------------------------------

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<float, access::write> 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 tilesMatrixDim = (gridSize - 1) / RENDERER_TILE_SIZE + 1;
-	const uint2 tilePos = tgid * threadsPerThreadgroup / RENDERER_TILE_SIZE;
+//	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<float, access::write> 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),

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->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;
 

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()
@@ -1799,6 +1812,51 @@ str32 ui_edit_perform_operation(ui_context* ui, ui_edit_op operation, ui_edit_mo
 	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_style* style = &frame->computedStyle;
-
+	mg_font_extents extents = mg_font_get_scaled_extents(style->font, style->fontSize);
-	ui_box_push(frame);
+	ui_box_set_size(frame, UI_AXIS_Y, UI_SIZE_PIXELS, extents.ascent+extents.descent+10, 1);
-
-	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_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)
 			}
 		}
 
-		result.text = utf8_push_from_codepoints(arena, codepoints);
+		//NOTE(martin): text box focus shortcuts
-		contents->string = str8_push_copy(&ui->frameArena, result.text);
+		if(mp_input_key_pressed(MP_KEY_ENTER))
-
+		{
-		//NOTE slide contents
+			//TODO(martin): extract in gui_edit_complete() (and use below)
-
+			ui_box_deactivate(frame);
-		//NOTE position caret
+			ui->focus = 0;
-
-		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();
 		}
 
-	ui_box_pop();
+		result.text = utf8_push_from_codepoints(arena, codepoints);
+
+		//TODO slide contents
+
+		//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);
+	}
 
 	return(result);
 }

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())