Metal kernel issues with 24-bit data...

I tried hard to find a bug in your code and couldn't. So it is either that something else is wrong (for example, the way you're invoking the kernel for 32/24 bit data, or 24 bit data order, or whatever else) or it is a Metal bug. What I'd do it to try to prove the latter. For example: 1) Test it on different GPU/device, if it works somewhere else, then it is a bug 2) Test result of that line 18 in "24" bit version, for example change line 19 to Uchar4 pixel = uchar4(ptr%256, ptr/256,255,255); and look at the patterns it produces. You can also inspect outBuffer contents in Metal debugger - remember to clear it to some distinct pattern first, so you'll see where your kernel really writes to and what? Try to write some constant/pattern from kernel, too or maybe sourceX/sourceY values. You have to break "black box" and try to determine what it is really doing. Good luck! Michał

I was having a very similar problem. I was trying to use Metal to convert an BGRA8unorm format to a compact BGR888 format.

I created a MTLBuffer for compact BGR888:

_tryBufferSize = _inputTexture.width * _inputTexture.height * 3;    // note *3 for BGR888
size_t aligned_size = ((_tryBufferSize + 4095) / 4096) * 4096;
int ret = posix_memalign(&_tryBufferRaw, 4096, aligned_size);
NSAssert(ret == 0 && _tryBufferRaw != NULL, @"posix_memalign failed.");
_tryBuffer = [_device newBufferWithBytesNoCopy:_tryBufferRaw
                                        length:aligned_size
                                       options:MTLResourceStorageModeShared
                                   deallocator:^(void * _Nonnull pointer, NSUInteger length) { free(pointer); }];

and compute with: (inTexture BGRA8888 -> outBuffer BGR888) (reduced demo code)

kernel void
tryKernel(texture2d<half, access::read>  inTexture  [[texture(0)]],
          device uchar3                  *outBuffer [[buffer(0)]],
          uint2                          gid        [[thread_position_in_grid]])
{
    if((gid.x >= inTexture.get_width()) || (gid.y >= inTexture.get_height()))
    {
        return;
    }

    uint index = gid.y * inTexture.get_width() + gid.x;
    outBuffer[index] = uchar3(1, 1, 1);    // note uchar3 for BGR888
}

This code failed.

The following code is OK to run:

_tryBufferSize = _inputTexture.width * _inputTexture.height * 4;
...
device uchar4 *outBuffer [[buffer(0)]],
...
outBuffer[index] = uchar4(1, 1, 1, 1); // uchar2 also worked fine.

Then I came upon this 4-year old question.

MSLS 2.2 Vector Data Types specified (Table 2.3):

uchar2 size 2B

uchar3 size 4B

uchar4 size 4B

which is not the way I thought it was.

It turns out Metal still supports a very straight-forward way to access non-2B-aligned element:

define the element layout first:

typedef struct __attribute__((__packed__))
{
    unsigned char r;
    unsigned char g;
    unsigned char b;
} TryRGB888;

then compute works just as charm:

_tryBufferSize = _inputTexture.width * _inputTexture.height * 3;    // CPU side memory alloc size, *3 for BGR888
...
device TryRGB888 *outBuffer [[buffer(0)]],    // shader side
...
uint index = gid.y * inTexture.get_width() + gid.x;
outBuffer[index].r = 1;
outBuffer[index].g = 1;
outBuffer[index].b = 1;

I found similar issues and I simplified the case to as below:

Given an array unit32_t* inA, I want to output an array with each element increased by 1. Every thing works until the array length becomes 1024 * 1024 * 4, when the output array becomes all 0. It works even when the array length is 1024 * 1024 * 4 - 1.

And, somehow I increase the array size to 1024 * 1024 * 4 + 128 * 128, it works again... as a really weird workround.

Could anyone explain why 1024 * 1024 * 4 is a special number?

Thanks

    kernel void increase_array(
        /* param idx 0 - setBuffer */
        device const uint32_t* inA,
        
        /* param idx 1 - setBuffer */
        device uint32_t* result,
        
        /* the thread index */
        uint index [[thread_position_in_grid]]
       )
       {
           // the for-loop is replaced with a collection of threads, each of which
           // calls this function.
           result[index] = index;
       }