Draft shader code for the new tessellation-based algorithm
This commit is contained in:
parent
bdf6ebab24
commit
f1ec3385de
|
@ -12,39 +12,26 @@
|
||||||
//
|
//
|
||||||
// This proceeds top to bottom for better data locality. For details on the algorithm, see [1].
|
// 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) {
|
__kernel void accum(__write_only image2d_t gTexture,
|
||||||
uint2 tile = point / TILE_SIZE, pointInTile = point % TILE_SIZE;
|
__read_only image2d_t gCoverage,
|
||||||
return &gPixels[(tile.y * widthInTiles + tile.x) * TILE_SIZE * TILE_SIZE +
|
uint kAtlasWidth,
|
||||||
pointInTile.y * TILE_SIZE +
|
|
||||||
pointInTile.x];
|
|
||||||
}
|
|
||||||
|
|
||||||
__kernel void accum(__global const int *gPixels,
|
|
||||||
__write_only image2d_t gTexture,
|
|
||||||
uint4 kAtlasRect,
|
|
||||||
uint kAtlasShelfHeight) {
|
uint kAtlasShelfHeight) {
|
||||||
// Compute our column.
|
// Determine the boundaries of the column we'll be traversing.
|
||||||
int globalID = get_global_id(0);
|
uint column = get_global_id(0) % kAtlasWidth, shelfIndex = get_global_id(0) / kAtlasWidth;
|
||||||
uint atlasWidth = kAtlasRect.z - kAtlasRect.x;
|
uint firstRow = shelfIndex * kAtlasShelfHeight, lastRow = (shelfIndex + 1) * kAtlasShelfHeight;
|
||||||
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;
|
|
||||||
|
|
||||||
// Sweep down the column, accumulating coverage as we go.
|
// Sweep down the column, accumulating coverage as we go.
|
||||||
float coverage = 0;
|
float coverage = 0.0f;
|
||||||
for (uint y = yStart; y < yEnd; y++) {
|
for (uint row = firstRow; row < lastRow; row++) {
|
||||||
coverage += as_float(*getPixel(gPixels, (uint2)(x, y), widthInTiles));
|
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));
|
uint gray = 255 - convert_uint(clamp(coverage, 0.0f, 1.0f) * 255.0f);
|
||||||
write_imageui(gTexture, (int2)((int)x, (int)y), (uint4)(grayscaleValue, 255, 255, 255));
|
write_imageui(gTexture, coord, (uint4)(gray, 255, 255, 255));
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
@ -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 <LICENSE-APACHE or
|
|
||||||
// http://www.apache.org/licenses/LICENSE-2.0> or the MIT license
|
|
||||||
// <LICENSE-MIT or http://opensource.org/licenses/MIT>, 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);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
|
@ -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 <LICENSE-APACHE or
|
||||||
|
// http://www.apache.org/licenses/LICENSE-2.0> or the MIT license
|
||||||
|
// <LICENSE-MIT or http://opensource.org/licenses/MIT>, 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);
|
||||||
|
}
|
||||||
|
|
|
@ -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 <LICENSE-APACHE or
|
||||||
|
// http://www.apache.org/licenses/LICENSE-2.0> or the MIT license
|
||||||
|
// <LICENSE-MIT or http://opensource.org/licenses/MIT>, 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;
|
||||||
|
}
|
||||||
|
|
|
@ -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 <LICENSE-APACHE or
|
||||||
|
// http://www.apache.org/licenses/LICENSE-2.0> or the MIT license
|
||||||
|
// <LICENSE-MIT or http://opensource.org/licenses/MIT>, 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);
|
||||||
|
}
|
||||||
|
|
|
@ -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 <LICENSE-APACHE or
|
||||||
|
// http://www.apache.org/licenses/LICENSE-2.0> or the MIT license
|
||||||
|
// <LICENSE-MIT or http://opensource.org/licenses/MIT>, 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);
|
||||||
|
}
|
||||||
|
|
Loading…
Reference in New Issue