MetalKernel1¶
-
MLModule¶ author Felix Ritterpackage FMEwork/ReleaseMeVisdll MLMetalKernel1definition MLMetalKernel1.def see also Convolutionkeywords 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.