Fix shader compilation errors; convert the coverage buffer to a texture
This commit is contained in:
parent
f1ec3385de
commit
0569831969
|
@ -11,7 +11,7 @@ extern crate pathfinder;
|
||||||
|
|
||||||
use compute_shader::buffer;
|
use compute_shader::buffer;
|
||||||
use compute_shader::instance::Instance;
|
use compute_shader::instance::Instance;
|
||||||
use compute_shader::texture::ExternalTexture;
|
use compute_shader::texture::{ExternalTexture, Format};
|
||||||
use euclid::{Point2D, Rect, Size2D};
|
use euclid::{Point2D, Rect, Size2D};
|
||||||
use gl::types::GLint;
|
use gl::types::GLint;
|
||||||
use glfw::{Action, Context, Key, OpenGlProfileHint, WindowEvent, WindowHint, WindowMode};
|
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 coverage_buffer = CoverageBuffer::new(&rasterizer.device, &atlas_size).unwrap();
|
||||||
|
|
||||||
let texture = rasterizer.device
|
let texture = rasterizer.device
|
||||||
.create_texture(buffer::Protection::WriteOnly, &atlas_size)
|
.create_texture(Format::R8, buffer::Protection::WriteOnly, &atlas_size)
|
||||||
.unwrap();
|
.unwrap();
|
||||||
|
|
||||||
rasterizer.draw_atlas(&Rect::new(Point2D::new(0, 0), atlas_size),
|
rasterizer.draw_atlas(&Rect::new(Point2D::new(0, 0), atlas_size),
|
||||||
|
|
|
@ -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,
|
__kernel void accum(__write_only image2d_t gTexture,
|
||||||
__read_only image2d_t gCoverage,
|
__read_only image2d_t gCoverage,
|
||||||
uint kAtlasWidth,
|
uint4 kAtlasRect,
|
||||||
uint kAtlasShelfHeight) {
|
uint kAtlasShelfHeight) {
|
||||||
// Determine the boundaries of the column we'll be traversing.
|
// 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;
|
uint firstRow = shelfIndex * kAtlasShelfHeight, lastRow = (shelfIndex + 1) * kAtlasShelfHeight;
|
||||||
|
|
||||||
// Sweep down the column, accumulating coverage as we go.
|
// 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;
|
coverage += read_imagef(gCoverage, SAMPLER, coord).r;
|
||||||
|
|
||||||
uint gray = 255 - convert_uint(clamp(coverage, 0.0f, 1.0f) * 255.0f);
|
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));
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
@ -10,9 +10,6 @@
|
||||||
|
|
||||||
#version 410
|
#version 410
|
||||||
|
|
||||||
#define OPERATION_LINE 1
|
|
||||||
#define OPERATION_QUAD_CURVE 2
|
|
||||||
|
|
||||||
#define CURVE_THRESHOLD 0.333f
|
#define CURVE_THRESHOLD 0.333f
|
||||||
#define CURVE_TOLERANCE 3.0f
|
#define CURVE_TOLERANCE 3.0f
|
||||||
|
|
||||||
|
@ -54,10 +51,10 @@ void main() {
|
||||||
// Quadratic curve.
|
// Quadratic curve.
|
||||||
vec2 dev = vpP0 - 2.0f * vpP1 + vpP2;
|
vec2 dev = vpP0 - 2.0f * vpP1 + vpP2;
|
||||||
float devSq = dot(dev, dev);
|
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
|
// Inverse square root is likely no slower and may be faster than regular square root
|
||||||
// (e.g. on x86).
|
// (e.g. on x86).
|
||||||
lineCount += floor(inversesqrt(inversesqrt(QUAD_CURVE_TOLERANCE * devSq)));
|
lineCount += floor(inversesqrt(inversesqrt(CURVE_TOLERANCE * devSq)));
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
@ -39,7 +39,7 @@ void main() {
|
||||||
// Work out how many lines made up this segment, which line we're working on, and which
|
// 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.
|
// endpoint of that line we're looking at.
|
||||||
uint tessPointCount = uint(gl_TessLevelInner[0] + 1.0f);
|
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;
|
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
|
// 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);
|
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);
|
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.
|
// 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;
|
float y = gl_TessCoord.y == 0.0f ? floor(vYMinMax.x) : ceil(vYMinMax.y) + 1.0f;
|
||||||
|
|
||||||
// Convert atlas space to device space.
|
// Convert atlas space to device space.
|
||||||
|
|
|
@ -18,11 +18,14 @@ layout(std140) struct ImageInfo {
|
||||||
ivec4 extents;
|
ivec4 extents;
|
||||||
// The font size in pixels.
|
// The font size in pixels.
|
||||||
float pointSize;
|
float pointSize;
|
||||||
}
|
};
|
||||||
|
|
||||||
// The size of the atlas in pixels.
|
// The size of the atlas in pixels.
|
||||||
uniform uvec2 uAtlasSize;
|
uniform uvec2 uAtlasSize;
|
||||||
|
|
||||||
|
// The number of ems per unit (reciprocal of units per em).
|
||||||
|
uniform float uEmsPerUnit;
|
||||||
|
|
||||||
layout(std140) uniform ubImageInfo {
|
layout(std140) uniform ubImageInfo {
|
||||||
ImageInfo uImageInfo[256];
|
ImageInfo uImageInfo[256];
|
||||||
};
|
};
|
||||||
|
@ -43,7 +46,7 @@ void main() {
|
||||||
|
|
||||||
ImageInfo imageInfo = uImageInfo[aImageIndex];
|
ImageInfo imageInfo = uImageInfo[aImageIndex];
|
||||||
vec2 glyphPos = vec2(aPosition.x - imageInfo.extents.x, imageInfo.extents.w - aPosition.y);
|
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);
|
gl_Position = vec4(atlasPos, 0.0f, 1.0f);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
@ -8,22 +8,22 @@
|
||||||
// option. This file may not be copied, modified, or distributed
|
// option. This file may not be copied, modified, or distributed
|
||||||
// except according to those terms.
|
// 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::device::Device;
|
||||||
|
use compute_shader::texture::{Format, Texture};
|
||||||
use euclid::size::Size2D;
|
use euclid::size::Size2D;
|
||||||
use std::mem;
|
use std::mem;
|
||||||
|
|
||||||
pub struct CoverageBuffer {
|
pub struct CoverageBuffer {
|
||||||
pub buffer: Buffer,
|
pub texture: Texture,
|
||||||
}
|
}
|
||||||
|
|
||||||
impl CoverageBuffer {
|
impl CoverageBuffer {
|
||||||
pub fn new(device: &Device, size: &Size2D<u32>) -> Result<CoverageBuffer, ()> {
|
pub fn new(device: &Device, size: &Size2D<u32>) -> Result<CoverageBuffer, ()> {
|
||||||
let size = size.width as usize * size.height as usize * mem::size_of::<u32>();
|
let texture = try!(device.create_texture(Format::R32F, Protection::ReadWrite, size)
|
||||||
let buffer = try!(device.create_buffer(Protection::ReadWrite,
|
.map_err(drop));
|
||||||
BufferData::Uninitialized(size)).map_err(drop));
|
|
||||||
Ok(CoverageBuffer {
|
Ok(CoverageBuffer {
|
||||||
buffer: buffer,
|
texture: texture,
|
||||||
})
|
})
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
|
@ -15,6 +15,7 @@ extern crate bitflags;
|
||||||
extern crate byteorder;
|
extern crate byteorder;
|
||||||
extern crate compute_shader;
|
extern crate compute_shader;
|
||||||
extern crate euclid;
|
extern crate euclid;
|
||||||
|
extern crate gl;
|
||||||
#[cfg(test)]
|
#[cfg(test)]
|
||||||
extern crate memmap;
|
extern crate memmap;
|
||||||
#[cfg(test)]
|
#[cfg(test)]
|
||||||
|
|
|
@ -16,31 +16,69 @@ use compute_shader::queue::{Queue, Uniform};
|
||||||
use compute_shader::texture::Texture;
|
use compute_shader::texture::Texture;
|
||||||
use coverage::CoverageBuffer;
|
use coverage::CoverageBuffer;
|
||||||
use euclid::rect::Rect;
|
use euclid::rect::Rect;
|
||||||
|
use gl::types::{GLchar, GLenum, GLint, GLsizei, GLuint};
|
||||||
|
use gl;
|
||||||
use glyph_buffer::GlyphBuffers;
|
use glyph_buffer::GlyphBuffers;
|
||||||
|
use std::ptr;
|
||||||
|
|
||||||
// TODO(pcwalton): Don't force that these be compiled in.
|
// TODO(pcwalton): Don't force that these be compiled in.
|
||||||
// TODO(pcwalton): GLSL version.
|
// TODO(pcwalton): GLSL version.
|
||||||
static ACCUM_CL_SHADER: &'static str = include_str!("../resources/shaders/accum.cl");
|
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 struct Rasterizer {
|
||||||
pub device: Device,
|
pub device: Device,
|
||||||
pub queue: Queue,
|
pub queue: Queue,
|
||||||
|
draw_program: GLuint,
|
||||||
accum_program: Program,
|
accum_program: Program,
|
||||||
draw_program: Program,
|
|
||||||
}
|
}
|
||||||
|
|
||||||
impl Rasterizer {
|
impl Rasterizer {
|
||||||
pub fn new(device: Device, queue: Queue) -> Result<Rasterizer, ()> {
|
pub fn new(device: Device, queue: Queue) -> Result<Rasterizer, ()> {
|
||||||
// TODO(pcwalton): GLSL version.
|
let draw_program;
|
||||||
// FIXME(pcwalton): Don't panic if these fail to compile; just return an error.
|
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 accum_program = device.create_program(ACCUM_CL_SHADER).unwrap();
|
||||||
let draw_program = device.create_program(DRAW_CL_SHADER).unwrap();
|
|
||||||
Ok(Rasterizer {
|
Ok(Rasterizer {
|
||||||
device: device,
|
device: device,
|
||||||
queue: queue,
|
queue: queue,
|
||||||
accum_program: accum_program,
|
|
||||||
draw_program: draw_program,
|
draw_program: draw_program,
|
||||||
|
accum_program: accum_program,
|
||||||
})
|
})
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -52,20 +90,7 @@ impl Rasterizer {
|
||||||
coverage_buffer: &CoverageBuffer,
|
coverage_buffer: &CoverageBuffer,
|
||||||
texture: &Texture)
|
texture: &Texture)
|
||||||
-> Result<Event, ()> {
|
-> Result<Event, ()> {
|
||||||
let draw_uniforms = [
|
// TODO(pcwalton)
|
||||||
(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 = [
|
let atlas_rect_uniform = [
|
||||||
atlas_rect.origin.x,
|
atlas_rect.origin.x,
|
||||||
|
@ -75,8 +100,8 @@ impl Rasterizer {
|
||||||
];
|
];
|
||||||
|
|
||||||
let accum_uniforms = [
|
let accum_uniforms = [
|
||||||
(0, Uniform::Buffer(&coverage_buffer.buffer)),
|
(0, Uniform::Texture(texture)),
|
||||||
(1, Uniform::Texture(texture)),
|
(1, Uniform::Texture(&coverage_buffer.texture)),
|
||||||
(2, Uniform::UVec4(atlas_rect_uniform)),
|
(2, Uniform::UVec4(atlas_rect_uniform)),
|
||||||
(3, Uniform::U32(atlas_shelf_height)),
|
(3, Uniform::U32(atlas_shelf_height)),
|
||||||
];
|
];
|
||||||
|
@ -86,7 +111,46 @@ impl Rasterizer {
|
||||||
self.queue.submit_compute(&self.accum_program,
|
self.queue.submit_compute(&self.accum_program,
|
||||||
&[accum_columns],
|
&[accum_columns],
|
||||||
&accum_uniforms,
|
&accum_uniforms,
|
||||||
&[draw_event]).map_err(drop)
|
&[]).map_err(drop)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
fn compile_gl_shader(shader_type: GLuint, description: &str, source: &str) -> Result<GLuint, ()> {
|
||||||
|
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(())
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
Loading…
Reference in New Issue