diff --git a/src/java.desktop/macosx/native/libawt_lwawt/awt/shaders.metal b/src/java.desktop/macosx/native/libawt_lwawt/awt/shaders.metal index 709b1825599..31e0c70ed67 100644 --- a/src/java.desktop/macosx/native/libawt_lwawt/awt/shaders.metal +++ b/src/java.desktop/macosx/native/libawt_lwawt/awt/shaders.metal @@ -741,6 +741,33 @@ kernel void xbgr_to_rgba(const device uchar *imageBuffer [[buffer(0)]], outputBuffer[4 * gid + 3] = imageBuffer[4 * gid]; // a } +kernel void copyArea_shader(texture2d inputTex[[texture(0)]], + texture2d outputTex[[texture(1)]], + constant unsigned int& src_x [[buffer(0)]], + constant unsigned int& src_y [[buffer(1)]], + constant unsigned int& dst_x [[buffer(2)]], + constant unsigned int& dst_y [[buffer(3)]], + uint2 threadgroup_pos_in_grid [[threadgroup_position_in_grid]], + uint2 thread_pos_in_threadgroup [[thread_position_in_threadgroup]], + //uint2 gid [[thread_position_in_grid]]) + uint2 threads_per_threadgroup [[threads_per_threadgroup]]) +{ + + uint2 thread_pos = (threadgroup_pos_in_grid * threads_per_threadgroup) + thread_pos_in_threadgroup; + + uint2 src_pos = thread_pos; + src_pos.x += src_x; + src_pos.y += src_y; + + half4 pixel = inputTex.read(src_pos); + //pixel = half4(1.0, 0.0, 0.0, 1.0); + + uint2 dst_pos = thread_pos; + dst_pos.x += dst_x; + dst_pos.y += dst_y; + outputTex.write(pixel, dst_pos); +} + // ---------------------------------------------------------------------------- // Shaders for rendering in XOR Mode // ---------------------------------------------------------------------------- diff --git a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLBlitLoops.m b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLBlitLoops.m index 5c5182255da..82fb04ccb9c 100644 --- a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLBlitLoops.m +++ b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLBlitLoops.m @@ -786,7 +786,56 @@ void copyFromMTLBuffer(void *pDst, id srcBuf, jint offset, jint len, dstOps, dstOps->pTexture, ((id)dstOps->pTexture).width, ((id)dstOps->pTexture).height, x, y, width, height, dx, dy); #endif //DEBUG - @autoreleasepool { +@autoreleasepool { + MTLTextureDescriptor *textureDescriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat: MTLPixelFormatBGRA8Unorm width: width height: height mipmapped: NO]; + textureDescriptor.storageMode = MTLStorageModePrivate; + id tempTex = [[mtlc.device newTextureWithDescriptor: textureDescriptor] autorelease]; + + id cb = [mtlc createCommandBuffer]; + id computeEncoder = [cb computeCommandEncoder]; + id computePipelineState = [mtlc.pipelineStateStorage + getComputePipelineState:@"copyArea_shader"]; + [computeEncoder setComputePipelineState:computePipelineState]; + + MTLSize _threadGroupSize = MTLSizeMake(8, 8, 1); + MTLSize _threadGroupCount = MTLSizeMake( (width + _threadGroupSize.width -1)/_threadGroupSize.width, + (height + _threadGroupSize.height -1)/_threadGroupSize.height, 1); + + unsigned int src_origin_x = x; + unsigned int src_origin_y = y; + unsigned int dst_origin_x = 0; + unsigned int dst_origin_y = 0; + + + // copy from texture to tempTex + [computeEncoder setTexture:dstOps->pTexture atIndex:0]; + [computeEncoder setTexture:tempTex atIndex:1]; + [computeEncoder setBytes:&src_origin_x length:sizeof(src_origin_x) atIndex:0]; + [computeEncoder setBytes:&src_origin_y length:sizeof(src_origin_x) atIndex:1]; + [computeEncoder setBytes:&dst_origin_x length:sizeof(src_origin_x) atIndex:2]; + [computeEncoder setBytes:&dst_origin_y length:sizeof(src_origin_x) atIndex:3]; + + [computeEncoder dispatchThreadgroups:_threadGroupCount threadsPerThreadgroup:_threadGroupSize]; + + // copy from tempTex to texture + src_origin_x = 0; + src_origin_y = 0; + dst_origin_x = x + dx; + dst_origin_y = y + dy; + [computeEncoder setTexture:tempTex atIndex:0]; + [computeEncoder setTexture:dstOps->pTexture atIndex:1]; + [computeEncoder setBytes:&src_origin_x length:sizeof(src_origin_x) atIndex:0]; + [computeEncoder setBytes:&src_origin_y length:sizeof(src_origin_x) atIndex:1]; + [computeEncoder setBytes:&dst_origin_x length:sizeof(src_origin_x) atIndex:2]; + [computeEncoder setBytes:&dst_origin_y length:sizeof(src_origin_x) atIndex:3]; + + [computeEncoder dispatchThreadgroups:_threadGroupCount threadsPerThreadgroup:_threadGroupSize]; + + [computeEncoder endEncoding]; + [cb commit]; + } + + /* @autoreleasepool { id cb = [mtlc createCommandBuffer]; id blitEncoder = [cb blitCommandEncoder]; @@ -804,7 +853,7 @@ void copyFromMTLBuffer(void *pDst, id srcBuf, jint offset, jint len, [blitEncoder endEncoding]; [cb commit]; - //[cb waitUntilCompleted]; + //[cb waitUntilCompleted];*/ /*[blitEncoder copyFromTexture:dstOps->pTexture @@ -812,7 +861,7 @@ void copyFromMTLBuffer(void *pDst, id srcBuf, jint offset, jint len, toTexture:dstOps->pTexture destinationSlice:0 destinationLevel:0 destinationOrigin:MTLOriginMake(x + dx, y + dy, 0)]; [blitEncoder endEncoding];*/ - } + //} // TODO: // 1. check rect bounds // 2. support CopyArea with extra-alpha (and with custom Composite if necessary) diff --git a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLLayer.h b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLLayer.h index e3e8416ec94..73deaaeccfa 100644 --- a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLLayer.h +++ b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLLayer.h @@ -70,6 +70,7 @@ - (void) blitCallback; - (void) display; - (void) redraw; +- (void) stopDisplayLink; @end #endif /* CGLLayer_h_Included */ diff --git a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLLayer.m b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLLayer.m index 6feb19577f3..5347576d241 100644 --- a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLLayer.m +++ b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLLayer.m @@ -121,11 +121,15 @@ - (void) blitTexture { - (void) dealloc { self.javaLayer = nil; - if (CVDisplayLinkIsRunning(self.displayLink)) + [self stopDisplayLink]; + [super dealloc]; +} + +- (void) stopDisplayLink { +if (CVDisplayLinkIsRunning(self.displayLink)) CVDisplayLinkStop(self.displayLink); CVDisplayLinkRelease(self.displayLink); self.displayLink = nil; - [super dealloc]; } - (void) blitCallback { diff --git a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLRenderQueue.m b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLRenderQueue.m index 53e94e7cef6..b31c7a1e789 100644 --- a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLRenderQueue.m +++ b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLRenderQueue.m @@ -896,6 +896,7 @@ void MTLRenderQueue_CheckPreviousOp(jint op) { if (mtlsdo != NULL) { CONTINUE_IF_NULL(mtlc); MTLSD_Delete(env, mtlsdo); + if (mtlsdo->privOps != NULL) { free(mtlsdo->privOps); }