2

I'm trying to very efficiently convert packed 24bpp RGB images to packed 32bpp RGBA. I've tried using vImageConvert_RGB888toRGBA8888 from Accelerate.framework, but was wondering if there was a faster way by using a compute kernel in Metal. I've tried several different approaches in Metal but the result is always considerably slower than with Accelerate.framework, even for large images with >1M pixels.

Here's what my compute kernel looks like:

kernel void rgb24_to_rgba32(texture2d<half, access::read> inTexture [[texture(0)]],
                     texture2d<half, access::write> outTexture [[texture(1)]],
                     uint2 id [[ thread_position_in_grid ]])
{   
    uint2 srcAddr1 = uint2(id.x * 3, id.y);
    uint2 srcAddr2 = uint2(id.x * 3 + 1, id.y);
    uint2 srcAddr3 = uint2(id.x * 3 + 2, id.y);

    outTexture.write(half4(inTexture.read(srcAddr1).r, inTexture.read(srcAddr2).r, inTexture.read(srcAddr3).r, 1), id);

    return;
}

I am defining the inTexture as a r8Unorm, and the outTexture as a bgra8Unorm. Both textures are loaded using .storageModeShared, so there shouldn't be any memory copies taking place.

The code works and the conversion is performed correctly, but the performance is unimpressive. I've tried different threadgroupsPerGrid and threadsPerThreadgroup settings, but none of those achieve comparable performance to Accelerate.framework.

For example, on an A7 (1st generation iPad Air), a 1024x1024 image takes around 32 ms, compared to 6 ms using Accelerate.framework. Interestingly, the difference is far smaller for a faster device such as an A9-based iPhone 6s (1.5 ms on the GPU vs. 1.1 ms using Accelerate), but the Metal implementation is always slower.

Is this just not a GPU-friendly operation (possibly due to countless unaligned memory accesses?) Might I be missing something fundamental in terms of maximizing the performance of my compute kernel?

UPDATE: I was eventually able to achieve significantly better performance than described above using the following implementation:

This approach utilizes 96-bit reads using packed_uint3, and 128-bit writes using packed_uint4 to significantly improve performance.

#define RGB24_TO_RGBA32_PIXEL1(myUint) (myUint | 0xff000000)

#define RGB24_TO_RGBA32_PIXEL2(myUint1, myUint2) (myUint1 >> 24 | \
                                                ((myUint2) << 8) | 0xff000000)


#define RGB24_TO_RGBA32_PIXEL3(myUint2, myUint3) (myUint2 >> 16 | \
                                                ((myUint3) << 16) | 0xff000000)

#define RGB24_TO_RGBA32_PIXEL4(myUint3) ((myUint3 >> 8) | 0xff000000)

inline packed_uint4 packed_rgb24_to_packed_rgba32(packed_uint3 src) {
    return uint4(RGB24_TO_RGBA32_PIXEL1(src[0]),
                 RGB24_TO_RGBA32_PIXEL2(src[0], src[1]),
                 RGB24_TO_RGBA32_PIXEL3(src[1], src[2]),
                 RGB24_TO_RGBA32_PIXEL4(src[2]));
}

kernel void rgb24_to_rgba32_textures(
                         constant packed_uint3 *src [[ buffer(0) ]],
                         device packed_uint4 *dest [[ buffer(1) ]],
                         uint2 id [[ thread_position_in_grid ]])
{
    // Process 8 pixels per thread (two packed_uint3s, each containing 4 pixels):
    uint index = id.x  * 2;
    dest[index] = packed_rgb24_to_packed_rgba32(src[index]);
    dest[index + 1] = packed_rgb24_to_packed_rgba32(src[index + 1]);
    return;
}

With this approach, the performance differential on older devices becomes far smaller (Accelerate is about 2x faster than the GPU), and on more modern (A9) devices, Metal actually winds up being about 40-50% faster.

I've tried processing one, two, or more packed_uint3 vectors per thread and the conclusion was that two vectors are the sweet spot for performance.

1
  • How to convert .bgra8Unorm texture to .r8Unorm? Jun 21, 2018 at 9:36

2 Answers 2

4

Just for the sake of closure, here is Apple's Developer Relations response to this question. The bottom line is that the GPU just doesn't offer any real advantages in this case because this conversion is not a computationally heavy operation.

After discussions with engineering, and evaluating more sample implementations, the verdict is out on Metal v.s. Accelerate performance for converting packed 24bpp RGB images to packed 32bpp RGBA images: on newer devices you can get close to the same performance using Metal but Accelerate will be faster for this operation. “vImage is an extremely well-tuned implementation and since this conversion operation is not compute heavy the best we can do is to be at parity.”

The proposed reasoning behind this is data locality and efficiently operating on multiple pixels at a time (something you’ve mentioned). The fastest Metal implementation tested processed two pixels per thread and still lagged behind vImageConvert_RGB888toRGBA8888.

There was an “optimized” implementation using Metal buffers rather than textures (something else that you’d mentioned exploring) and surprisingly this approach was slightly less performant.

Lastly, adjustment of thread groups came into discussion as well as tuning by adding code to the kernel to handle the case where the thread position in grid is outside the destination image. Again, despite these considerations Accelerate remained as the fastest implementation.

I should add that one real advantage to using Metal is CPU usage, while it's no faster, it does significantly reduce the CPU's workload. For applications where the CPU is heavily loaded, the Metal approach might actually make sense.

2
  • We did eventually get Metal to run faster than the CPU for this task, but it wasn't straightforward. Oct 14, 2016 at 1:37
  • 1
    @IanOllmann, any chance you could share that code here? I have a specific application that could greatly benefit from that.
    – ldoogy
    Oct 14, 2016 at 2:03
2

There are a couple of avenues to explore here. I can't guarantee you'll get Metal to beat Accelerate on your target devices, but perhaps there's a chance for a bit of a speed-up.

  • Consider using buffers instead of textures. Your input buffer could be of type packed_char3 and your output buffer could be of type packed_char4. Then, rather than having to do three texture reads per write, you can index into the source buffer just once per pixel. As you observe, most of these reads will be unaligned, but this approach might save you some format conversions and bandwidth.

  • Consider doing more work per kernel invocation. If your image dimensions are a multiple of 4 or 8 (for example), you can use a loop (which should get unrolled by the compiler) to process that many pixels in the kernel, thereby reducing the number of threadgroups you need to dispatch.

Accelerate is a good fit for your use case, so you might want to stick with it unless you're tight on CPU time or you can tolerate the latency of dispatching the work to the GPU and waiting for the result.

2
  • Those are good comments. I have already tried using regular buffers with packed_char (performance was about the same), but I haven't tried processing more than one pixel per invocation, that's a good idea. As you point out, the one advantage with Metal (even though it is slower) is that the CPU remains fully available.
    – ldoogy
    Oct 1, 2016 at 1:35
  • load 8 pixels using 3 128-bit loads. Reshuffle the contents to make 8 ARGB pixels. Write out using 4 uint4 stores. Data must be properly aligned, etc., etc. Oct 14, 2016 at 1:39

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Not the answer you're looking for? Browse other questions tagged or ask your own question.