From f00a64999c197e6a96e65fd00f64224a6f22c9fa Mon Sep 17 00:00:00 2001 From: Eric Anholt Date: Fri, 16 Nov 2007 16:43:45 -0800 Subject: [intel] Add 965 support to shared intel_blit.c This requires that regions grow a marker of whether they are tiled or not, because fence (surface) registers are ignored by the 965 2D engine. --- src/mesa/drivers/dri/i915/Makefile | 2 +- src/mesa/drivers/dri/i915/intel_pixel_copy.c | 4 +- src/mesa/drivers/dri/i915/intel_pixel_draw.c | 6 +- src/mesa/drivers/dri/i915/intel_pixel_read.c | 5 +- src/mesa/drivers/dri/intel/intel_blit.c | 91 +++++++++++++++++++--------- src/mesa/drivers/dri/intel/intel_blit.h | 3 + src/mesa/drivers/dri/intel/intel_regions.c | 55 ++++++++--------- src/mesa/drivers/dri/intel/intel_regions.h | 6 +- src/mesa/drivers/dri/intel/intel_screen.c | 9 ++- src/mesa/drivers/dri/intel/intel_screen.h | 6 ++ src/mesa/drivers/dri/intel/intel_tex_copy.c | 3 +- src/mesa/drivers/dri/intel/intel_tex_image.c | 4 +- 12 files changed, 119 insertions(+), 75 deletions(-) diff --git a/src/mesa/drivers/dri/i915/Makefile b/src/mesa/drivers/dri/i915/Makefile index bff09ee8424..7ef055ccad2 100644 --- a/src/mesa/drivers/dri/i915/Makefile +++ b/src/mesa/drivers/dri/i915/Makefile @@ -61,7 +61,7 @@ C_SOURCES = \ ASM_SOURCES = -DRIVER_DEFINES = -I../intel -I../intel/server \ +DRIVER_DEFINES = -I../intel -I../intel/server -DI915 \ $(shell pkg-config libdrm --atleast-version=2.3.1 \ && echo "-DDRM_VBLANK_FLIP=DRM_VBLANK_FLIP") diff --git a/src/mesa/drivers/dri/i915/intel_pixel_copy.c b/src/mesa/drivers/dri/i915/intel_pixel_copy.c index 9d478283e47..629cdb979d9 100644 --- a/src/mesa/drivers/dri/i915/intel_pixel_copy.c +++ b/src/mesa/drivers/dri/i915/intel_pixel_copy.c @@ -342,8 +342,8 @@ do_blit_copypixels(GLcontext * ctx, intelEmitCopyBlit(intel, dst->cpp, - src->pitch, src->buffer, 0, - dst->pitch, dst->buffer, 0, + src->pitch, src->buffer, 0, src->tiled, + dst->pitch, dst->buffer, 0, dst->tiled, rect.x1 + delta_x, rect.y1 + delta_y, /* srcx, srcy */ rect.x1, rect.y1, /* dstx, dsty */ diff --git a/src/mesa/drivers/dri/i915/intel_pixel_draw.c b/src/mesa/drivers/dri/i915/intel_pixel_draw.c index afb586b4a33..8349f4c7484 100644 --- a/src/mesa/drivers/dri/i915/intel_pixel_draw.c +++ b/src/mesa/drivers/dri/i915/intel_pixel_draw.c @@ -312,10 +312,8 @@ do_blit_drawpixels(GLcontext * ctx, intelEmitCopyBlit(intel, dest->cpp, - rowLength, - src_buffer, src_offset, - dest->pitch, - dest->buffer, 0, + rowLength, src_buffer, src_offset, GL_FALSE, + dest->pitch, dest->buffer, 0, dest->tiled, rect.x1 - dest_rect.x1, rect.y2 - dest_rect.y2, rect.x1, diff --git a/src/mesa/drivers/dri/i915/intel_pixel_read.c b/src/mesa/drivers/dri/i915/intel_pixel_read.c index a22844926c9..2e31656e984 100644 --- a/src/mesa/drivers/dri/i915/intel_pixel_read.c +++ b/src/mesa/drivers/dri/i915/intel_pixel_read.c @@ -264,9 +264,8 @@ do_blit_readpixels(GLcontext * ctx, intelEmitCopyBlit(intel, src->cpp, - src->pitch, src->buffer, 0, - rowLength, - dst_buffer, dst_offset, + src->pitch, src->buffer, 0, src->tiled, + rowLength, dst_buffer, dst_offset, GL_FALSE, rect.x1, rect.y1, rect.x1 - src_rect.x1, diff --git a/src/mesa/drivers/dri/intel/intel_blit.c b/src/mesa/drivers/dri/intel/intel_blit.c index 2761136f47f..09b45e43a12 100644 --- a/src/mesa/drivers/dri/intel/intel_blit.c +++ b/src/mesa/drivers/dri/intel/intel_blit.c @@ -86,8 +86,9 @@ intelCopyBuffer(const __DRIdrawablePrivate * dPriv, = intel_get_rb_region(&intel_fb->Base, BUFFER_BACK_LEFT); const int nbox = dPriv->numClipRects; const drm_clip_rect_t *pbox = dPriv->pClipRects; - const int pitch = frontRegion->pitch; const int cpp = frontRegion->cpp; + int src_pitch = backRegion->pitch * cpp; + int dst_pitch = frontRegion->pitch * cpp; int BR13, CMD; int i; @@ -99,14 +100,25 @@ intelCopyBuffer(const __DRIdrawablePrivate * dPriv, ASSERT(frontRegion->cpp == backRegion->cpp); if (cpp == 2) { - BR13 = (pitch * cpp) | (0xCC << 16) | (1 << 24); + BR13 = (0xCC << 16) | (1 << 24); CMD = XY_SRC_COPY_BLT_CMD; } else { - BR13 = (pitch * cpp) | (0xCC << 16) | (1 << 24) | (1 << 25); - CMD = (XY_SRC_COPY_BLT_CMD | XY_BLT_WRITE_ALPHA | XY_BLT_WRITE_RGB); + BR13 = (0xCC << 16) | (1 << 24) | (1 << 25); + CMD = XY_SRC_COPY_BLT_CMD | XY_BLT_WRITE_ALPHA | XY_BLT_WRITE_RGB; } +#ifndef I915 + if (backRegion->tiled) { + CMD |= XY_SRC_TILED; + src_pitch /= 4; + } + if (frontRegion->tiled) { + CMD |= XY_DST_TILED; + dst_pitch /= 4; + } +#endif + for (i = 0; i < nbox; i++, pbox++) { drm_clip_rect_t box; @@ -133,14 +145,14 @@ intelCopyBuffer(const __DRIdrawablePrivate * dPriv, BEGIN_BATCH(8, INTEL_BATCH_NO_CLIPRECTS); OUT_BATCH(CMD); - OUT_BATCH(BR13); + OUT_BATCH(BR13 | dst_pitch); OUT_BATCH((pbox->y1 << 16) | pbox->x1); OUT_BATCH((pbox->y2 << 16) | pbox->x2); OUT_RELOC(frontRegion->buffer, DRM_BO_FLAG_MEM_TT | DRM_BO_FLAG_WRITE, 0); OUT_BATCH((pbox->y1 << 16) | pbox->x1); - OUT_BATCH(BR13 & 0xffff); + OUT_BATCH(src_pitch); OUT_RELOC(backRegion->buffer, DRM_BO_FLAG_MEM_TT | DRM_BO_FLAG_READ, 0); @@ -166,6 +178,7 @@ intelEmitFillBlit(struct intel_context *intel, GLshort dst_pitch, dri_bo *dst_buffer, GLuint dst_offset, + GLboolean dst_tiled, GLshort x, GLshort y, GLshort w, GLshort h, GLuint color) { GLuint BR13, CMD; @@ -177,16 +190,22 @@ intelEmitFillBlit(struct intel_context *intel, case 1: case 2: case 3: - BR13 = dst_pitch | (0xF0 << 16) | (1 << 24); + BR13 = (0xF0 << 16) | (1 << 24); CMD = XY_COLOR_BLT_CMD; break; case 4: - BR13 = dst_pitch | (0xF0 << 16) | (1 << 24) | (1 << 25); - CMD = (XY_COLOR_BLT_CMD | XY_BLT_WRITE_ALPHA | XY_BLT_WRITE_RGB); + BR13 = (0xF0 << 16) | (1 << 24) | (1 << 25); + CMD = XY_COLOR_BLT_CMD | XY_BLT_WRITE_ALPHA | XY_BLT_WRITE_RGB; break; default: return; } +#ifndef I915 + if (dst_tiled) { + CMD |= XY_DST_TILED; + dst_pitch /= 4; + } +#endif DBG("%s dst:buf(%p)/%d+%d %d,%d sz:%dx%d\n", __FUNCTION__, dst_buffer, dst_pitch, dst_offset, x, y, w, h); @@ -194,7 +213,7 @@ intelEmitFillBlit(struct intel_context *intel, BEGIN_BATCH(6, INTEL_BATCH_NO_CLIPRECTS); OUT_BATCH(CMD); - OUT_BATCH(BR13); + OUT_BATCH(BR13 | dst_pitch); OUT_BATCH((y << 16) | x); OUT_BATCH(((y + h) << 16) | (x + w)); OUT_RELOC(dst_buffer, DRM_BO_FLAG_MEM_TT | DRM_BO_FLAG_WRITE, dst_offset); @@ -235,9 +254,11 @@ intelEmitCopyBlit(struct intel_context *intel, GLshort src_pitch, dri_bo *src_buffer, GLuint src_offset, + GLboolean src_tiled, GLshort dst_pitch, dri_bo *dst_buffer, GLuint dst_offset, + GLboolean dst_tiled, GLshort src_x, GLshort src_y, GLshort dst_x, GLshort dst_y, GLshort w, GLshort h, @@ -257,25 +278,34 @@ intelEmitCopyBlit(struct intel_context *intel, src_pitch *= cpp; dst_pitch *= cpp; + BR13 = (translate_raster_op(logic_op) << 16); + switch (cpp) { case 1: case 2: case 3: - BR13 = (((GLint) dst_pitch) & 0xffff) | - (translate_raster_op(logic_op) << 16) | (1 << 24); + BR13 |= (1 << 24); CMD = XY_SRC_COPY_BLT_CMD; break; case 4: - BR13 = - (((GLint) dst_pitch) & 0xffff) | - (translate_raster_op(logic_op) << 16) | (1 << 24) | (1 << 25); - CMD = - (XY_SRC_COPY_BLT_CMD | XY_BLT_WRITE_ALPHA | XY_BLT_WRITE_RGB); + BR13 |= (1 << 24) | (1 << 25); + CMD = XY_SRC_COPY_BLT_CMD | XY_BLT_WRITE_ALPHA | XY_BLT_WRITE_RGB; break; default: return; } +#ifndef I915 + if (dst_tiled) { + CMD |= XY_DST_TILED; + dst_pitch /= 4; + } + if (src_tiled) { + CMD |= XY_SRC_TILED; + src_pitch /= 4; + } +#endif + if (dst_y2 < dst_y || dst_x2 < dst_x) { return; } @@ -290,25 +320,25 @@ intelEmitCopyBlit(struct intel_context *intel, if (dst_pitch > 0 && src_pitch > 0) { BEGIN_BATCH(8, INTEL_BATCH_NO_CLIPRECTS); OUT_BATCH(CMD); - OUT_BATCH(BR13); + OUT_BATCH(BR13 | dst_pitch); OUT_BATCH((dst_y << 16) | dst_x); OUT_BATCH((dst_y2 << 16) | dst_x2); OUT_RELOC(dst_buffer, DRM_BO_FLAG_MEM_TT | DRM_BO_FLAG_WRITE, dst_offset); OUT_BATCH((src_y << 16) | src_x); - OUT_BATCH(((GLint) src_pitch & 0xffff)); + OUT_BATCH(src_pitch); OUT_RELOC(src_buffer, DRM_BO_FLAG_MEM_TT | DRM_BO_FLAG_READ, src_offset); ADVANCE_BATCH(); } else { BEGIN_BATCH(8, INTEL_BATCH_NO_CLIPRECTS); OUT_BATCH(CMD); - OUT_BATCH(BR13); + OUT_BATCH(BR13 | dst_pitch); OUT_BATCH((0 << 16) | dst_x); OUT_BATCH((h << 16) | dst_x2); OUT_RELOC(dst_buffer, DRM_BO_FLAG_MEM_TT | DRM_BO_FLAG_WRITE, dst_offset + dst_y * dst_pitch); OUT_BATCH((0 << 16) | src_x); - OUT_BATCH(((GLint) src_pitch & 0xffff)); + OUT_BATCH(src_pitch); OUT_RELOC(src_buffer, DRM_BO_FLAG_MEM_TT | DRM_BO_FLAG_READ, src_offset + src_y * src_pitch); ADVANCE_BATCH(); @@ -435,12 +465,13 @@ intelClearWithBlit(GLcontext * ctx, GLbitfield mask) irb_region->draw_offset, b.x1, b.y1, b.x2 - b.x1, b.y2 - b.y1); + BR13 = 0xf0 << 16; + CMD = XY_COLOR_BLT_CMD; /* Setup the blit command */ if (cpp == 4) { - BR13 = (0xF0 << 16) | (pitch * cpp) | (1 << 24) | (1 << 25); + BR13 |= (1 << 24) | (1 << 25); if (buf == BUFFER_DEPTH || buf == BUFFER_STENCIL) { - CMD = XY_COLOR_BLT_CMD; if (clearMask & BUFFER_BIT_DEPTH) CMD |= XY_BLT_WRITE_RGB; if (clearMask & BUFFER_BIT_STENCIL) @@ -448,16 +479,22 @@ intelClearWithBlit(GLcontext * ctx, GLbitfield mask) } else { /* clearing RGBA */ - CMD = XY_COLOR_BLT_CMD | - XY_BLT_WRITE_ALPHA | XY_BLT_WRITE_RGB; + CMD |= XY_BLT_WRITE_ALPHA | XY_BLT_WRITE_RGB; } } else { ASSERT(cpp == 2 || cpp == 0); - BR13 = (0xF0 << 16) | (pitch * cpp) | (1 << 24); - CMD = XY_COLOR_BLT_CMD; + BR13 |= (1 << 24); } +#ifndef I915 + if (irb_region->tiled) { + CMD |= XY_DST_TILED; + pitch /= 4; + } +#endif + BR13 |= (pitch * cpp); + if (buf == BUFFER_DEPTH || buf == BUFFER_STENCIL) { clearVal = clear_depth; } diff --git a/src/mesa/drivers/dri/intel/intel_blit.h b/src/mesa/drivers/dri/intel/intel_blit.h index a66af863591..35cc8868d94 100644 --- a/src/mesa/drivers/dri/intel/intel_blit.h +++ b/src/mesa/drivers/dri/intel/intel_blit.h @@ -42,9 +42,11 @@ extern void intelEmitCopyBlit(struct intel_context *intel, GLshort src_pitch, dri_bo *src_buffer, GLuint src_offset, + GLboolean src_tiled, GLshort dst_pitch, dri_bo *dst_buffer, GLuint dst_offset, + GLboolean dst_tiled, GLshort srcx, GLshort srcy, GLshort dstx, GLshort dsty, GLshort w, GLshort h, @@ -55,6 +57,7 @@ extern void intelEmitFillBlit(struct intel_context *intel, GLshort dst_pitch, dri_bo *dst_buffer, GLuint dst_offset, + GLboolean dst_tiled, GLshort x, GLshort y, GLshort w, GLshort h, GLuint color); diff --git a/src/mesa/drivers/dri/intel/intel_regions.c b/src/mesa/drivers/dri/intel/intel_regions.c index 009e8a3e3e2..8d9fb5c1f70 100644 --- a/src/mesa/drivers/dri/intel/intel_regions.c +++ b/src/mesa/drivers/dri/intel/intel_regions.c @@ -152,7 +152,8 @@ intel_region_create_static(intelScreenPrivate *intelScreen, unsigned int bo_handle, GLuint offset, void *virtual, - GLuint cpp, GLuint pitch, GLuint height) + GLuint cpp, GLuint pitch, GLuint height, + GLboolean tiled) { struct intel_region *region = calloc(sizeof(*region), 1); DBG("%s\n", __FUNCTION__); @@ -161,6 +162,7 @@ intel_region_create_static(intelScreenPrivate *intelScreen, region->pitch = pitch; region->height = height; /* needed? */ region->refcount = 1; + region->tiled = tiled; if (intelScreen->ttm) { assert(bo_handle != -1); @@ -188,13 +190,15 @@ intel_region_update_static(intelScreenPrivate *intelScreen, unsigned int bo_handle, GLuint offset, void *virtual, - GLuint cpp, GLuint pitch, GLuint height) + GLuint cpp, GLuint pitch, GLuint height, + GLboolean tiled) { DBG("%s\n", __FUNCTION__); region->cpp = cpp; region->pitch = pitch; region->height = height; /* needed? */ + region->tiled = tiled; /* * We use a "shared" buffer type to indicate buffers created and @@ -329,8 +333,8 @@ intel_region_copy(intelScreenPrivate *intelScreen, intelEmitCopyBlit(intel, dst->cpp, - src->pitch, src->buffer, src_offset, - dst->pitch, dst->buffer, dst_offset, + src->pitch, src->buffer, src_offset, src->tiled, + dst->pitch, dst->buffer, dst_offset, dst->tiled, srcx, srcy, dstx, dsty, width, height, GL_COPY); } @@ -362,7 +366,7 @@ intel_region_fill(intelScreenPrivate *intelScreen, intelEmitFillBlit(intel, dst->cpp, - dst->pitch, dst->buffer, dst_offset, + dst->pitch, dst->buffer, dst_offset, dst->tiled, dstx, dsty, width, height, color); } @@ -425,6 +429,7 @@ intel_region_cow(intelScreenPrivate *intelScreen, struct intel_region *region) { struct intel_context *intel = intelScreenContext(intelScreen); struct intel_buffer_object *pbo = region->pbo; + GLboolean was_locked = intel->locked; if (intel == NULL) return; @@ -440,34 +445,22 @@ intel_region_cow(intelScreenPrivate *intelScreen, struct intel_region *region) intel_batchbuffer_flush(intel->batch); - if (!intel->locked) { + was_locked = intel->locked; + if (intel->locked) LOCK_HARDWARE(intel); - intelEmitCopyBlit(intel, - region->cpp, - region->pitch, - region->buffer, 0, - region->pitch, - pbo->buffer, 0, - 0, 0, 0, 0, - region->pitch, region->height, - GL_COPY); - - intel_batchbuffer_flush(intel->batch); + + intelEmitCopyBlit(intel, + region->cpp, + region->pitch, region->buffer, 0, region->tiled, + region->pitch, pbo->buffer, 0, region->tiled, + 0, 0, 0, 0, + region->pitch, region->height, + GL_COPY); + + intel_batchbuffer_flush(intel->batch); + + if (was_locked) UNLOCK_HARDWARE(intel); - } - else { - intelEmitCopyBlit(intel, - region->cpp, - region->pitch, - region->buffer, 0, - region->pitch, - pbo->buffer, 0, - 0, 0, 0, 0, - region->pitch, region->height, - GL_COPY); - - intel_batchbuffer_flush(intel->batch); - } } dri_bo * diff --git a/src/mesa/drivers/dri/intel/intel_regions.h b/src/mesa/drivers/dri/intel/intel_regions.h index 77bb6916d16..a0d9a9005f4 100644 --- a/src/mesa/drivers/dri/intel/intel_regions.h +++ b/src/mesa/drivers/dri/intel/intel_regions.h @@ -53,6 +53,7 @@ struct intel_region GLuint map_refcount; /**< Reference count for mapping */ GLuint draw_offset; /**< Offset of drawing address within the region */ + GLboolean tiled; /**< True if the region is X or Y-tiled. Used on 965. */ struct intel_buffer_object *pbo; /* zero-copy uploads */ }; @@ -78,7 +79,7 @@ extern struct intel_region GLuint offset, void *virtual, GLuint cpp, - GLuint pitch, GLuint height); + GLuint pitch, GLuint height, GLboolean tiled); extern void intel_region_update_static(intelScreenPrivate *intelScreen, struct intel_region *region, @@ -87,7 +88,8 @@ intel_region_update_static(intelScreenPrivate *intelScreen, unsigned int bo_handle, GLuint offset, void *virtual, - GLuint cpp, GLuint pitch, GLuint height); + GLuint cpp, GLuint pitch, GLuint height, + GLboolean tiled); void intel_region_idle(intelScreenPrivate *intelScreen, diff --git a/src/mesa/drivers/dri/intel/intel_screen.c b/src/mesa/drivers/dri/intel/intel_screen.c index 0a118c742b7..00ad4b14caf 100644 --- a/src/mesa/drivers/dri/intel/intel_screen.c +++ b/src/mesa/drivers/dri/intel/intel_screen.c @@ -182,14 +182,15 @@ intel_recreate_static(intelScreenPrivate *intelScreen, region_desc->bo_handle, region_desc->offset, region_desc->map, intelScreen->cpp, region_desc->pitch / intelScreen->cpp, - intelScreen->height); + intelScreen->height, region_desc->tiled); } else { region = intel_region_create_static(intelScreen, name, mem_type, region_desc->bo_handle, region_desc->offset, region_desc->map, intelScreen->cpp, region_desc->pitch / intelScreen->cpp, - intelScreen->height); + intelScreen->height, + region_desc->tiled); } assert(region->buffer != NULL); @@ -337,23 +338,27 @@ intelUpdateScreenFromSAREA(intelScreenPrivate * intelScreen, intelScreen->front.pitch = sarea->pitch * intelScreen->cpp; intelScreen->front.handle = sarea->front_handle; intelScreen->front.size = sarea->front_size; + intelScreen->front.tiled = sarea->front_tiled; intelScreen->back.offset = sarea->back_offset; intelScreen->back.pitch = sarea->pitch * intelScreen->cpp; intelScreen->back.handle = sarea->back_handle; intelScreen->back.size = sarea->back_size; + intelScreen->back.tiled = sarea->back_tiled; if (intelScreen->driScrnPriv->ddx_version.minor >= 8) { intelScreen->third.offset = sarea->third_offset; intelScreen->third.pitch = sarea->pitch * intelScreen->cpp; intelScreen->third.handle = sarea->third_handle; intelScreen->third.size = sarea->third_size; + intelScreen->third.tiled = sarea->third_tiled; } intelScreen->depth.offset = sarea->depth_offset; intelScreen->depth.pitch = sarea->pitch * intelScreen->cpp; intelScreen->depth.handle = sarea->depth_handle; intelScreen->depth.size = sarea->depth_size; + intelScreen->depth.tiled = sarea->depth_tiled; if (intelScreen->driScrnPriv->ddx_version.minor >= 9) { intelScreen->front.bo_handle = sarea->front_bo_handle; diff --git a/src/mesa/drivers/dri/intel/intel_screen.h b/src/mesa/drivers/dri/intel/intel_screen.h index 26197b760a7..3a1a969b23f 100644 --- a/src/mesa/drivers/dri/intel/intel_screen.h +++ b/src/mesa/drivers/dri/intel/intel_screen.h @@ -45,6 +45,12 @@ typedef struct int offset; /* from start of video mem, in bytes */ int pitch; /* row stride, in bytes */ unsigned int bo_handle; /* buffer object id if available, or -1 */ + /** + * Flags if the region is tiled. + * + * Not included is Y versus X tiling. + */ + GLboolean tiled; } intelRegion; typedef struct diff --git a/src/mesa/drivers/dri/intel/intel_tex_copy.c b/src/mesa/drivers/dri/intel/intel_tex_copy.c index b85a25642a2..f1a455a04cd 100644 --- a/src/mesa/drivers/dri/intel/intel_tex_copy.c +++ b/src/mesa/drivers/dri/intel/intel_tex_copy.c @@ -40,7 +40,6 @@ #include "intel_fbo.h" #include "intel_tex.h" #include "intel_blit.h" -#include "intel_pixel.h" #define FILE_DEBUG_FLAG DEBUG_TEXTURE @@ -142,9 +141,11 @@ do_copy_texsubimage(struct intel_context *intel, -src->pitch, src->buffer, src->height * src->pitch * src->cpp, + GL_FALSE, intelImage->mt->pitch, intelImage->mt->region->buffer, image_offset, + intelImage->mt->region->tiled, x, y + height, dstx, dsty, width, height, GL_COPY); /* ? */ diff --git a/src/mesa/drivers/dri/intel/intel_tex_image.c b/src/mesa/drivers/dri/intel/intel_tex_image.c index 197cf35ebe6..44772e8588a 100644 --- a/src/mesa/drivers/dri/intel/intel_tex_image.c +++ b/src/mesa/drivers/dri/intel/intel_tex_image.c @@ -229,8 +229,8 @@ try_pbo_upload(struct intel_context *intel, intelEmitCopyBlit(intel, intelImage->mt->cpp, - src_stride, src_buffer, src_offset, - dst_stride, dst_buffer, dst_offset, + src_stride, src_buffer, src_offset, GL_FALSE, + dst_stride, dst_buffer, dst_offset, GL_FALSE, 0, 0, 0, 0, width, height, GL_COPY); -- cgit v1.2.3