diff --git a/Cargo.toml b/Cargo.toml index 33e4a74d..b3954497 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -19,3 +19,6 @@ quickcheck = "0.4" [dev-dependencies.glfw] git = "https://github.com/bjz/glfw-rs.git" +[dev-dependencies.lord-drawquaad] +git = "https://github.com/pcwalton/lord-drawquaad.git" + diff --git a/examples/generate-atlas.rs b/examples/generate-atlas.rs index 86a5ab50..3bde1951 100644 --- a/examples/generate-atlas.rs +++ b/examples/generate-atlas.rs @@ -2,16 +2,23 @@ * http://creativecommons.org/publicdomain/zero/1.0/ */ extern crate compute_shader; +extern crate euclid; extern crate gl; extern crate glfw; +extern crate lord_drawquaad; extern crate memmap; extern crate pathfinder; +use compute_shader::buffer; use compute_shader::instance::Instance; -use glfw::{Context, OpenGlProfileHint, WindowHint, WindowMode}; +use compute_shader::texture::ExternalTexture; +use euclid::{Point2D, Rect, Size2D}; +use gl::types::GLint; +use glfw::{Action, Context, Key, OpenGlProfileHint, WindowEvent, WindowHint, WindowMode}; use memmap::{Mmap, Protection}; use pathfinder::batch::{BatchBuilder, GlyphRange}; use pathfinder::charmap::CodepointRange; +use pathfinder::coverage::CoverageBuffer; use pathfinder::glyph_buffer::GlyphBufferBuilder; use pathfinder::otf::FontData; use pathfinder::rasterizer::Rasterizer; @@ -30,7 +37,7 @@ fn main() { glfw.window_hint(WindowHint::OpenGlProfile(OpenGlProfileHint::Core)); let context = glfw.create_window(WIDTH, HEIGHT, "generate-atlas", WindowMode::Windowed); - let (mut window, _) = context.expect("Couldn't create a window!"); + let (mut window, events) = context.expect("Couldn't create a window!"); window.make_current(); gl::load_with(|symbol| window.get_proc_address(symbol) as *const c_void); @@ -59,8 +66,55 @@ fn main() { } } - let glyph_buffer = glyph_buffer_builder.finish(&rasterizer.device).unwrap(); + let glyph_buffers = glyph_buffer_builder.finish(&rasterizer.device).unwrap(); let batch = batch_builder.finish(&rasterizer.device).unwrap(); - // TODO(pcwalton): ... + + let atlas_size = Size2D::new(WIDTH, HEIGHT); + let coverage_buffer = CoverageBuffer::new(&rasterizer.device, &atlas_size).unwrap(); + + let texture = rasterizer.device + .create_texture(buffer::Protection::WriteOnly, &atlas_size) + .unwrap(); + + rasterizer.draw_atlas(&Rect::new(Point2D::new(0, 0), atlas_size), + SHELF_HEIGHT, + &glyph_buffers, + &batch, + &coverage_buffer, + &texture).unwrap().wait().unwrap(); + + let draw_context = lord_drawquaad::Context::new(); + + let mut gl_texture = 0; + unsafe { + gl::GenTextures(1, &mut gl_texture); + texture.bind_to(&ExternalTexture::Gl(gl_texture)).unwrap(); + + gl::BindTexture(gl::TEXTURE_RECTANGLE, gl_texture); + gl::TexParameteri(gl::TEXTURE_RECTANGLE, gl::TEXTURE_MIN_FILTER, gl::LINEAR as GLint); + gl::TexParameteri(gl::TEXTURE_RECTANGLE, gl::TEXTURE_MAG_FILTER, gl::LINEAR as GLint); + gl::TexParameteri(gl::TEXTURE_RECTANGLE, gl::TEXTURE_WRAP_S, gl::CLAMP_TO_EDGE as GLint); + gl::TexParameteri(gl::TEXTURE_RECTANGLE, gl::TEXTURE_WRAP_T, gl::CLAMP_TO_EDGE as GLint); + } + + unsafe { + gl::ClearColor(1.0, 1.0, 1.0, 1.0); + gl::Clear(gl::COLOR_BUFFER_BIT); + } + + draw_context.draw(gl_texture); + window.swap_buffers(); + + while !window.should_close() { + glfw.poll_events(); + for (_, event) in glfw::flush_messages(&events) { + match event { + WindowEvent::Key(Key::Escape, _, Action::Press, _) => { + window.set_should_close(true) + } + _ => {} + } + } + } } diff --git a/resources/shaders/draw.cl b/resources/shaders/draw.cl index ece87dfe..355d06ba 100644 --- a/resources/shaders/draw.cl +++ b/resources/shaders/draw.cl @@ -51,8 +51,14 @@ 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, float coverage) { - __global int *pixel = getPixel(gPixels, point, widthInTiles); +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); @@ -68,7 +74,8 @@ __kernel void draw(__global const ImageDescriptor *gImages, __global const short2 *gCoordinates, __global const uchar *gOperations, __global const uint *gIndices, - __global int *gPixels) { + __global int *gPixels, + uint atlasWidth) { // Find the image. int batchID = get_global_id(0); uint imageID = gIndices[batchID / POINTS_PER_SEGMENT]; @@ -115,5 +122,120 @@ __kernel void draw(__global const ImageDescriptor *gImages, 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/src/batch.rs b/src/batch.rs index 83db8648..4b063d40 100644 --- a/src/batch.rs +++ b/src/batch.rs @@ -71,6 +71,7 @@ impl BatchBuilder { Ok(Batch { indices: try!(device.create_buffer(Protection::ReadOnly, indices).map_err(drop)), images: try!(device.create_buffer(Protection::ReadOnly, images).map_err(drop)), + point_count: self.point_count, }) } } @@ -78,6 +79,7 @@ impl BatchBuilder { pub struct Batch { pub indices: Buffer, pub images: Buffer, + pub point_count: u32, } #[derive(Clone, Copy, Debug)] diff --git a/src/coverage.rs b/src/coverage.rs new file mode 100644 index 00000000..f88106a4 --- /dev/null +++ b/src/coverage.rs @@ -0,0 +1,30 @@ +// 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. + +use compute_shader::buffer::{Buffer, BufferData, Protection}; +use compute_shader::device::Device; +use euclid::size::Size2D; +use std::mem; + +pub struct CoverageBuffer { + pub buffer: Buffer, +} + +impl CoverageBuffer { + pub fn new(device: &Device, size: &Size2D) -> Result { + let size = size.width as usize * size.height as usize * mem::size_of::(); + let buffer = try!(device.create_buffer(Protection::ReadWrite, + BufferData::Uninitialized(size)).map_err(drop)); + Ok(CoverageBuffer { + buffer: buffer, + }) + } +} + diff --git a/src/lib.rs b/src/lib.rs index 890153df..17c4ae9a 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -26,6 +26,7 @@ extern crate test; pub mod atlas; pub mod batch; pub mod charmap; +pub mod coverage; pub mod glyph_buffer; pub mod otf; pub mod rasterizer; diff --git a/src/rasterizer.rs b/src/rasterizer.rs index 6f260eb7..84ff0162 100644 --- a/src/rasterizer.rs +++ b/src/rasterizer.rs @@ -8,9 +8,15 @@ // option. This file may not be copied, modified, or distributed // except according to those terms. +use batch::Batch; use compute_shader::device::Device; +use compute_shader::event::Event; use compute_shader::program::Program; -use compute_shader::queue::Queue; +use compute_shader::queue::{Queue, Uniform}; +use compute_shader::texture::Texture; +use coverage::CoverageBuffer; +use euclid::rect::Rect; +use glyph_buffer::GlyphBuffers; // TODO(pcwalton): Don't force that these be compiled in. // TODO(pcwalton): GLSL version. @@ -37,5 +43,50 @@ impl Rasterizer { draw_program: draw_program, }) } + + pub fn draw_atlas(&self, + atlas_rect: &Rect, + atlas_shelf_height: u32, + glyph_buffers: &GlyphBuffers, + batch: &Batch, + coverage_buffer: &CoverageBuffer, + texture: &Texture) + -> Result { + let draw_uniforms = [ + (0, Uniform::Buffer(&batch.images)), + (1, Uniform::Buffer(&glyph_buffers.descriptors)), + (2, Uniform::Buffer(&glyph_buffers.coordinates)), + (3, Uniform::Buffer(&glyph_buffers.operations)), + (4, Uniform::Buffer(&batch.indices)), + (5, Uniform::Buffer(&coverage_buffer.buffer)), + (6, Uniform::U32(try!(texture.width().map_err(drop)))), + ]; + + let draw_event = try!(self.queue.submit_compute(&self.draw_program, + &[batch.point_count], + &draw_uniforms, + &[]).map_err(drop)); + + let atlas_rect_uniform = [ + atlas_rect.origin.x, + atlas_rect.origin.y, + atlas_rect.max_x(), + atlas_rect.max_y() + ]; + + let accum_uniforms = [ + (0, Uniform::Buffer(&coverage_buffer.buffer)), + (1, Uniform::Texture(texture)), + (2, Uniform::UVec4(atlas_rect_uniform)), + (3, Uniform::U32(atlas_shelf_height)), + ]; + + let accum_columns = atlas_rect.size.width * (atlas_rect.size.height / atlas_shelf_height); + + self.queue.submit_compute(&self.accum_program, + &[accum_columns], + &accum_uniforms, + &[draw_event]).map_err(drop) + } }