"Indirect, writeable textures are not supported"

I'm getting this error when calling makeComputePipelineState. What does it mean and how can I get around it?


My compute kernel is writing to the texture in computed locations (not just based on thread_position_in_grid). Is that was is meant by "indirect"?


thanks!

Replies

I don't know for sure, and don't have time to check, but quick guess is that you have some kind of not standard texture. Like created by MTLTexture's makeTextureView, or perhaps MTLBuffer's makeTexture? Because writing to arbitrary locations in texture should definitely be possible in compute kernels.

Regards

Michal

Thanks for your reply Michal!


I'm afraid this error occurs before I've created the MTLTexture. It occurs when I call makeComputePipelineState.


The texture is a texture3d<float, access::write>

Are you able to post the Metal kernel source that is specified to create the compute pipeline state that is failing? Even a stripped down version of the kernel source that reproduces the problem should be fine.

I'm having a little trouble reprodicing that error in a simple example which I can post, but here's another error:


struct Test {

    device float* foo;

    void write() {
        foo[0] = 0;
    }
};
kernel void testError() {

    Test t;

    t.write();

}


When calling makeComputePipelineState, I get:

Error Domain=CompilerError Code=1 "Unhandled atom in intrinsic: llvm.trap" UserInfo={NSLocalizedDescription=Unhandled atom in intrinsic: llvm.trap}


on my mac and


Error Domain=AGXMetalA9 Code=1 "Compiler encountered an internal error" UserInfo={NSLocalizedDescription=Compiler encountered an internal error}


on my iPad.


Initializing foo seems to be a work-around.


Shall I file a bug?

This is not a bug. You are getting the compilation error because the variable T (and therefore foo) is unitialized.

Hope I'm not being pedantic here, but if the issue is as described I'd expect a proper compilation error message describing the lack of initialization vs some form of "internal compiler error".