Proper way to read from a texture in a compute shader which uses a storage texture

Question

I've been searching for different resources and it seems to not be possible to read from a storage texture and then write back to it. That's why I've tried to use 2 separate textures: one as texture_storage_2d<rgba16float, write> and a second one as texture_2d<f32>. However, when I use them in the same (or even in 2 separate) bind group(s), I got this error I can't quite understand:

[Texture "Framebuffer Texture"] usage (TextureUsage::
(TextureBinding|TextureUsage::8000000)) includes writable usage and another usage in the same synchronization scope.
 - While validating compute pass usage.

[Invalid CommandBuffer] is invalid.
 - While calling [Queue].Submit([[Invalid CommandBuffer]])

According to this example, it should be possible even though a comment reported an error which might be similar..?

In another example an array<f32> is used as a storage buffer and then as an input to read pixel data info from.

And in webgpu-samples, an array<atomic<u32>> seems to be used for the same purpose, although it looks much more complicated for some reason.

Also, according to the webgpufundamentals blog post, there IS a way to have a read_write storage texture, but only 3 formats r32* can be used and I wasn't able to find any example with those.

I'm quite confused right now, so hopefully someone can bring some clarity to this whole read-write texture business. Any help would be much appreciated, thanks! Anyway, this is how I'm currently doing it:

    const framebuffer = this.device.createTexture({
      size: [Config.width, Config.height],
      label: 'Framebuffer Texture',
      format: 'rgba16float',
    
      usage:
        GPUTextureUsage.RENDER_ATTACHMENT |
        GPUTextureUsage.STORAGE_BINDING |
        GPUTextureUsage.TEXTURE_BINDING
    });
    
    const computeBindGroupLayout = this.device.createBindGroupLayout({
      label: 'Compute Bind Group Layout',
    
      entries: [{
        binding: 0,
        visibility: GPUShaderStage.COMPUTE,
        storageTexture: { format: framebuffer.format }
      }, {
        binding: 1,
        texture: {},
        visibility: GPUShaderStage.COMPUTE
      }]
    });
    
    
    this.computeBindGroup = this.device.createBindGroup({
      layout: computeBindGroupLayout,
      label: 'Compute Bind Group',
    
      entries: [{
        binding: 0,
        resource: framebuffer.createView()
      }, {
        binding: 1,
        resource: framebuffer.createView()
      }]
    });

    // Workgroup size:
    override size: u32;
    
    @group(0) @binding(0)
    var framebuffer: texture_storage_2d<rgba16float, write>;
    
    @group(0) @binding(1)
    var texture: texture_2d<f32>;
    
    @compute @workgroup_size(size, size)
    fn mainCompute(@builtin(global_invocation_id) globalInvocation: vec3u)
    {
      let color = textureLoad(texture, globalInvocation.xy, 0);
      textureStore(framebuffer, globalInvocation.xy, color);
    }

And this is how I'd like to sample the output texture in the fragment shader (works well with the framebuffer when texture is not passed and textureLoad is not used):

    @group(0) @binding(0) var texture: texture_2d<f32>;
    @group(0) @binding(1) var textureSampler: sampler;
    
    @fragment
    fn mainFragment(@location(0) coords: vec2f) -> @location(0) vec4f
    {
      return textureSample(texture, textureSampler, coords);
    }

Answer 1

I ended up using a storage buffer to duplicate output color information after each compute pass as an input for the next pass. As far I understand, the use of a framebuffer storage texture could be avoided now in the fragment shader, but that will require to modify that shader in order for it to map this storage buffer values to pixel colors and add a frame resolution uniform to achieve that. Since I wanted to reuse an already implemented render pipeline (which leverages a "framebuffer" texture), I've opted for the first option.

const computeBindGroupLayout = this.device.createBindGroupLayout({
  label: 'Compute Bind Group Layout',

  entries: [/*
    Bindings 0 - 2...
  */ {
    binding: 3,
    buffer: { type: 'storage' },
    visibility: GPUShaderStage.COMPUTE
  }, {
    binding: 4,
    visibility: GPUShaderStage.COMPUTE,
    storageTexture: { format: this.imageTexture.format }
  }]
});

const colorBuffer = this.device.createBuffer({
  label: 'Compute Color Buffer',
  usage: GPUBufferUsage.STORAGE,
  size: Image.width * Image.height * 16
});

this.computeBindGroup = this.device.createBindGroup({
  layout: computeBindGroupLayout,
  label: 'Compute Bind Group',

  entries: [/*
    Entries 0 - 2...
  */ {
    binding: 3,
    resource: { buffer: colorBuffer }
  }, {
    binding: 4,
    resource: this.imageTexture.createView()
  }]
});

@group(0) @binding(3)
var<storage, read_write> colorBuffer: array<vec3f>;

@group(0) @binding(4)
var framebuffer: texture_storage_2d<rgba16float, write>;

@compute @workgroup_size(size, size)
fn mainCompute(@builtin(global_invocation_id) globalInvocation: vec3u)
{
  let res = vec2f(textureDimensions(framebuffer));
  let coord = vec2f(globalInvocation.xy);

  if (all(coord < res))
  {
    let bufferIndex = u32(coord.x + coord.y * res.x);
    var color = colorBuffer[bufferIndex];

    colorBuffer[bufferIndex] = doSomeMagic(color);
    textureStore(framebuffer, globalInvocation.xy, vec4f(colorBuffer[bufferIndex], 1.0));
  }
}