1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
|
#import <Metal/Metal.h>
#import <Foundation/Foundation.h>
#include "metal_context.h"
#include "core/log.h"
#include <vector>
namespace Donut
{
bool metal_probe()
{
@autoreleasepool
{
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
if (device == nil)
{
DONUT_ERROR("Metal: no default device available");
return false;
}
id<MTLCommandQueue> queue = [device newCommandQueue];
const char* name = [[device name] UTF8String];
DONUT_INFO("Metal device: {}", name ? name : "(unknown)");
DONUT_INFO("Metal: unified memory = {}, max threads/threadgroup = {}",
device.hasUnifiedMemory ? "yes" : "no",
(unsigned long)device.maxThreadsPerThreadgroup.width);
if (queue == nil)
{
DONUT_WARN("Metal: failed to create command queue");
return false;
}
return true;
}
}
// Compiles a trivial compute kernel, dispatches it over a small texture, and
// reads the result back to confirm the full Metal compute path works: source
// compilation, pipeline state, command encoding, dispatch, and shared-memory
// read-back. This is the mechanism the geodesic ray tracer will run on.
bool metal_compute_self_test()
{
@autoreleasepool
{
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
id<MTLCommandQueue> queue = [device newCommandQueue];
if (device == nil || queue == nil)
return false;
NSString* src =
@"#include <metal_stdlib>\n"
"using namespace metal;\n"
"kernel void selfTest(texture2d<float, access::write> outTex [[texture(0)]],\n"
" uint2 gid [[thread_position_in_grid]])\n"
"{\n"
" uint w = outTex.get_width();\n"
" uint h = outTex.get_height();\n"
" if (gid.x >= w || gid.y >= h) return;\n"
" outTex.write(float4(float(gid.x) / float(w - 1),\n"
" float(gid.y) / float(h - 1), 0.5, 1.0), gid);\n"
"}\n";
NSError* err = nil;
id<MTLLibrary> lib = [device newLibraryWithSource:src options:nil error:&err];
if (lib == nil)
{
DONUT_ERROR("Metal self-test: kernel compile failed: {}",
err ? [[err localizedDescription] UTF8String] : "unknown");
return false;
}
id<MTLFunction> fn = [lib newFunctionWithName:@"selfTest"];
id<MTLComputePipelineState> pipeline =
[device newComputePipelineStateWithFunction:fn error:&err];
if (pipeline == nil)
{
DONUT_ERROR("Metal self-test: pipeline creation failed");
return false;
}
const uint32_t W = 64, H = 64;
MTLTextureDescriptor* desc =
[MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA8Unorm
width:W
height:H
mipmapped:NO];
desc.usage = MTLTextureUsageShaderWrite | MTLTextureUsageShaderRead;
desc.storageMode = MTLStorageModeShared;
id<MTLTexture> tex = [device newTextureWithDescriptor:desc];
id<MTLCommandBuffer> cb = [queue commandBuffer];
id<MTLComputeCommandEncoder> enc = [cb computeCommandEncoder];
[enc setComputePipelineState:pipeline];
[enc setTexture:tex atIndex:0];
MTLSize tg = MTLSizeMake(16, 16, 1);
MTLSize grid = MTLSizeMake(W, H, 1);
[enc dispatchThreads:grid threadsPerThreadgroup:tg];
[enc endEncoding];
[cb commit];
[cb waitUntilCompleted];
// Read back the far corner; the kernel writes (~1, ~1, 0.5, 1) there.
std::vector<uint8_t> px(W * H * 4);
[tex getBytes:px.data()
bytesPerRow:W * 4
fromRegion:MTLRegionMake2D(0, 0, W, H)
mipmapLevel:0];
size_t corner = ((size_t)(H - 1) * W + (W - 1)) * 4;
DONUT_INFO("Metal compute self-test: corner pixel RGBA = ({}, {}, {}, {})",
(int)px[corner + 0], (int)px[corner + 1],
(int)px[corner + 2], (int)px[corner + 3]);
bool ok = px[corner + 0] > 250 && px[corner + 1] > 250 &&
px[corner + 3] == 255;
DONUT_INFO("Metal compute self-test: {}", ok ? "PASS" : "FAIL");
return ok;
}
}
}
|
