From f1ec3385de10003441524401e4f95fcb1e086dde Mon Sep 17 00:00:00 2001 From: Patrick Walton Date: Mon, 23 Jan 2017 15:33:12 -0800 Subject: [PATCH] Draft shader code for the new tessellation-based algorithm --- resources/shaders/accum.cl | 41 ++---- resources/shaders/draw.cl | 241 -------------------------------- resources/shaders/draw.fs.glsl | 116 +++++++++++++++ resources/shaders/draw.tcs.glsl | 112 +++++++++++++++ resources/shaders/draw.tes.glsl | 67 +++++++++ resources/shaders/draw.vs.glsl | 49 +++++++ 6 files changed, 358 insertions(+), 268 deletions(-) delete mode 100644 resources/shaders/draw.cl create mode 100644 resources/shaders/draw.fs.glsl create mode 100644 resources/shaders/draw.tcs.glsl create mode 100644 resources/shaders/draw.tes.glsl create mode 100644 resources/shaders/draw.vs.glsl diff --git a/resources/shaders/accum.cl b/resources/shaders/accum.cl index b0c4bb13..001f2bd1 100644 --- a/resources/shaders/accum.cl +++ b/resources/shaders/accum.cl @@ -12,39 +12,26 @@ // // This proceeds top to bottom for better data locality. For details on the algorithm, see [1]. // -// [1]: http://nothings.org/gamedev/rasterize/ +// [1]: https://medium.com/@raphlinus/inside-the-fastest-font-renderer-in-the-world-75ae5270c445 -#define TILE_SIZE 4 +const sampler_t SAMPLER = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; -__global const int *getPixel(__global const int *gPixels, uint2 point, uint widthInTiles) { - uint2 tile = point / TILE_SIZE, pointInTile = point % TILE_SIZE; - return &gPixels[(tile.y * widthInTiles + tile.x) * TILE_SIZE * TILE_SIZE + - pointInTile.y * TILE_SIZE + - pointInTile.x]; -} - -__kernel void accum(__global const int *gPixels, - __write_only image2d_t gTexture, - uint4 kAtlasRect, +__kernel void accum(__write_only image2d_t gTexture, + __read_only image2d_t gCoverage, + uint kAtlasWidth, uint kAtlasShelfHeight) { - // Compute our column. - int globalID = get_global_id(0); - uint atlasWidth = kAtlasRect.z - kAtlasRect.x; - uint x = globalID % atlasWidth; - uint widthInTiles = atlasWidth / TILE_SIZE; - - // Compute the row range we'll traverse. - uint shelf = globalID % atlasWidth; - uint yStart = shelf * kAtlasShelfHeight; - uint yEnd = yStart + kAtlasShelfHeight; + // Determine the boundaries of the column we'll be traversing. + uint column = get_global_id(0) % kAtlasWidth, shelfIndex = get_global_id(0) / kAtlasWidth; + uint firstRow = shelfIndex * kAtlasShelfHeight, lastRow = (shelfIndex + 1) * kAtlasShelfHeight; // Sweep down the column, accumulating coverage as we go. - float coverage = 0; - for (uint y = yStart; y < yEnd; y++) { - coverage += as_float(*getPixel(gPixels, (uint2)(x, y), widthInTiles)); + float coverage = 0.0f; + for (uint row = firstRow; row < lastRow; row++) { + int2 coord = (int2)((int)column, (int)row); + coverage += read_imagef(gCoverage, SAMPLER, coord).r; - uint grayscaleValue = 255 - convert_uint(clamp(coverage, 0.0f, 255.0f)); - write_imageui(gTexture, (int2)((int)x, (int)y), (uint4)(grayscaleValue, 255, 255, 255)); + uint gray = 255 - convert_uint(clamp(coverage, 0.0f, 1.0f) * 255.0f); + write_imageui(gTexture, coord, (uint4)(gray, 255, 255, 255)); } } diff --git a/resources/shaders/draw.cl b/resources/shaders/draw.cl deleted file mode 100644 index 355d06ba..00000000 --- a/resources/shaders/draw.cl +++ /dev/null @@ -1,241 +0,0 @@ -// Copyright 2017 The Servo Project Developers. See the COPYRIGHT -// file at the top-level directory of this distribution and at -// http://rust-lang.org/COPYRIGHT. -// -// Licensed under the Apache License, Version 2.0 or the MIT license -// , at your -// option. This file may not be copied, modified, or distributed -// except according to those terms. - -// Computes exact area coverage for lines, breaking Bézier curves down into them as necessary. -// Proceeds top to bottom for better data locality during the subsequent accumulation stage. For -// details on the algorithm, see [1]. -// -// [1]: http://nothings.org/gamedev/rasterize/ - -#define POINTS_PER_SEGMENT 32 -#define TILE_SIZE 4 - -#define OPERATION_MOVE 0 -#define OPERATION_ON_CURVE 1 -#define OPERATION_OFF_CURVE 2 - -struct GlyphDescriptor { - short4 rect; - ushort unitsPerEm; - ushort pointCount; - uint startPoint; -}; - -typedef struct GlyphDescriptor GlyphDescriptor; - -struct ImageDescriptor { - uint2 atlasPosition; - float pointSize; - uint glyphIndex; - uint startPointInBatch; - uint pointCount; -}; - -typedef struct ImageDescriptor ImageDescriptor; - -__global int *getPixel(__global int *gPixels, uint2 point, uint widthInTiles) { - uint2 tile = point / TILE_SIZE, pointInTile = point % TILE_SIZE; - return &gPixels[(tile.y * widthInTiles + tile.x) * TILE_SIZE * TILE_SIZE + - pointInTile.y * TILE_SIZE + - pointInTile.x]; -} - -uchar getOperation(uint globalPointIndex, __global const uchar *gOperations) { - return (gOperations[globalPointIndex / 4] >> (globalPointIndex % 4 * 2)) & 0x3; -} - -void plot(__global int *gPixels, - uint2 point, - uint widthInTiles, - uint imageHeight, - float coverage) { - __global int *pixel = getPixel(gPixels, - (uint2)(point.x, imageHeight - point.y - 1), - widthInTiles); - int oldCoverage = as_int(*pixel); - while (true) { - int newCoverage = as_int(as_float(oldCoverage) + coverage); - int existingCoverage = atomic_cmpxchg(pixel, oldCoverage, newCoverage); - if (existingCoverage == oldCoverage) - break; - oldCoverage = existingCoverage; - } -} - -__kernel void draw(__global const ImageDescriptor *gImages, - __global const GlyphDescriptor *gGlyphs, - __global const short2 *gCoordinates, - __global const uchar *gOperations, - __global const uint *gIndices, - __global int *gPixels, - uint atlasWidth) { - // Find the image. - int batchID = get_global_id(0); - uint imageID = gIndices[batchID / POINTS_PER_SEGMENT]; - __global const ImageDescriptor *image = &gImages[imageID]; - while (batchID >= image->startPointInBatch + image->pointCount) { - imageID++; - image = &gImages[imageID]; - } - - // Find the glyph. - uint glyphIndex = image->glyphIndex; - __global const GlyphDescriptor *glyph = &gGlyphs[glyphIndex]; - - // Unpack glyph and image. - uint2 atlasPosition = image->atlasPosition; - float pixelsPerUnit = image->pointSize * convert_float(glyph->unitsPerEm); - uint pointIndexInGlyph = batchID - image->startPointInBatch; - uint globalPointIndex = glyph->startPoint + pointIndexInGlyph; - - // Stop here if this is a move operation. - uchar curOperation = getOperation(globalPointIndex, gOperations); - if (curOperation == OPERATION_MOVE) - return; - - // Unpack the points that make up this line or curve. - short2 p0, p1, p2; - float t0, t1; - uchar prevOperation = getOperation(globalPointIndex - 1, gOperations); - short2 prevPoint = gCoordinates[globalPointIndex - 1]; - short2 curPoint = gCoordinates[globalPointIndex]; - if (prevOperation == OPERATION_OFF_CURVE) { - p0 = gCoordinates[globalPointIndex - 2]; - p1 = prevPoint; - p2 = curPoint; - t0 = 0.0f; - t1 = 0.5f; - } else if (curOperation == OPERATION_OFF_CURVE) { - p0 = prevPoint; - p1 = curPoint; - p2 = gCoordinates[globalPointIndex + 1]; - t0 = 0.5f; - t1 = 1.0f; - } else { - p0 = prevPoint; - p2 = curPoint; - } - - // Convert units to pixels. - float2 pP0 = convert_float2(p0) * pixelsPerUnit; - float2 pP1 = convert_float2(p1) * pixelsPerUnit; - float2 pP2 = convert_float2(p2) * pixelsPerUnit; - - // Determine the direction we're going. - float2 direction = copysign((float2)(1.0f, 1.0f), pP0 - pP2); - - // Set up plotting. - uint widthInTiles = atlasWidth / TILE_SIZE; - short4 glyphRect = glyph->rect; - uint imageHeight = convert_uint(glyphRect.w - glyphRect.y); - - // Loop over each line segment. - float t = t0; - while (t < t1) { - // Compute endpoints. - float2 lP0, lP1; - if (direction.x >= 0.0f) { - lP0 = pP0; - lP1 = pP2; - } else { - lP0 = pP2; - lP1 = pP0; - } - - // Compute the slope. - float dXdY = fabs(lP1.x - lP0.x / lP1.y - lP0.y); - - // Initialize the current point. Determine how long the segment extends across the first - // pixel column. - int2 p = (int2)((int)p0.x, 0); - float dX = min(convert_float(p.x) + 1.0f, lP1.x) - lP0.x; - - // Initialize `yLeft` and `yRight`, the intercepts of Y with the current pixel. - float yLeft = lP0.y; - float yRight = yLeft + direction.y * dX / dXdY; - - // Iterate over columns. - while (p.x < (int)ceil(lP1.x)) { - // Flip `yLeft` and `yRight` around if necessary so that the slope is positive. - float y0, y1; - if (yLeft <= yRight) { - y0 = yLeft; - y1 = yRight; - } else { - y0 = yRight; - y1 = yLeft; - } - - // Split `y0` into fractional and whole parts, and split `y1` into remaining fractional - // and whole parts. - float y0R, y1R; - float y0F = fract(y0, &y0R), y1F = fract(y1, &y1R); - int y0I = convert_int(y0R), y1I = convert_int(y1R); - if (y1F != 0.0f) - y1I++; - - // Compute area coverage for the first pixel. - float coverage; - if (y1I <= y0I + 1) { - // The line is less than one pixel. This is a trapezoid. - coverage = 1.0f - mix(y0F, y1F, 0.5f); - } else { - // Usual case: This is a triangle. - coverage = 0.5f * dXdY * (1.0f - y0F) * (1.0f - y0F); - } - - // Plot the first pixel of this column. - plot(gPixels, as_uint2(p), widthInTiles, imageHeight, dX * direction.x * coverage); - - // Since the coverage of this row must sum to 1, we keep track of the total coverage. - float coverageLeft = coverage; - - // Plot the pixels between the first and the last. - if (p.y + 1 < y1I) { - // Compute coverage for and plot the second pixel in the column. - p.y++; - if (p.y + 1 == y1I) - coverage = 1.0f - (0.5f * dXdY * y1F * y1F) - coverage; - else - coverage = dXdY * (1.5f - y0F) - coverage; - coverageLeft += coverage; - plot(gPixels, as_uint2(p), widthInTiles, imageHeight, dX * direction.x * coverage); - - // Iterate over any remaining pixels. - p.y++; - coverage = dXdY; - while (p.y < y1I) { - coverageLeft += coverage; - plot(gPixels, - as_uint2(p), - widthInTiles, - imageHeight, - dX * direction.x * coverage); - p.y++; - } - } - - // Plot the remaining coverage. - coverage = 1.0f - coverageLeft; - plot(gPixels, as_uint2(p), widthInTiles, imageHeight, dX * direction.x * coverage); - - // Move to the next column. - p.x++; - - // Compute Y intercepts for the next column. - yLeft = yRight; - float yRight = yLeft + direction.y * dX / dXdY; - - // Determine how long the segment extends across the next pixel column. - dX = min(convert_float(p.x) + 1.0f, lP1.x) - convert_float(p.x); - } - } -} - diff --git a/resources/shaders/draw.fs.glsl b/resources/shaders/draw.fs.glsl new file mode 100644 index 00000000..2aa9d1ce --- /dev/null +++ b/resources/shaders/draw.fs.glsl @@ -0,0 +1,116 @@ +// Copyright 2017 The Servo Project Developers. See the COPYRIGHT +// file at the top-level directory of this distribution and at +// http://rust-lang.org/COPYRIGHT. +// +// Licensed under the Apache License, Version 2.0 or the MIT license +// , at your +// option. This file may not be copied, modified, or distributed +// except according to those terms. + +#version 410 + +// The size of the atlas in pixels. +uniform uvec2 uAtlasSize; + +// The starting point of the segment. +flat in vec2 vP0; +// The endpoint of this segment. +flat in vec2 vP1; +// 1.0 if this segment runs left to right; -1.0 otherwise. +flat in float vDirection; +// The slope of this line. +flat in float vSlope; +// Minimum and maximum vertical extents, unrounded. +flat in vec2 vYMinMax; + +out vec4 oFragColor; + +void main() { + // Compute the X boundaries of this pixel. + float xMin = floor(gl_FragCoord.x), xMax = ceil(gl_FragCoord.x); + + // Compute the horizontal span that the line segment covers across this pixel. + float dX = min(xMax, vP1.x) - max(xMin, vP0.x); + + // Compute the Y-intercepts of the portion of the line crossing this pixel. + float yMin = clamp(vP0.y + (xMin - vP0.x) * vSlope, vYMinMax.x, vYMinMax.y); + float yMax = clamp(yMin + vSlope, vYMinMax.x, vYMinMax.y); + if (yMin > yMax) { + float tmp = yMin; + yMin = yMax; + yMax = tmp; + } + + // Round the Y-intercepts out to the nearest pixel. + int yMinI = int(floor(yMin)), yMaxI = int(ceil(yMax)); + + // Determine which vertical pixel we're looking at. + int yI = int(floor(gl_FragCoord.y)); + + // Compute trapezoidal area coverage. + // + // It may be helpful to follow along with this explanation, keeping in mind that we compute + // downward coverage rather than rightward coverage: + // + // http://nothings.org/gamedev/rasterize/ + // + // Note that the algorithm above computes total area coverage for each pixel, while here we + // compute *delta* coverage: that is, the *difference* in the area covered between this pixel + // and the pixel above it. In general, this means that, in contrast to the stb_truetype + // algorithm, we have to specially handle the first fully covered pixel, in order to account + // for the residual area difference between that pixel and the one above it. + float coverage = 0.0f; + if (yMaxI <= yMinI + 1) { + // The line touches only one pixel (case 1). Compute the area of that trapezoid (or the + // residual area for the pixel right after that trapezoid). + float trapArea = 0.5f * (yMin + yMax) - float(yMinI); + if (yI == yMinI) + coverage = 1.0f - trapArea; + else if (yI == yMinI + 1) + coverage = trapArea; + } else { + // The line touches multiple pixels (case 2). There are several subcases to handle here. + + // Compute the area of the topmost triangle. + float yMinF = fract(yMin); + float dXdY = 1.0f / (yMax - yMin); + float triAreaMin = 0.5f * dXdY * (1.0f - yMinF) * (1.0f - yMinF); + + if (yI == yMinI) { + // We're looking at the pixel that triangle covers, so we're done. + coverage = triAreaMin; + } else { + // Compute the area of the bottommost triangle. + float yMaxF = yMax - ceil(yMax) + 1.0f; + float triAreaMax = 0.5f * dXdY * yMaxF * yMaxF; + + bool lineTouchesThreePixels = yMaxI == yMinI + 2; + if (lineTouchesThreePixels && yI == yMinI + 1) { + // The line touches exactly 3 pixels, and we're looking at the middle one. + coverage = 1.0f - triAreaMin - triAreaMax; + } else if (!lineTouchesThreePixels && yI < yMaxI) { + // The line touches more than 3 pixels. Compute the area of the first trapezoid. + float trapAreaMin = dXdY * (1.5f - yMinF); + if (yI == yMinI + 1) { + // We're looking at that trapezoid, so we're done. + coverage = trapAreaMin - triAreaMin; + } else if (yI == yMaxI - 1) { + // We're looking at the last trapezoid. Compute its area. + float trapAreaMax = trapAreaMin + float(yMaxI - yMinI - 3) * dXdY; + coverage = 1.0f - trapAreaMax - triAreaMax; + } else if (yI > yMinI + 1 && yI < yMaxI - 1) { + // We're looking at one of the pixels in between the two trapezoids. The delta + // coverage in this case is simply the inverse slope. + coverage = dXdY; + } + } else if (yI == yMaxI) { + // We're looking at the final pixel in the column. + coverage = triAreaMax; + } + } + } + + oFragColor = vec4(dX * vDirection * coverage, 1.0f, 1.0f, 1.0f); +} + diff --git a/resources/shaders/draw.tcs.glsl b/resources/shaders/draw.tcs.glsl new file mode 100644 index 00000000..b97936bb --- /dev/null +++ b/resources/shaders/draw.tcs.glsl @@ -0,0 +1,112 @@ +// Copyright 2017 The Servo Project Developers. See the COPYRIGHT +// file at the top-level directory of this distribution and at +// http://rust-lang.org/COPYRIGHT. +// +// Licensed under the Apache License, Version 2.0 or the MIT license +// , at your +// option. This file may not be copied, modified, or distributed +// except according to those terms. + +#version 410 + +#define OPERATION_LINE 1 +#define OPERATION_QUAD_CURVE 2 + +#define CURVE_THRESHOLD 0.333f +#define CURVE_TOLERANCE 3.0f + +layout(vertices = 1) out; + +// The size of the atlas in pixels. +uniform uvec2 uAtlasSize; + +// The vertex ID, passed into this shader. +flat in uint vVertexID[]; + +// The starting point of the segment. +patch out vec2 vpP0; +// The control point, if this is a curve. If this is a line, this value must be ignored. +patch out vec2 vpP1; +// The endpoint of this segment. +patch out vec2 vpP2; +// 1.0 if this segment runs left to right; -1.0 otherwise. +patch out float vpDirection; + +void main() { + vpP0 = gl_in[0].gl_Position.xy; + vpP1 = gl_in[1].gl_Position.xy; + vpP2 = gl_in[2].gl_Position.xy; + + // Compute direction. Flip around if necessary so that p0 is to the left of p2. + if (vpP0.x < vpP2.x) { + vpDirection = 1.0f; + } else { + vpDirection = -1.0f; + vec2 tmp = vpP0; + vpP0 = vpP2; + vpP2 = tmp; + } + + // Divide into lines. + float lineCount = 1.0f; + if (vVertexID[1] > 0) { + // Quadratic curve. + vec2 dev = vpP0 - 2.0f * vpP1 + vpP2; + float devSq = dot(dev, dev); + if (devSq >= QUAD_CURVE_THRESHOLD) { + // Inverse square root is likely no slower and may be faster than regular square root + // (e.g. on x86). + lineCount += floor(inversesqrt(inversesqrt(QUAD_CURVE_TOLERANCE * devSq))); + } + } + + // Tessellate into lines. This is subtle, so a diagram may help. + // + // Suppose we decided to divide this curve up into 4 lines. Then our abstract tessellated patch + // space will look like this: + // + // x₀ x₁ x₂ x₃ x₄ x₅ x₆ x₇ + // ┌──┬──┬──┬──┬──┬──┬──┐ + // │▒▒│ │▒▒│ │▒▒│ │▒▒│ + // │▒▒│ │▒▒│ │▒▒│ │▒▒│ + // └──┴──┴──┴──┴──┴──┴──┘ + // + // The shaded areas are the only areas that will actually be drawn. They might look like this: + // + // x₅ + // x₆ x₇ + // x₃ ┌───────┐ + // x₄ │▒▒▒▒▒▒▒│ + // x₁ ┌─────┼───────┘ + // x₂ │▒▒▒▒▒│ + // ┌──┼─────┘ + // │▒▒│ + // │▒▒│ + // x₀ │▒▒│ + // ┌──┼──┘ + // │▒▒│ + // │▒▒│ + // └──┘ + // + // In this way, the unshaded areas become zero-size and are discarded by the rasterizer. + // + // Note that, in reality, it will often be the case that the quads overlap vertically by one + // pixel in the horizontal direction. In fact, this will occur whenever a line segment endpoint + // does not precisely straddle a pixel boundary. However, observe that we can guarantee that + // x₂ ≤ x₁, x₄ ≤ x₃, and so on, because there is never any horizontal space between endpoints. + // This means that all triangles inside the unshaded areas are guaranteed to be wound in the + // opposite direction from those inside the shaded areas. Because the OpenGL spec guarantees + // that, by default, all tessellated triangles are wound counterclockwise in abstract patch + // space, the triangles within the unshaded areas must be wound clockwise and are therefore + // candidates for backface culling. Backface culling is always enabled when running Pathfinder, + // so we're in the clear: the rasterizer will always discard the unshaded areas and render only + // the shaded ones. + + gl_TessLevelInner[0] = vpP0.x == vpP2.x ? 0.0f : lineCount * 2.0f - 1.0f; + gl_TessLevelOuter[1] = gl_TessLevelOuter[3] = gl_TessLevelInner[0]; + + // Don't split vertically at all. We only tessellate horizontally. + gl_TessLevelInner[1] = gl_TessLevelOuter[0] = gl_TessLevelOuter[2] = 1.0f; +} + diff --git a/resources/shaders/draw.tes.glsl b/resources/shaders/draw.tes.glsl new file mode 100644 index 00000000..cd4524b5 --- /dev/null +++ b/resources/shaders/draw.tes.glsl @@ -0,0 +1,67 @@ +// Copyright 2017 The Servo Project Developers. See the COPYRIGHT +// file at the top-level directory of this distribution and at +// http://rust-lang.org/COPYRIGHT. +// +// Licensed under the Apache License, Version 2.0 or the MIT license +// , at your +// option. This file may not be copied, modified, or distributed +// except according to those terms. + +#version 410 + +layout(quads) in; + +// The size of the atlas in pixels. +uniform uvec2 uAtlasSize; + +// The starting point of the segment. +patch in vec2 vpP0; +// The control point, if this is a curve. If this is a line, this value must be ignored. +patch in vec2 vpP1; +// The endpoint of this segment. +patch in vec2 vpP2; +// 1.0 if this segment runs left to right; -1.0 otherwise. +patch in float vpDirection; + +// The starting point of the segment. +flat out vec2 vP0; +// The endpoint of this segment. +flat out vec2 vP1; +// 1.0 if this segment runs left to right; -1.0 otherwise. +flat out float vDirection; +// The slope of this line. +flat out float vSlope; +// Minimum and maximum vertical extents, unrounded. +flat out vec2 vYMinMax; + +void main() { + // Work out how many lines made up this segment, which line we're working on, and which + // endpoint of that line we're looking at. + uint tessPointCount = uint(gl_TessLevelInner[0] + 1.0f); + uint tessIndex = uint(round(gl_TessCoord.x * float(tessPointCount - 1))) + uint lineCount = tessPointCount / 2, lineIndex = tessIndex / 2, endpoint = tessIndex % 2; + + // Compute our endpoints (trivial if this is part of a line, less trivial if this is part of a + // curve). + if (lineCount == 1) { + vP0 = vpP0; + vP1 = vpP2; + } else { + float t0 = float(lineIndex + 0) / float(lineCount); + float t1 = float(lineIndex + 1) / float(lineCount); + vP0 = mix(mix(vpP0, vpP1, t0), mix(vpP1, vpP2, t0), t0); + vP1 = mix(mix(vpP0, vpP1, t1), mix(vpP1, vpP2, t0), t1); + } + + // Compute Y extents. + vYMinMax = vP0.y <= vP1.y ? vec2(vP0.y, vP1.y) : vec2(vP1.y, vP0.y); + + // Compute our final position in atlas space, rounded out to the next pixel. + float x = pointIndex == 0 ? floor(vP0.x) : ceil(vP1.x); + float y = gl_TessCoord.y == 0.0f ? floor(vYMinMax.x) : ceil(vYMinMax.y) + 1.0f; + + // Convert atlas space to device space. + gl_Position = vec4(vec2(x, y) / vec2(uAtlasSize) * 2.0f - 1.0f, 0.0f, 1.0f); +} + diff --git a/resources/shaders/draw.vs.glsl b/resources/shaders/draw.vs.glsl new file mode 100644 index 00000000..b8d6366c --- /dev/null +++ b/resources/shaders/draw.vs.glsl @@ -0,0 +1,49 @@ +// Copyright 2017 The Servo Project Developers. See the COPYRIGHT +// file at the top-level directory of this distribution and at +// http://rust-lang.org/COPYRIGHT. +// +// Licensed under the Apache License, Version 2.0 or the MIT license +// , at your +// option. This file may not be copied, modified, or distributed +// except according to those terms. + +#version 410 + +// Information about the position of each glyph in the atlas. +layout(std140) struct ImageInfo { + // The left/top/right/bottom positions of the glyph in the atlas. + uvec4 atlasRect; + // The left/top/right/bottom offsets of the glyph from point (0, 0) in glyph space. + ivec4 extents; + // The font size in pixels. + float pointSize; +} + +// The size of the atlas in pixels. +uniform uvec2 uAtlasSize; + +layout(std140) uniform ubImageInfo { + ImageInfo uImageInfo[256]; +}; + +// The position of each vertex in glyph space. +in ivec2 aPosition; + +// Which image the vertex belongs to. +// +// TODO(pcwalton): See if this is faster as a binary search on the vertex ID. +in uint aImageIndex; + +// The vertex ID, passed along onto the TCS. +flat out uint vVertexID; + +void main() { + vVertexID = gl_VertexID; + + ImageInfo imageInfo = uImageInfo[aImageIndex]; + vec2 glyphPos = vec2(aPosition.x - imageInfo.extents.x, imageInfo.extents.w - aPosition.y); + vec2 atlasPos = glyphPos * EMS_PER_UNIT * pointSize + vec2(imageInfo.atlasRect.xy); + gl_Position = vec4(atlasPos, 0.0f, 1.0f); +} +