write method of texture2d<int, access:write> do not work in metal shader function

Question

As mentioned in Apple's document, texture2d of shading language could be of int type. I have tried to use texture2d of int type as parameter of shader language, but the write method of texture2d failed to work.

kernel void dummy(texture2d<int, access::write> outTexture [[ texture(0) ]],

                               uint2 gid [[ thread_position_in_grid ]])
{       

     outTexture.write( int4( 2, 4, 6, 8 ), gid );

}

However, if I replace the int with float, it worked.

kernel void dummy(texture2d<float, access::write> outTexture [[ texture(0) ]],

                               uint2 gid [[ thread_position_in_grid ]])
{       

     outTexture.write( float4( 1.0, 0, 0, 1.0 ), gid );

}

Could other types of texture2d, such texture2d of int, texture2d of short and so on, be used as shader function parameters, and how to use them? Thanks for reviewing my question.

The related host codes:

 MTLTextureDescriptor *desc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA8Unorm 
 desc.usage = MTLTextureUsageShaderWrite;
 id<MTLTexture> texture = [device newTextureWithDescriptor:desc];
 [commandEncoder setTexture:texture    atIndex:0];

The code to show the output computed by GPU, w and h represents width and height of textrue, respectively.

uint8_t* imageBytes = malloc(w*h*4);
memset( imageBytes, 0, w*h*4 );
MTLRegion region = MTLRegionMake2D(0, 0, [texture width], [texture height]);
[texture getBytes:imageBytes bytesPerRow:[texture width]*4 fromRegion:region mipmapLevel:0];

for( int j = 0; j < h; j++ )
{
  printf("%3d: ", j);
  for( int i = 0; i < w*pixel_size; i++ )
  {
       printf(" %3d",imageBytes[j*w*pixel_size+i] );
  }
   printf("\n")
 }

gpu3d gpu3d · Accepted Answer · 2016-08-17T10:09:19

The Metal Shading Language Guide states that:

Note: If T is int or short, the data associated with the texture must use a signed integer format. If T is uint or ushort, the data associated with the texture must use an unsigned integer format.

All you have to do is make sure the texture you write to in the API (host code) matches what you have in the kernel function. Alternatively, you can also cast the int values into float before writing to the outTexture.

write method of texture2d<int, access:write> do not work in metal shader function

1 Answers