MetalKernel1

MLModule

author	Felix Ritter
package	FMEwork/ReleaseMeVis
dll	MLMetalKernel1
definition	MLMetalKernel1.def
see also	Convolution
keywords	GPGPU, glsl, cuda, opencl, metal, convolution, kernel, arithmetic

Purpose

The MetalKernel1 module applies a Metal kernel to an ML image and may be used to process the image data on Metal-supported devices.

Usage

Connect the module to the image pipeline and type a Metal kernel into the kernel text field.

Details

The kernel below examplifies the integration of an image processing algorithm using Metal:

#include <metal_stdlib>

using namespace metal;

typedef struct {
  // X component of the vector.
  uint32_t x;
  // Y component of the vector.
  uint32_t y;
  // Z component of the vector.
  uint32_t z;
  // Color component of the vector.
  uint32_t c;
  // Time component of the vector.
  uint32_t t;
  // Unit/Modality/User component of the vector.
  uint32_t u;
} ImageVector;


#define KERNEL_SIZE 9

constant float cKernel[KERNEL_SIZE] = {
  0.0,  1.0,  0.0,
  1.0, -4.0,  1.0,
  0.0,  1.0,  0.0
};

kernel void kernelFunc(const device float*        in [[buffer(0)]],
                       device float*             out [[buffer(1)]],
                       constant ImageVector& imgSize [[buffer(2)]],
                       const uint3               gid [[thread_position_in_grid]])
{
  const int offset[KERNEL_SIZE] = {
    -(int)imgSize.x - 1, -(int)imgSize.x + 0, -(int)imgSize.x + 1,
                     -1,                  +0,                  +1,
     (int)imgSize.x - 1,  (int)imgSize.x + 0,  (int)imgSize.x + 1
  };

  const uint32_t addr = gid.z * imgSize.y * imgSize.x +
                                    gid.y * imgSize.x +
                                                gid.x;

  unsigned int o;
  float sum = 0;

  for (int ki = 0; ki < KERNEL_SIZE; ++ki) {
    o = addr + offset[ki];
    if (o > 0 && o < imgSize.x*imgSize.y*imgSize.z) {
      sum += in[o] * cKernel[ki];
    }
  }

  out[addr] = sum;
}

The name of the function is of no particular importance, the module uses the first kernel function it discovers in the code. In and out pointers address the input and output buffers which store single precision floating point data. The third, ImageVector parameter stores the extends of the ML image.

Input Fields

input0

name: input0, type: Image

Output Fields

output0

name: output0, type: Image

Parameter Fields

Visible Fields

Compute Device

name: computeDevice, type: String

The chosen computing device.

Kernel

name: kernel, type: String

The kernel source code.

Hidden Fields

computeDeviceList

name: computeDeviceList, type: String, persistent: no

Comma-separated list of available computing devices.