Skip to content

Instantly share code, notes, and snippets.

@mattrajca
Created October 11, 2015 19:55
Show Gist options
  • Save mattrajca/a2d46e43d57fc01bc32d to your computer and use it in GitHub Desktop.
Save mattrajca/a2d46e43d57fc01bc32d to your computer and use it in GitHub Desktop.
MetalComputeOSX
//
// 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
// 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);
}
@mattrajca
Copy link
Author

This never prints "Non-zero value" on OS X but works on iOS.

Copy link

ghost commented Oct 11, 2015

Textures on OS X use the Managed storage mode by default. Your code seems reasonable overall, but you should add a call to synchronizeResource using a blit command encoder before you commit your buffer:

id <MTLBlitCommandEncoder> blitEncoder = [buffer blitCommandEncoder];
[blitEncoder synchronizeResource:outTexture];
[blitEncoder endEncoding];

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment