Move direction calculation from the TCS to the TES to properly handle cusps

resources/shaders/accum.cl

@@ -21,9 +21,10 @@ __kernel void accum(__write_only image2d_t gImage,
                     uint4 kAtlasRect,
                     uint kAtlasShelfHeight) {
     // Determine the boundaries of the column we'll be traversing.
-    uint atlasWidth = kAtlasRect.z - kAtlasRect.x;
+    uint atlasWidth = kAtlasRect.z - kAtlasRect.x, atlasHeight = kAtlasRect.w - kAtlasRect.y;
     uint column = get_global_id(0) % atlasWidth, shelfIndex = get_global_id(0) / atlasWidth;
-    uint firstRow = shelfIndex * kAtlasShelfHeight, lastRow = (shelfIndex + 1) * kAtlasShelfHeight;
+    uint firstRow = min(shelfIndex * kAtlasShelfHeight, atlasHeight);
+    uint lastRow = min((shelfIndex + 1) * kAtlasShelfHeight, atlasHeight);
 
     // Sweep down the column, accumulating coverage as we go.
     float coverage = 0.0f;

resources/shaders/draw.tcs.glsl

@@ -18,39 +18,25 @@ layout(vertices = 1) out;
 // The vertex ID, passed into this shader.
 flat in int vVertexID[];
 
+// These outputs really should be patch outs, but that causes problems in Apple drivers.
+
 // The starting point of the segment.
-patch out vec2 vpP0;
+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;
+out vec2 vpP1[];
 // The endpoint of this segment.
-patch out vec2 vpP2;
-// x: 1.0 if this segment runs left to right; -1.0 otherwise.
-// y: The tessellation level.
+out vec2 vpP2[];
+// The tessellation level.
 //
-// This is packed together into a single vec2 to work around an Apple Intel driver bug whereby
-// patch outputs beyond the first 4 are forced to 0.
-//
-// And in case you're wondering why the tessellation level is passed along in a patch out instead
-// of having the TES read it directly, that's another Apple bug workaround, this time in the Radeon
-// driver.
-patch out vec2 vpDirectionTessLevel;
+// This is passed along explicitly instead of having the TES read it from `gl_TessLevelInner` in
+// order to work around an Apple bug in the Radeon driver.
+out float vpTessLevel[];
 
 void main() {
     vec2 p0 = gl_in[0].gl_Position.xy;
     vec2 p1 = gl_in[1].gl_Position.xy;
     vec2 p2 = gl_in[2].gl_Position.xy;
 
-    // Compute direction. Flip around if necessary so that p0 is to the left of p2.
-    float direction;
-    if (p0.x < p2.x) {
-        direction = 1.0f;
-    } else {
-        direction = -1.0f;
-        vec2 tmp = p0;
-        p0 = p2;
-        p2 = tmp;
-    }
-
     // Divide into lines.
     float lineCount = 1.0f;
     if (vVertexID[1] > 0) {
@@ -116,9 +102,9 @@ void main() {
 
     // NB: These per-patch outputs must be assigned in this order, or Apple's compiler will
     // miscompile us.
-    vpP0 = p0;
-    vpP1 = p1;
-    vpP2 = p2;
-    vpDirectionTessLevel = vec2(direction, tessLevel);
+    vpP0[gl_InvocationID] = p0;
+    vpP1[gl_InvocationID] = p1;
+    vpP2[gl_InvocationID] = p2;
+    vpTessLevel[gl_InvocationID] = tessLevel;
 }
 

resources/shaders/draw.tes.glsl

@@ -16,21 +16,16 @@ layout(quads) in;
 uniform uvec2 uAtlasSize;
 
 // The starting point of the segment.
-patch in vec2 vpP0;
+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;
+in vec2 vpP1[];
 // The endpoint of this segment.
-patch in vec2 vpP2;
-// x: 1.0 if this segment runs left to right; -1.0 otherwise.
-// y: The tessellation level.
+in vec2 vpP2[];
+// The tessellation level.
 //
-// This is packed together into a single vec2 to work around an Apple Intel driver bug whereby
-// patch outputs beyond the first 4 are forced to 0.
-//
-// And in case you're wondering why the tessellation level is passed along in a patch out instead
-// of having the TES read it directly, that's another Apple bug workaround, this time in the Radeon
-// driver.
-patch in vec2 vpDirectionTessLevel;
+// This is passed along explicitly instead of having the TES read it from `gl_TessLevelInner` in
+// order to work around an Apple bug in the Radeon driver.
+in float vpTessLevel[];
 
 // The starting point of the segment.
 flat out vec2 vP0;
@@ -44,33 +39,51 @@ flat out float vSlope;
 flat out vec2 vYMinMax;
 
 void main() {
+    // Read in curve points.
+    vec2 cP0 = vpP0[0], cP1 = vpP1[0], cP2 = vpP2[0];
+
     // 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(vpDirectionTessLevel.y + 1.0f);
+    uint tessPointCount = uint(vpTessLevel[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).
+    vec2 p0, p1;
     if (lineCount == 1) {
-        vP0 = vpP0;
-        vP1 = vpP2;
+        p0 = cP0;
+        p1 = cP2;
     } 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, t1), t1);
+        p0 = mix(mix(cP0, cP1, t0), mix(cP1, cP2, t0), t0);
+        p1 = mix(mix(cP0, cP1, t1), mix(cP1, cP2, t1), t1);
     }
 
-    // Compute Y extents and slope.
-    vYMinMax = vP0.y <= vP1.y ? vec2(vP0.y, vP1.y) : vec2(vP1.y, vP0.y);
-    vSlope = (vP1.y - vP0.y) / (vP1.x - vP0.x);
+    // Compute direction. Flip the two points around so that p0 is on the left and p1 is on the
+    // right if necessary.
+    float direction;
+    if (p0.x < p1.x) {
+        direction = 1.0f;
+    } else {
+        direction = -1.0f;
+        vec2 tmp = p0;
+        p0 = p1;
+        p1 = tmp;
+    }
 
-    // Forward direction onto the fragment shader.
-    vDirection = vpDirectionTessLevel.x;
+    // Forward points and direction onto the fragment shader.
+    vP0 = p0;
+    vP1 = p1;
+    vDirection = direction;
+
+    // Compute Y extents and slope.
+    vSlope = (p1.y - p0.y) / (p1.x - p0.x);
+    vYMinMax = p0.y <= p1.y ? vec2(p0.y, p1.y) : vec2(p1.y, p0.y);
 
     // Compute our final position in atlas space, rounded out to the next pixel.
-    float x = endpoint == 0 ? floor(vP0.x) : ceil(vP1.x);
+    float x = endpoint == 0 ? floor(p0.x) : ceil(p1.x);
     float y = gl_TessCoord.y == 0.0f ? floor(vYMinMax.x) : ceil(vYMinMax.y) + 1.0f;
 
     // Convert atlas space to device space.