MetalComputeOSX

AppDelegate.m

//
//  AppDelegate.m
//  MetalComputeOSX
//

#import "AppDelegate.h"

@import Metal;

#define IMAGE_SIZE 128

@interface AppDelegate ()
@property (weak) IBOutlet NSWindow *window;
@end

@implementation AppDelegate

- (void)applicationDidFinishLaunching:(NSNotification *)aNotification {
	id <MTLDevice> device = MTLCreateSystemDefaultDevice();

	id <MTLLibrary> library = [device newDefaultLibrary];
	id <MTLFunction> function = [library newFunctionWithName:@"fill"];

	NSError *error;
	id<MTLComputePipelineState> state = [device newComputePipelineStateWithFunction:function error:&error];

	NSLog(@"Pipeline state creation error: %@", error);

	id <MTLCommandQueue> queue = [device newCommandQueue];

	MTLTextureDescriptor *descriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA32Float width:IMAGE_SIZE height:IMAGE_SIZE mipmapped:NO];
	descriptor.usage = MTLTextureUsageShaderWrite | MTLTextureUsageShaderRead;

	id <MTLTexture> outTexture = [device newTextureWithDescriptor:descriptor];

	id <MTLCommandBuffer> buffer = queue.commandBuffer;
	id <MTLComputeCommandEncoder> encoder = buffer.computeCommandEncoder;
	[encoder setComputePipelineState:state];
	[encoder setTexture:outTexture atIndex:0];

	MTLSize threadGroupSize = MTLSizeMake(8, 8, 1);
	MTLSize threadGroups = MTLSizeMake(outTexture.width / threadGroupSize.width, outTexture.height / threadGroupSize.height, 1);

	[encoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadGroupSize];
	[encoder endEncoding];

	NSLog(@"Command buffer error: %@", buffer.error);

	[buffer commit];
	NSLog(@"Command buffer error: %@", buffer.error);

	[buffer waitUntilCompleted];
	NSLog(@"Command buffer error: %@", buffer.error);

	NSUInteger imageByteCount = outTexture.width * outTexture.height * 4;
	NSUInteger bytesPerRow = 4 * outTexture.width * sizeof(float);
	float *imageBytes = calloc(imageByteCount, sizeof(float));
	MTLRegion region = MTLRegionMake2D(0, 0, outTexture.width, outTexture.height);
	[outTexture getBytes:imageBytes bytesPerRow:bytesPerRow fromRegion:region mipmapLevel:0];

	for (NSUInteger x = 0; x < imageByteCount; x++) {
		if (imageBytes[x] != 0) {
			NSLog(@"Non-zero value");
		}
	}

	free(imageBytes);
}

@end

Shaders.metal

//
//  Shaders.metal
//  MetalComputeOSX
//

#include <metal_stdlib>

using namespace metal;

kernel void fill(texture2d<float, access::write> outTexture [[texture(0)]],
				 uint2 gid [[thread_position_in_grid]])
{
	outTexture.write(float4(1, 0, 0, 1), gid);
}