Computer shaders cannot load textures that they have previously created themselves

[Spiri0]

Description

If I write a compute shader that store textures, I can load and use those textures in another compute Shader.
But when I try to pass a texture to a compute shader that I previously created with the same shader, then I get these warnings and it doesn't work:

Tint WGSL reader failure: :64:43 error: type mismatch for argument 3 in call to 'computeWGSL', expected 'texture_2d<f32>', got 'texture_storage_2d<rgba8unorm, write>'

  computeWGSL( nodeUniform0, nodeUniform1, nodeUniform0, instanceIndex, 512.0, 0.01, 250.0, f32( true ) );

                                           ^^^^^^^^^^^^
- While validating [ShaderModuleDescriptor "compute"]

- While calling [Device].CreateShaderModule([ShaderModuleDescriptor "compute"]).

[127.0.0.1/:1](http://127.0.0.1/:1) 1 error(s) generated while compiling the shader:

:64:43 error: type mismatch for argument 3 in call to 'computeWGSL', expected 'texture_2d<f32>', got 'texture_storage_2d<rgba8unorm, write>'

  computeWGSL( nodeUniform0, nodeUniform1, nodeUniform0, instanceIndex, 512.0, 0.01, 250.0, f32( true ) );

                                           ^^^^^^^^^^^^
166[Invalid ShaderModule "compute"] is invalid.

- While validating compute stage ([Invalid ShaderModule "compute"], entryPoint: main).

- While calling [Device].CreateComputePipeline([ComputePipelineDescriptor]).

166[Invalid ComputePipeline] is invalid.

- While encoding [ComputePassEncoder].SetPipeline([Invalid ComputePipeline]).

166[Invalid CommandBuffer] is invalid.

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

[127.0.0.1/:1](http://127.0.0.1/:1) Tint WGSL reader failure: :64:43 error: type mismatch for argument 3 in call to 'computeWGSL', expected 'texture_2d<f32>', got 'texture_storage_2d<rgba8unorm, write>'

  computeWGSL( nodeUniform0, nodeUniform1, nodeUniform1, instanceIndex, 512.0, 0.01, 250.0, f32( false ) );

                                           ^^^^^^^^^^^^
- While validating [ShaderModuleDescriptor "compute"]

- While calling [Device].CreateShaderModule([ShaderModuleDescriptor "compute"]).

[127.0.0.1/:1](http://127.0.0.1/:1) 1 error(s) generated while compiling the shader:

:64:43 error: type mismatch for argument 3 in call to 'computeWGSL', expected 'texture_2d<f32>', got 'texture_storage_2d<rgba8unorm, write>'

  computeWGSL( nodeUniform0, nodeUniform1, nodeUniform1, instanceIndex, 512.0, 0.01, 250.0, f32( false ) );

                                           ^^^^^^^^^^^^

[127.0.0.1/:1](http://127.0.0.1/:1) WebGPU: too many warnings, no more warnings will be reported to the console for this GPUDevice.

Reproduction steps

1.) Create a compute shader that stores a texture and reads a texture.

2.) In the first run, use an initial texture for the texture input of the compute shader.

3.) In the next run, pass the texture created by the compute shader itself to the texture input of the compute shader.

Unfortunately, in the CodePen example that I created, I only get a white image so far.

Since I can use a texture created by another computer shader, I suspect that it might have something to do with the bindings and locations.

The compute shaders are really great for IFFT (inverse fast Fourier transformations). IFFT requires computer shaders that can load their own previously created textures. Considering that the textureStorage function is brand new in 3js, it runs very satisfactorily.

Code

import * as THREE from "three";
import {texture, textureStore, instanceIndex, MeshBasicNodeMaterial, attribute, uniform, vec2, vec3, vec4, wgslFn } from 'three/nodes';
import {OrbitControls} from "three/addons/controls/OrbitControls.js";
import Stats from "three/addons/libs/stats.module.js";
import WebGPURenderer from 'three/addons/renderers/webgpu/WebGPURenderer.js';
import {EffectComposer} from "three/addons/postprocessing/EffectComposer.js";
import {RenderPass} from "three/addons/postprocessing/RenderPass.js";
import {ShaderPass} from "three/addons/postprocessing/ShaderPass.js";


const initialSpectrumWGSL = wgslFn(`

    fn computeWGSL( 
        storageTex: texture_storage_2d<rgba8unorm, write>, 
        index: u32,
        resolution: f32,
        L: f32,
        wind: vec2<f32>
    ) -> void {

        let posX = index % u32(resolution);
        let posY = index / u32(resolution);
        let indexUV = vec2u( posX, posY );
        let uv = vec2f( f32( posX ) / resolution, f32( posY ) / resolution );
                    

        let K: vec2<f32> = (2 * PI * (uv - vec2<f32>(0.5))) / L;
        let k: f32 = length(K);
        let l_wind: f32 = length(wind);
        let Omega: f32 = 0.84;
        let kp: f32 = g * square(Omega/l_wind);
        let c: f32 = omega(k)/k;

        let cp: f32 = omega(kp)/kp;
        let Lpm: f32 = exp(-1.25 * square(kp/k));
        let gamma: f32 = 1.7;
        let sigma: f32 = 0.08 * (1 + 4*pow(Omega, -3));
        let Gamma: f32 = exp(-square(sqrt(k/kp) - 1)/2 * square(sigma));
        let Jp: f32 = pow(gamma, Gamma);
        let Fp: f32 = Lpm * Jp * exp(-Omega/sqrt(10) * (sqrt(k/kp) - 1));
        let alphap: f32 = 0.006*sqrt(Omega);
        let Bl: f32 = 0.5*alphap*cp/c*Fp;
        let z0: f32 = 0.000037 * square(l_wind)/g*pow(l_wind/cp, 0.9);
        let uStar: f32 = 0.41*l_wind/log(10.0/z0);
        let alpham: f32 = alphaM(uStar);
        let Fm: f32 = exp(-0.25*square(k/KM-1.0));
        let Bh: f32 = 0.5*alpham*CM/c*Fm*Lpm;
        let a0: f32 = log(2.0)/4.0;
        let am: f32 = 0.13*uStar/CM;
        let Delta: f32 = tanH(a0+4.0*pow(c/cp, 2.5)+am*pow(CM/c, 2.5));
        let cosPhi: f32 = dot(normalize(wind), normalize(K));
        let S: f32 = (1.0/(2.0*PI))*pow(k,-4.0)*(Bl+Bh)*(1.0+Delta*(2.0*cosPhi*cosPhi-1.0));
        let dk: f32 = 2.0*PI/L;
        let h: f32 = spectrum( S, K, dk );

        textureStore( storageTex, indexUV, vec4f( h, 0, 0, 0 ) );
    }

    const PI: f32 = 3.141592653;
    const g: f32 = 9.81;
    const KM: f32 = 370;
    const CM: f32 = 0.23;
    

    fn square( x: f32 ) -> f32 {
        return x * x;
    }
            
    fn omega( k: f32 ) -> f32 {
        return sqrt(g * k * (1 + square(k / KM)));
    }

    fn tanH( x: f32 ) -> f32 {
        return (1 - exp(-2 * x)) / (1 + exp(-2 * x));
    }

    fn alphaM( uStar: f32 ) -> f32 {
        if(uStar < CM){
            return 0.01 * (1.0+log(uStar/CM));
        }
        else{
            return 0.01 * (1.0+3.0*log(uStar/CM));
        }
    }

    fn spectrum( S: f32, K: vec2<f32>, dk: f32 ) -> f32 {
        if(K.x == 0 && K.y == 0){
            return 0;
        }
        else{
            return sqrt(S/2.0)*dk;
        }
    } 
                
`);


const phaseWGSL = wgslFn(`

    fn computeWGSL( 
        storageTexPing: texture_storage_2d<rgba8unorm, write>,
        storageTexPong: texture_storage_2d<rgba8unorm, write>,
        phase: texture_2d<f32>,
        index: u32, 
        resolution: f32,
        time: f32,
        L: f32,
        pingPhase: f32
    ) -> void {

        let posX = index % u32(resolution);
        let posY = index / u32(resolution);
        let indexUV = vec2u( posX, posY );
        let uv = vec2f( f32( posX ) / resolution, f32( posY ) / resolution );

        
		    let K: vec2<f32> = (2 * PI * (uv - vec2<f32>(0.5))) / L;
        let ph = textureLoad(phase, indexUV, 0).r;
        let dph = omega(length(K)) * time;
        let new_ph = (ph + dph) % (2.0 * PI);
 
        
        if(u32(pingPhase) == 1){
            textureStore(storageTexPong, indexUV, vec4f(new_ph, 0, 0, 0));
        }
        if(u32(pingPhase) == 0){
            textureStore(storageTexPing, indexUV, vec4f(new_ph, 0, 0, 0));
        }

    }


    const PI: f32 = 3.141592653;
    const g: f32 = 9.81;
    const KM: f32 = 370.0;
    

    fn omega(k: f32) -> f32 {
		   return sqrt(g*k*(1.0+k*k/KM*KM));
	  }
         
`);




let scene = new THREE.Scene();
scene.background = new THREE.Color(0x00001f);
let renderer = new WebGPURenderer();
renderer.setPixelRatio(window.devicePixelRatio);
renderer.setSize(window.innerWidth,window.innerHeight);
document.body.appendChild(renderer.domElement);
let camera = new THREE.PerspectiveCamera(50.0, window.innerWidth/window.innerHeight, 0.5, 10000);

let composer = new EffectComposer(renderer);
composer.setSize(window.innerWidth,window.innerHeight);
const renderPass = new RenderPass(scene, camera);
composer.addPass(renderPass);

camera.position.set(3, 3, 3);

let controls = new OrbitControls(camera, renderer.domElement);


window.addEventListener("resize", onWindowResize, false);



var changed = true;
var initial = true;
var pingPhase = true;
			
var initialSpectrumMaterial = new MeshBasicNodeMaterial();	
var pingMaterial = new MeshBasicNodeMaterial();
var pongMaterial = new MeshBasicNodeMaterial();



init();
render();

var initialSpectrumParams, initialSpectrumTexture, initialSpectrumCall;
var phaseParams, pingTexture, pongTexture, phaseCall, computePhase;
    

function init() {
  
              			
  initialSpectrumParams = {
    resolution: 512,
    workgroup_size: [8, 8, 1],
    L: 100,
    wind: vec2(10, 10)
  }
	initialSpectrumTexture = new THREE.Texture();
	initialSpectrumTexture.image = {width: initialSpectrumParams.resolution, height: initialSpectrumParams.resolution};
  initialSpectrumTexture.magFilter = initialSpectrumTexture.minFilter = THREE.NearestFilter;

  initialSpectrumCall = initialSpectrumWGSL({ 
    storageTex: textureStore(initialSpectrumTexture), 
    index: instanceIndex,
    resolution: initialSpectrumParams.resolution,
    L: initialSpectrumParams.L,
    wind: initialSpectrumParams.wind,
  });



  phaseParams = {
    resolution: 512,
    workgroup_size: [8, 8, 1],
    time: 0,
    L: 250,
  }
  pingTexture = new THREE.Texture();
  pongTexture = new THREE.Texture();    
  pingTexture.image = {width: phaseParams.resolution, height: phaseParams.resolution};
  pongTexture.image = {width: phaseParams.resolution, height: phaseParams.resolution};
  pingTexture.magFilter = pingTexture.minFilter = THREE.NearestFilter;
  pongTexture.magFilter = pongTexture.minFilter = THREE.NearestFilter;

  phaseCall = phaseWGSL({ 
    storageTexPing: textureStore(pingTexture), 
    storageTexPong: textureStore(pongTexture), 
    index: instanceIndex,
    resolution: phaseParams.resolution,
    phase: null,
    time: phaseParams.time,
    L: phaseParams.L,
    pingPhase: pingPhase
  });
  computePhase = phaseCall.compute(phaseParams.resolution * phaseParams.resolution); 
  
  
  const geometry = new THREE.PlaneGeometry( 1, 1 );
  const plane = new THREE.Mesh( geometry, initialSpectrumMaterial );
  scene.add( plane );
  
  

}


function update() {
  
  if(changed) {
    var computeInitialSpectrum = initialSpectrumCall.compute(initialSpectrumParams.resolution * initialSpectrumParams.resolution);
    renderer.compute(computeInitialSpectrum, initialSpectrumParams.workgroup_size);
    changed = false;
  }  
  
  
  if(initial) {
    computePhase.computeNode.parameters.phase = texture(InitPhaseDataTexture());
		initial = false;
	}
  else{
    //computePhase.computeNode.parameters.phase = pingPhase ? texture(pingTexture) : texture(pongTexture); //here i get strange warnings and it doesn't work

    computePhase.computeNode.parameters.phase = texture(initialSpectrumTexture); //this works (just for test)
  }
  computePhase.computeNode.parameters.time.value = 0.01;
  computePhase.computeNode.parameters.pingPhase.value = pingPhase;
  computePhase = phaseCall.compute(phaseParams.resolution * phaseParams.resolution);
  renderer.compute(computePhase, phaseParams.workgroup_size);
  pingPhase = !pingPhase;
   
  
  
  initialSpectrumMaterial.colorNode = texture(InitPhaseDataTexture());
  
  
}

function render() {
  requestAnimationFrame(render);
  
  update();
  
  renderer.render(scene, camera);
  //composer.render();
}


function onWindowResize() {
  camera.aspect = window.innerWidth / window.innerHeight;
  camera.updateProjectionMatrix();
  renderer.setSize(window.innerWidth, window.innerHeight);
}



function InitPhaseDataTexture(){
  var Res = 512;
	let phaseArray = new Float32Array(Res * Res * 4);
	for (let y = 0; y < Res; y++) {
		for (let x = 0; x < Res; x++) {
			phaseArray[y*Res*4 + x*4] = Math.random() * 2.0 * Math.PI;
			phaseArray[y*Res*4 + x*4+1] = 0.0;
			phaseArray[y*Res*4 + x*4+2] = 0.0;
			phaseArray[y*Res*4 + x*4+3] = 0.0;
		}
	}
	const pingPhaseTexture = new THREE.DataTexture(phaseArray, Res, Res, THREE.RGBAFormat, THREE.FloatType);
			
	pingPhaseTexture.minFilter = THREE.NearestFilter;
	pingPhaseTexture.magFilter = THREE.NearestFilter;
	pingPhaseTexture.wrapS = THREE.ClampToEdgeWrapping;
	pingPhaseTexture.wrapT = THREE.ClampToEdgeWrapping;
	pingPhaseTexture.needsUpdate = true;
			
	return pingPhaseTexture;
}

Live example

https://codepen.io/Spiri0/pen/abPwOoG

Screenshots

No response

Version

156

Device

Desktop

Browser

Chrome

OS

Windows

[Spiri0]

I created a much shorter example.

https://codepen.io/Spiri0/pen/wvReJKR

I have no idea why it doesn't work in CodePen, but I have the same code locally and use composer to control the textures. Basically, this is secondary because the error message in the console is independent of it. I have left out all unnecessary ballast in the example. When I want to use line 115 so that the shader should read in its previously created texture, the error message comes up.

I have no idea how the bindings and locations are processed in the background, but based on the error message I suspect the problem is there.
When trying to use textures previously created by the shader on its texture_2d port, due console complains that a data type texture_storage_2d is present.
I hope the much shorter example helps

[sunag]

This seems confusing to me, using textureStore for a ping texture, these changes should be made directly to textureNode.value like you did in phase.

storageTexPing: textureStore(pingTexture), 
storageTexPong: textureStore(pongTexture), 
//
computePhase.computeNode.parameters.phase = pingPhase ? texture(pingTexture) : texture(pongTexture);

There is a lot in this code that is not part of the issue, I would suggest if possible to make it more minimalist to facilitate interpretation.

[Spiri0]

I have reduced this significantly. I also didn't like the scope of my example.

Have you seen this yet? I deleted everything unnecessary:
https://codepen.io/Spiri0/pen/wvReJKR

The interesting thing is that the computer shader understand textures from other computer shaders. It only seems to have a problem with the ones which it stored self before.

[sunag]

Have you seen this yet? I deleted everything unnecessary:

Thanks! But still about the code, I don't think this is a bug, the same texture intended to be written and read in the same pass, it seems unlikely, that's why there is the ping/pong method, try remove phase and do something like:

readTexPing: texture(pingTexture), 
storeTexPong: textureStore(pongTexture), 

and change the uniforms values between the frames and ping/pong as I said here.

[Spiri0]

I thought I would avoid exactly the conflict with reading and writing at the same time with my ping-pong logic.

I initialize the pingtexture with a datatexture. In the shader, when the pingtexture is at the input, only the condition for saving the pongtexture is met. And vice versa. This then changes after every interval. In the further course, two more computer shaders need the ping or pong texture.

You're talking about uniforms "changing the uniforms values" and "textureNode.value". I only know the parameters in the node system that I perceived as uniforms. I'll take a closer look at the whole thing, because if you're sure that I did something wrong, I don't want to keep you any further. That's a bit embarrassing for me.
I'm sorry and thank you for took the time.

[sunag]

I thought I would avoid exactly the conflict with reading and writing at the same time with my ping-pong logic.

If you try to use the same texture to write/read in the same render-pass, like this code below, it's unlikely to work:

someFn( {
  storageTexPing: textureStore( pingTexture ), 
  storageTexPong: textureStore( pongTexture ), 
  //...
  phase: texture( pingTexture ) // pingTexture already being used in the pass to write
} );

Try something like this:

const readTexNode = texture( pingTexture );
const writeTex = textureStore( pongTexture );

someFn( {
  readTex : readTexNode, 
  writeTex : writeTexNode
} );

// frame update
readTexNode.value = frame % 2 === 0 ? pingTexture : pongTexture;
writeTexNode.value = frame % 2 === 0 ? pongTexture : pingTexture;

// compute ...

[Spiri0]

Now I understand your argument with the textureNode and the uniforms. That's pretty elegant. I still tend to fall into the webgl way of thinking from time to time. I tested your recommendation and then updated the small CodePen example.

https://codepen.io/Spiri0/pen/wvReJKR

This has now become quite manageable. On my computer I can now see the uv test image in the pingTexture and the pongTexture. But I get a lot of warnings:

vAttributeNode.js:88 AttributeNode: Attribute "normal" not found.
generate @ AttributeNode.js:88
167[Texture] usage (TextureUsage::(CopySrc|CopyDst|TextureBinding|RenderAttachment)) doesn't include TextureUsage::StorageBinding.
 - While validating entries[0] as a Texture.
Expected entry layout: { binding: 0, visibility: ShaderStage::Compute, storageTexture: { access: StorageTextureAccess::WriteOnly, format: TextureFormat::RGBA8Unorm, viewDimension: TextureViewDimension::e2D } }
 - While validating [BindGroupDescriptor] against [BindGroupLayout]
 - While calling [Device].CreateBindGroup([BindGroupDescriptor]).

167[Invalid BindGroup] is invalid.
 - While encoding [ComputePassEncoder].SetBindGroup(0, [Invalid BindGroup], 0, ...).

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

127.0.0.1/:1 WebGPU: too many warnings, no more warnings will be reported to the console for this GPUDevice.

I'm trying to find out what all the warnings are about.

[Spiri0]

I didn't want to bother you any longer because I see there's enough to do and I could do the whole thing with two separate computer shaders. I admit I don't understand what the warning waterfall in the console is all about with the bindings. I can find very little or nothing about such topics on the web. I saw that in babylon.js the bindings are specifically reconfigured in the pingpong compute shader. This doesn't seem to be a trivial matter - swapping textures and textureStores. Looks like you have now done something that brings a solution with r157. If time permits, I would be happy if you could explain what the problem is or was. The better I understand what the causes are or were, the more precisely I can analyze it in advance if I come across something unusual again in order to save work for you.

[sunag]

The implementation of the StorageTexture class will simplify and facilitate the implementation of the ping/pong technique, I added a new example here #26769, it is better than the examples I sketched here, I find it more controllable to treat each compute pass as a function, in this principle I have a compute pass to initialize the texture with a noise, and in the loop, I obtain the hash of the noise from the previous noise frame to create a new noise . Ping/Pong is only really necessary if you intend to read and write to the "same texture"... I would recommend starting with the new example to continue and check if there are still any issues left.

[Spiri0]

Will the StorageTexture class understand the other texel formats in compute shaders in addition to 8unorm? I'm thinking of 16float and 32float. Unfortunately, the 8unorm is a bit too low for some data textures, 16float would be very helpful.
The StorageTexture is a big improvement. I already mentioned to mrdoob that in my opinion the StorageTexture class replaces the DataTexture class.