#import <Foundation/Foundation.h>
#import <Metal/Metal.h>

#include <algorithm>
#include <cmath>
#include <cstdint>
#include <exception>
#include <iostream>
#include <vector>

namespace {
const char* error_message(NSError* error) {
    const char* text = error.localizedDescription.UTF8String;
    return text ? text : "no NSError description provided";
}
} // namespace

// A self-contained SDK probe. ARC owns Objective-C objects in this .mm file.
// The course renderer uses metal-cpp instead: its C++ pointers are NOT ARC-owned.
int main() try {
    @autoreleasepool {
        id<MTLDevice> device = MTLCreateSystemDefaultDevice();
        if (!device) { std::cerr << "No Metal device available\n"; return 1; }
        std::cout << "Device: " << device.name.UTF8String
                  << "\nunified memory: " << device.hasUnifiedMemory
                  << "\nmax buffer length: " << device.maxBufferLength
                  << "\nrecommended working set: " << device.recommendedMaxWorkingSetSize
                  << "\nApple7: " << [device supportsFamily:MTLGPUFamilyApple7]
                  << "\nApple8: " << [device supportsFamily:MTLGPUFamilyApple8]
                  << "\nApple9: " << [device supportsFamily:MTLGPUFamilyApple9] << '\n';
        constexpr std::uint32_t count = 1003; // Deliberately not a multiple of a group.
        // Capacity includes every possible rounded tail at widths up to 64.
        // The extra output values are sentinels, not valid simulation records.
        constexpr std::uint32_t capacity = count + 63;
        constexpr float sentinel = -12345.25f;
        constexpr std::size_t bytes = capacity * sizeof(float);
        if (bytes > device.maxBufferLength) { std::cerr << "buffer limit\n"; return 1; }
        std::vector<float> input(capacity, 0.0f);
        std::vector<float> initial_output(capacity, sentinel);
        for (std::uint32_t i = 0; i < count; ++i) input[i] = static_cast<float>(i) * 0.25f;
        id<MTLBuffer> src = [device newBufferWithBytes:input.data() length:bytes
                                             options:MTLResourceStorageModeShared];
        id<MTLBuffer> dst = [device newBufferWithBytes:initial_output.data() length:bytes
                                             options:MTLResourceStorageModeShared];
        id<MTLCommandQueue> queue = [device newCommandQueue];
        if (!src || !dst || !queue) { std::cerr << "resource allocation failed\n"; return 1; }
        NSString* source = @R"metal(
            #include <metal_stdlib>
            using namespace metal;
            kernel void orchard_gain(device const float* input [[buffer(0)]],
                                     device float* output [[buffer(1)]],
                                     constant uint& count [[buffer(2)]],
                                     uint tid [[thread_position_in_grid]]) {
                if (tid >= count) return;
                output[tid] = input[tid] * 0.5f + 0.125f;
            }
        )metal";
        NSError* error = nil;
        id<MTLLibrary> library = [device newLibraryWithSource:source options:nil error:&error];
        if (!library) {
            std::cerr << "shader compile: " << error_message(error) << '\n';
            return 1;
        }
        id<MTLFunction> function = [library newFunctionWithName:@"orchard_gain"];
        if (!function) { std::cerr << "kernel missing\n"; return 1; }
        id<MTLComputePipelineState> pipeline = [device newComputePipelineStateWithFunction:function
                                                                                     error:&error];
        if (!pipeline) {
            std::cerr << "compute pipeline: " << error_message(error) << '\n';
            return 1;
        }
        id<MTLCommandBuffer> commands = [queue commandBuffer];
        id<MTLComputeCommandEncoder> encoder = [commands computeCommandEncoder];
        if (!commands || !encoder) { std::cerr << "command allocation failed\n"; return 1; }
        commands.label = @"Copper Orchard bounded compute check";
        [encoder setComputePipelineState:pipeline];
        [encoder setBuffer:src offset:0 atIndex:0];
        [encoder setBuffer:dst offset:0 atIndex:1];
        [encoder setBytes:&count length:sizeof(count) atIndex:2];
        const NSUInteger width = std::min<NSUInteger>({64, pipeline.maxTotalThreadsPerThreadgroup,
                                                       device.maxThreadsPerThreadgroup.width});
        if (width == 0) { std::cerr << "zero threadgroup limit\n"; return 1; }
        const NSUInteger groups = (count + width - 1) / width;
        [encoder dispatchThreadgroups:MTLSizeMake(groups, 1, 1)
                threadsPerThreadgroup:MTLSizeMake(width, 1, 1)];
        [encoder endEncoding];
        [commands commit];
        // Deliberate blocking boundary for this one-shot test, never a render loop recipe.
        [commands waitUntilCompleted];
        if (commands.status != MTLCommandBufferStatusCompleted) {
            std::cerr << "GPU failed: " << error_message(commands.error) << '\n';
            return 1;
        }
        const auto* actual = static_cast<const float*>(dst.contents);
        for (std::uint32_t i = 0; i < count; ++i) {
            const float expected = input[i] * 0.5f + 0.125f;
            if (!std::isfinite(actual[i]) || std::abs(actual[i] - expected) > 1e-5f) {
                std::cerr << "result mismatch at " << i << '\n'; return 1;
            }
        }
        for (std::uint32_t i = count; i < capacity; ++i) {
            if (actual[i] != sentinel) {
                std::cerr << "tail guard overwrote sentinel at " << i << '\n'; return 1;
            }
        }
        std::cout << "Verified " << count << " GPU results and " << capacity - count
                  << " untouched padding values; dispatched " << groups * width
                  << " threads (width " << width << ").\n";
        return 0;
    }
} catch (const std::exception& error) {
    std::cerr << "Copper Orchard: " << error.what() << '\n';
    return 1;
}