diff --git a/examples/generate-atlas.rs b/examples/generate-atlas.rs index 3bde1951..76761fa0 100644 --- a/examples/generate-atlas.rs +++ b/examples/generate-atlas.rs @@ -11,7 +11,7 @@ extern crate pathfinder; use compute_shader::buffer; use compute_shader::instance::Instance; -use compute_shader::texture::ExternalTexture; +use compute_shader::texture::{ExternalTexture, Format}; use euclid::{Point2D, Rect, Size2D}; use gl::types::GLint; use glfw::{Action, Context, Key, OpenGlProfileHint, WindowEvent, WindowHint, WindowMode}; @@ -73,7 +73,7 @@ fn main() { let coverage_buffer = CoverageBuffer::new(&rasterizer.device, &atlas_size).unwrap(); let texture = rasterizer.device - .create_texture(buffer::Protection::WriteOnly, &atlas_size) + .create_texture(Format::R8, buffer::Protection::WriteOnly, &atlas_size) .unwrap(); rasterizer.draw_atlas(&Rect::new(Point2D::new(0, 0), atlas_size), diff --git a/resources/shaders/accum.cl b/resources/shaders/accum.cl index 001f2bd1..a6022900 100644 --- a/resources/shaders/accum.cl +++ b/resources/shaders/accum.cl @@ -18,10 +18,11 @@ const sampler_t SAMPLER = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_F __kernel void accum(__write_only image2d_t gTexture, __read_only image2d_t gCoverage, - uint kAtlasWidth, + uint4 kAtlasRect, uint 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 atlasWidth = kAtlasRect.z - kAtlasRect.x; + uint column = get_global_id(0) % atlasWidth, shelfIndex = get_global_id(0) / atlasWidth; uint firstRow = shelfIndex * kAtlasShelfHeight, lastRow = (shelfIndex + 1) * kAtlasShelfHeight; // Sweep down the column, accumulating coverage as we go. @@ -31,7 +32,7 @@ __kernel void accum(__write_only image2d_t gTexture, coverage += read_imagef(gCoverage, SAMPLER, coord).r; uint gray = 255 - convert_uint(clamp(coverage, 0.0f, 1.0f) * 255.0f); - write_imageui(gTexture, coord, (uint4)(gray, 255, 255, 255)); + write_imageui(gTexture, coord + (int2)kAtlasRect.xy, (uint4)(gray, 255, 255, 255)); } } diff --git a/resources/shaders/draw.tcs.glsl b/resources/shaders/draw.tcs.glsl index b97936bb..6a69321e 100644 --- a/resources/shaders/draw.tcs.glsl +++ b/resources/shaders/draw.tcs.glsl @@ -10,9 +10,6 @@ #version 410 -#define OPERATION_LINE 1 -#define OPERATION_QUAD_CURVE 2 - #define CURVE_THRESHOLD 0.333f #define CURVE_TOLERANCE 3.0f @@ -54,10 +51,10 @@ void main() { // Quadratic curve. vec2 dev = vpP0 - 2.0f * vpP1 + vpP2; float devSq = dot(dev, dev); - if (devSq >= QUAD_CURVE_THRESHOLD) { + if (devSq >= 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))); + lineCount += floor(inversesqrt(inversesqrt(CURVE_TOLERANCE * devSq))); } } diff --git a/resources/shaders/draw.tes.glsl b/resources/shaders/draw.tes.glsl index cd4524b5..6b50546f 100644 --- a/resources/shaders/draw.tes.glsl +++ b/resources/shaders/draw.tes.glsl @@ -39,7 +39,7 @@ 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 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 @@ -54,11 +54,15 @@ void main() { vP1 = mix(mix(vpP0, vpP1, t1), mix(vpP1, vpP2, t0), t1); } - // Compute Y extents. + // 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); + + // Forward direction onto the fragment shader. + vDirection = vpDirection; // Compute our final position in atlas space, rounded out to the next pixel. - float x = pointIndex == 0 ? floor(vP0.x) : ceil(vP1.x); + float x = endpoint == 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. diff --git a/resources/shaders/draw.vs.glsl b/resources/shaders/draw.vs.glsl index b8d6366c..ec551807 100644 --- a/resources/shaders/draw.vs.glsl +++ b/resources/shaders/draw.vs.glsl @@ -18,11 +18,14 @@ layout(std140) struct ImageInfo { ivec4 extents; // The font size in pixels. float pointSize; -} +}; // The size of the atlas in pixels. uniform uvec2 uAtlasSize; +// The number of ems per unit (reciprocal of units per em). +uniform float uEmsPerUnit; + layout(std140) uniform ubImageInfo { ImageInfo uImageInfo[256]; }; @@ -43,7 +46,7 @@ void main() { 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); + vec2 atlasPos = glyphPos * uEmsPerUnit * imageInfo.pointSize + vec2(imageInfo.atlasRect.xy); gl_Position = vec4(atlasPos, 0.0f, 1.0f); } diff --git a/src/coverage.rs b/src/coverage.rs index f88106a4..9be4c571 100644 --- a/src/coverage.rs +++ b/src/coverage.rs @@ -8,22 +8,22 @@ // 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::buffer::Protection; use compute_shader::device::Device; +use compute_shader::texture::{Format, Texture}; use euclid::size::Size2D; use std::mem; pub struct CoverageBuffer { - pub buffer: Buffer, + pub texture: Texture, } 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)); + let texture = try!(device.create_texture(Format::R32F, Protection::ReadWrite, size) + .map_err(drop)); Ok(CoverageBuffer { - buffer: buffer, + texture: texture, }) } } diff --git a/src/lib.rs b/src/lib.rs index 17c4ae9a..224ded2f 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -15,6 +15,7 @@ extern crate bitflags; extern crate byteorder; extern crate compute_shader; extern crate euclid; +extern crate gl; #[cfg(test)] extern crate memmap; #[cfg(test)] diff --git a/src/rasterizer.rs b/src/rasterizer.rs index 84ff0162..1b8e42d9 100644 --- a/src/rasterizer.rs +++ b/src/rasterizer.rs @@ -16,31 +16,69 @@ use compute_shader::queue::{Queue, Uniform}; use compute_shader::texture::Texture; use coverage::CoverageBuffer; use euclid::rect::Rect; +use gl::types::{GLchar, GLenum, GLint, GLsizei, GLuint}; +use gl; use glyph_buffer::GlyphBuffers; +use std::ptr; // TODO(pcwalton): Don't force that these be compiled in. // TODO(pcwalton): GLSL version. static ACCUM_CL_SHADER: &'static str = include_str!("../resources/shaders/accum.cl"); -static DRAW_CL_SHADER: &'static str = include_str!("../resources/shaders/draw.cl"); + +static DRAW_VERTEX_SHADER: &'static str = include_str!("../resources/shaders/draw.vs.glsl"); +static DRAW_TESS_CONTROL_SHADER: &'static str = include_str!("../resources/shaders/draw.tcs.glsl"); +static DRAW_TESS_EVALUATION_SHADER: &'static str = + include_str!("../resources/shaders/draw.tes.glsl"); +static DRAW_FRAGMENT_SHADER: &'static str = include_str!("../resources/shaders/draw.fs.glsl"); pub struct Rasterizer { pub device: Device, pub queue: Queue, + draw_program: GLuint, accum_program: Program, - draw_program: Program, } impl Rasterizer { pub fn new(device: Device, queue: Queue) -> Result { - // TODO(pcwalton): GLSL version. - // FIXME(pcwalton): Don't panic if these fail to compile; just return an error. + let draw_program; + unsafe { + let shaders = [ + try!(compile_gl_shader(gl::VERTEX_SHADER, + "Vertex shader", + DRAW_VERTEX_SHADER)), + try!(compile_gl_shader(gl::TESS_CONTROL_SHADER, + "Tessellation control shader", + DRAW_TESS_CONTROL_SHADER)), + try!(compile_gl_shader(gl::TESS_EVALUATION_SHADER, + "Tessellation evaluation shader", + DRAW_TESS_EVALUATION_SHADER)), + try!(compile_gl_shader(gl::FRAGMENT_SHADER, + "Fragment shader", + DRAW_FRAGMENT_SHADER)), + ]; + + draw_program = gl::CreateProgram(); + for &shader in &shaders { + gl::AttachShader(draw_program, shader); + } + + gl::LinkProgram(draw_program); + + try!(check_gl_object_status(draw_program, + gl::LINK_STATUS, + "Program", + gl::GetProgramiv, + gl::GetProgramInfoLog)) + } + + // FIXME(pcwalton): Don't panic if this fails to compile; just return an error. let accum_program = device.create_program(ACCUM_CL_SHADER).unwrap(); - let draw_program = device.create_program(DRAW_CL_SHADER).unwrap(); + Ok(Rasterizer { device: device, queue: queue, - accum_program: accum_program, draw_program: draw_program, + accum_program: accum_program, }) } @@ -52,20 +90,7 @@ impl Rasterizer { 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)); + // TODO(pcwalton) let atlas_rect_uniform = [ atlas_rect.origin.x, @@ -75,8 +100,8 @@ impl Rasterizer { ]; let accum_uniforms = [ - (0, Uniform::Buffer(&coverage_buffer.buffer)), - (1, Uniform::Texture(texture)), + (0, Uniform::Texture(texture)), + (1, Uniform::Texture(&coverage_buffer.texture)), (2, Uniform::UVec4(atlas_rect_uniform)), (3, Uniform::U32(atlas_shelf_height)), ]; @@ -86,7 +111,46 @@ impl Rasterizer { self.queue.submit_compute(&self.accum_program, &[accum_columns], &accum_uniforms, - &[draw_event]).map_err(drop) + &[]).map_err(drop) + } +} + +fn compile_gl_shader(shader_type: GLuint, description: &str, source: &str) -> Result { + unsafe { + let shader = gl::CreateShader(shader_type); + gl::ShaderSource(shader, 1, &(source.as_ptr() as *const GLchar), &(source.len() as GLint)); + gl::CompileShader(shader); + try!(check_gl_object_status(shader, + gl::COMPILE_STATUS, + description, + gl::GetShaderiv, + gl::GetShaderInfoLog)); + Ok(shader) + } +} + +fn check_gl_object_status(object: GLuint, + parameter: GLenum, + description: &str, + get_status: unsafe fn(GLuint, GLenum, *mut GLint), + get_log: unsafe fn(GLuint, GLsizei, *mut GLsizei, *mut GLchar)) + -> Result<(), ()> { + unsafe { + let mut status = 0; + get_status(object, parameter, &mut status); + if status == gl::TRUE as i32 { + return Ok(()) + } + + let mut info_log_length = 0; + get_status(object, gl::INFO_LOG_LENGTH, &mut info_log_length); + + let mut info_log = vec![0; info_log_length as usize]; + get_log(object, info_log_length, ptr::null_mut(), info_log.as_mut_ptr() as *mut GLchar); + if let Ok(string) = String::from_utf8(info_log) { + println!("{} error:\n{}", description, string); + } + Err(()) } }