aboutsummaryrefslogtreecommitdiffhomepage
diff options
context:
space:
mode:
authorGravatar Denis Konoplev <dkonoplev@openjdk.org> 2021-04-17 06:07:44 +0000
committerGravatar Jayathirth D V <jdv@openjdk.org> 2021-04-17 06:07:44 +0000
commit926e3bc0c1c93a89666b77a39515689dd29e0121 (patch)
tree04bac10f379d44b914982012107b2fb855f3e495
parentff4997014fe5462dca2b313f3f483400ffee5b62 (diff)
downloadjdk-926e3bc0c1c93a89666b77a39515689dd29e0121.tar.gz
jdk-926e3bc0c1c93a89666b77a39515689dd29e0121.zip
8264143: Lanai: RenderPerfTest.BgrSwBlitImage has artefacts on apple M1
Reviewed-by: jdv
-rw-r--r--src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/MTLBlitLoops.m34
-rw-r--r--src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/common.h5
-rw-r--r--src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/shaders.metal52
3 files changed, 47 insertions, 44 deletions
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 22fbd2aa5d7..7601aacfa62 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
@@ -49,9 +49,15 @@ typedef struct {
// Consider deleting this field, since it's always MTLPixelFormatBGRA8Unorm
jboolean hasAlpha;
jboolean isPremult;
- NSString* swizzleKernel;
+ const uint8_t* swizzleMap;
} MTLRasterFormatInfo;
+
+const uint8_t rgb_to_rgba[4] = {0, 1, 2, 3};
+const uint8_t xrgb_to_rgba[4] = {1, 2, 3, 0};
+const uint8_t bgr_to_rgba[4] = {2, 1, 0, 3};
+const uint8_t xbgr_to_rgba[4] = {3, 2, 1, 0};
+
/**
* This table contains the "pixel formats" for all system memory surfaces
* that Metal is capable of handling, indexed by the "PF_" constants defined
@@ -62,10 +68,10 @@ typedef struct {
MTLRasterFormatInfo RasterFormatInfos[] = {
{ 1, 0, nil }, /* 0 - IntArgb */ // Argb (in java notation)
{ 1, 1, nil }, /* 1 - IntArgbPre */
- { 0, 1, @"rgb_to_rgba" }, /* 2 - IntRgb */
- { 0, 1, @"xrgb_to_rgba" }, /* 3 - IntRgbx */
- { 0, 1, @"bgr_to_rgba" }, /* 4 - IntBgr */
- { 0, 1, @"xbgr_to_rgba" }, /* 5 - IntBgrx */
+ { 0, 1, rgb_to_rgba }, /* 2 - IntRgb */
+ { 0, 1, xrgb_to_rgba }, /* 3 - IntRgbx */
+ { 0, 1, bgr_to_rgba }, /* 4 - IntBgr */
+ { 0, 1, xbgr_to_rgba }, /* 5 - IntBgrx */
// TODO: support 2-byte formats
// { GL_BGRA, GL_UNSIGNED_SHORT_1_5_5_5_REV,
@@ -152,6 +158,12 @@ void drawTex2Tex(MTLContext *mtlc,
[encoder drawPrimitives:MTLPrimitiveTypeTriangle vertexStart:0 vertexCount:6];
}
+static void fillSwizzleUniforms(struct SwizzleUniforms *uniforms, const MTLRasterFormatInfo *rfi) {
+ const size_t SWIZZLE_MAP_SIZE = 4;
+ memcpy(&uniforms->swizzle, rfi->swizzleMap, SWIZZLE_MAP_SIZE);
+ uniforms->hasAlpha = rfi->hasAlpha;
+}
+
static void
replaceTextureRegion(MTLContext *mtlc, id<MTLTexture> dest, const SurfaceDataRasInfo *srcInfo,
const MTLRasterFormatInfo *rfi,
@@ -181,24 +193,30 @@ replaceTextureRegion(MTLContext *mtlc, id<MTLTexture> dest, const SurfaceDataRas
}
[buff didModifyRange:NSMakeRange(0, buff.length)];
- if (rfi->swizzleKernel != nil) {
+ if (rfi->swizzleMap != nil) {
id <MTLBuffer> swizzled = [[mtlc.device newBufferWithLength:(sw * sh * srcInfo->pixelStride) options:MTLResourceStorageModeManaged] autorelease];
// this should be cheap, since data is already on GPU
id<MTLCommandBuffer> cb = [mtlc createCommandBuffer];
id<MTLComputeCommandEncoder> computeEncoder = [cb computeCommandEncoder];
id<MTLComputePipelineState> computePipelineState = [mtlc.pipelineStateStorage
- getComputePipelineState:rfi->swizzleKernel];
+ getComputePipelineState:@"swizzle_to_rgba"];
[computeEncoder setComputePipelineState:computePipelineState];
[computeEncoder setBuffer:buff offset:0 atIndex:0];
[computeEncoder setBuffer:swizzled offset:0 atIndex:1];
+ struct SwizzleUniforms uniforms;
+ fillSwizzleUniforms(&uniforms, rfi);
+ [computeEncoder setBytes:&uniforms length:sizeof(struct SwizzleUniforms) atIndex:2];
+
+ NSUInteger pixelCount = buff.length / srcInfo->pixelStride;
+ [computeEncoder setBytes:&pixelCount length:sizeof(NSUInteger) atIndex:3];
+
NSUInteger threadGroupSize = computePipelineState.maxTotalThreadsPerThreadgroup;
if (threadGroupSize == 0) {
threadGroupSize = 1;
}
- NSUInteger pixelCount = buff.length / srcInfo->pixelStride;
MTLSize threadsPerGroup = MTLSizeMake(threadGroupSize, 1, 1);
MTLSize threadGroups = MTLSizeMake((pixelCount + threadGroupSize - 1) / threadGroupSize,
1, 1);
diff --git a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/common.h b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/common.h
index 22b3b0c22ce..556bbf5d88e 100644
--- a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/common.h
+++ b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/common.h
@@ -156,4 +156,9 @@ struct LCDFrameUniforms {
vector_float3 gamma;
vector_float3 invgamma;
};
+
+struct SwizzleUniforms {
+ unsigned char swizzle[4];
+ unsigned char hasAlpha;
+};
#endif
diff --git a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/shaders.metal b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/shaders.metal
index 44f79c868a9..efc791bb6f8 100644
--- a/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/shaders.metal
+++ b/src/java.desktop/macosx/native/libawt_lwawt/java2d/metal/shaders.metal
@@ -640,45 +640,25 @@ kernel void stencil2tex(const device uchar *imageBuffer [[buffer(0)]],
// work item deals with 4 byte pixel
// assuming that data is aligned
-kernel void rgb_to_rgba(const device uchar *imageBuffer [[buffer(0)]],
- device uchar *outputBuffer [[buffer(1)]],
- uint gid [[thread_position_in_grid]])
+kernel void swizzle_to_rgba(const device uchar *imageBuffer [[buffer(0)]],
+ device uchar *outputBuffer [[buffer(1)]],
+ constant SwizzleUniforms& uniforms [[buffer(2)]],
+ constant uint& size [[buffer(3)]],
+ uint gid [[thread_position_in_grid]])
{
- outputBuffer[4 * gid] = imageBuffer[4 * gid]; // r
- outputBuffer[4 * gid + 1] = imageBuffer[4 * gid + 1]; // g
- outputBuffer[4 * gid + 2] = imageBuffer[4 * gid + 2]; // b
- outputBuffer[4 * gid + 3] = 255; // a
-}
-
-kernel void bgr_to_rgba(const device uchar *imageBuffer [[buffer(0)]],
- device uchar *outputBuffer [[buffer(1)]],
- uint gid [[thread_position_in_grid]])
-{
- outputBuffer[4 * gid] = imageBuffer[4 * gid + 2]; // r
- outputBuffer[4 * gid + 1] = imageBuffer[4 * gid + 1]; // g
- outputBuffer[4 * gid + 2] = imageBuffer[4 * gid]; // b
- outputBuffer[4 * gid + 3] = 255; // a
-}
-
-kernel void xrgb_to_rgba(const device uchar *imageBuffer [[buffer(0)]],
- device uchar *outputBuffer [[buffer(1)]],
- uint gid [[thread_position_in_grid]])
-{
- outputBuffer[4 * gid] = imageBuffer[4 * gid + 1]; // r
- outputBuffer[4 * gid + 1] = imageBuffer[4 * gid + 2]; // g
- outputBuffer[4 * gid + 2] = imageBuffer[4 * gid + 3]; // b
- outputBuffer[4 * gid + 3] = imageBuffer[4 * gid]; // a
-}
+ if (gid > size) {
+ return;
+ }
+ outputBuffer[4 * gid] = imageBuffer[4 * gid + uniforms.swizzle[0]]; // r
+ outputBuffer[4 * gid + 1] = imageBuffer[4 * gid + uniforms.swizzle[1]]; // g
+ outputBuffer[4 * gid + 2] = imageBuffer[4 * gid + uniforms.swizzle[2]]; // b
-kernel void xbgr_to_rgba(const device uchar *imageBuffer [[buffer(0)]],
- device uchar *outputBuffer [[buffer(1)]],
- uint gid [[thread_position_in_grid]])
-{
- outputBuffer[4 * gid] = imageBuffer[4 * gid + 3]; // r
- outputBuffer[4 * gid + 1] = imageBuffer[4 * gid + 2]; // g
- outputBuffer[4 * gid + 2] = imageBuffer[4 * gid + 1]; // b
- outputBuffer[4 * gid + 3] = imageBuffer[4 * gid]; // a
+ if (uniforms.hasAlpha) {
+ outputBuffer[4 * gid + 3] = imageBuffer[4 * gid + uniforms.swizzle[3]];
+ } else {
+ outputBuffer[4 * gid + 3] = 255;
+ }
}
// ----------------------------------------------------------------------------