Reading from read_write texture2d gives black on macOS 10.13

Hello,

I have a compute shader that draws in a read_write texture2d. Running this shader on macOS 10.14 gives the expected result, but on macOS 10.13 I notice that reading the texture gives black instead of the texture's color. I checked macOS 10.14 release notes & Metal Feature Set Tables and didn't notice anything that should affect read_write textures.

The Metal function is as follow:

kernel void drawDot(texture2d<half, access::read_write=""> texture [[texture(kComputeImageIndex)]],
                    constant int2& gidOrigin [[buffer(kComputeGidOriginIndex)]],
                    uint2 gid [[thread_position_in_grid]])
{
    const uint2 pixPos = gid + uint2(gidOrigin);
    texture.write(texture.read(pixPos), pixPos);
//    texture.write(half4(1.0, 0.0, 0.0, 1.0), pixPos);
}

If I use line 7 instead of line 6, output image is white with a red square around position given by gidOrigin. If I do the texture read instead, output contains a black square, although the input texture is fully white (I checked in Xcode GPU debugger). So something wrong is happening with the texture read() call.

Am I using read_write texture access in a bad way? Adding texture.fence() does not change anything, as I would expect: I'm never trying to read a pixel written by another thread of the grid.

And in Metal Shading Language Specification pdf I didn't see any restriction there could be on using read_write texture.

Thanks!

Lucas

===============================================================

Edit: I just came upon release notes of macOS 10.12 where read_write textures were introduced: Function Texture Read-Writes

Apple is giving the following usage example:

kernel void my_kernel(texture2d<float, access::read_write=""> texA [[ texture(0) ]],
                      ushort2 gid [[ thread_position_in_grid ]])
{
    float4 color = texA.read(gid);
    color = processColor(color);
    texA.write(color, gid);
}

Which is the same kind of operation I am doing. So I suppose there is a bug in macOS 10.13…

===============================================================

Edit 2 : just to make sure I didn't miss something obvious, I took "Hello Compute" example provided by Apple and modified it just enough to use read_write texture. I observed the same behavior: all good on 10.14 but black output on 10.13. Going to file a bug report although I don't have much hope…

Below is the full diff to Hello Compute project:

index 3c1d1c8..b0307af 100644
--- a/Renderer/AAPLRenderer.m
+++ b/Renderer/AAPLRenderer.m
@@ -118,7 +118,7 @@ Implementation of renderer class which performs Metal setup and per frame render
         textureDescriptor.pixelFormat = MTLPixelFormatBGRA8Unorm;
         textureDescriptor.width = image.width;
         textureDescriptor.height = image.height;
-        textureDescriptor.usage = MTLTextureUsageShaderRead;
+        textureDescriptor.usage = MTLTextureUsageShaderWrite | MTLTextureUsageShaderRead;

         // Create an input and output texture with similar descriptors.  We'll only
         //   fill in the inputTexture however.  And we'll set the output texture's descriptor
@@ -192,13 +192,17 @@ Implementation of renderer class which performs Metal setup and per frame render
     id commandBuffer = [_commandQueue commandBuffer];
     commandBuffer.label = @"MyCommand";

+    id blitEnc = [commandBuffer blitCommandEncoder];
+    MTLOrigin orig = { 0, 0, 0};
+    MTLSize size = { _inputTexture.width, _inputTexture.height, _inputTexture.depth };
+    [blitEnc copyFromTexture:_inputTexture sourceSlice:0 sourceLevel:0 sourceOrigin:orig sourceSize:size
+                   toTexture:_outputTexture destinationSlice:0 destinationLevel:0 destinationOrigin:orig];
+    [blitEnc endEncoding];
+    
     id computeEncoder = [commandBuffer computeCommandEncoder];

     [computeEncoder setComputePipelineState:_computePipelineState];

-    [computeEncoder setTexture:_inputTexture
-                       atIndex:AAPLTextureIndexInput];
-
     [computeEncoder setTexture:_outputTexture
                        atIndex:AAPLTextureIndexOutput];

diff --git a/Renderer/AAPLShaderTypes.h b/Renderer/AAPLShaderTypes.h
index 7a57f9d..897e6d9 100644
--- a/Renderer/AAPLShaderTypes.h
+++ b/Renderer/AAPLShaderTypes.h
@@ -22,7 +22,6 @@ typedef enum AAPLVertexInputIndex
 //   Metal API texture set calls
 typedef enum AAPLTextureIndex
 {
-    AAPLTextureIndexInput  = 0,
     AAPLTextureIndexOutput = 1,
 } AAPLTextureIndex;

diff --git a/Renderer/AAPLShaders.metal b/Renderer/AAPLShaders.metal
index 47a5e29..8d0e2de 100644
--- a/Renderer/AAPLShaders.metal
+++ b/Renderer/AAPLShaders.metal
@@ -90,8 +90,7 @@ constant half3 kRec709Luma = half3(0.2126, 0.7152, 0.0722);

 // Grayscale compute kernel
 kernel void
-grayscaleKernel(texture2d<half, access::read="">  inTexture  [[texture(AAPLTextureIndexInput)]],
-                texture2d<half, access::write=""> outTexture [[texture(AAPLTextureIndexOutput)]],
+grayscaleKernel(texture2d<half, access::read_write=""> outTexture [[texture(AAPLTextureIndexOutput)]],
                 uint2                          gid         [[thread_position_in_grid]])
 {
     // Check if the pixel is within the bounds of the output texture
@@ -101,7 +100,7 @@ grayscaleKernel(texture2d<half, access::read="">  inTexture  [[texture(AAPLTextureI
         return;
     }

-    half4 inColor  = inTexture.read(gid);
+    half4 inColor  = outTexture.read(gid);
     half  gray     = dot(inColor.rgb, kRec709Luma);
     outTexture.write(half4(gray, gray, gray, 1.0), gid);
 }