// Copyright 2017 GFX developers
//
// Licensed under the Apache License, Version 2.0, <LICENSE-APACHE or
// http://apache.org/licenses/LICENSE-2.0> or the MIT license <LICENSE-MIT or
// http://opensource.org/licenses/MIT>, at your option. This file may not be
// copied, modified, or distributed except according to those terms.

extern crate metal_rs as metal;
extern crate cocoa;
#[macro_use] extern crate objc;
extern crate objc_id;
extern crate objc_foundation;

use metal::*;

use cocoa::foundation::NSAutoreleasePool;

use std::mem;

static LIBRARY_SRC: &str = include_str!("compute-argument-buffer.metal");

fn main() {
    let pool = unsafe { NSAutoreleasePool::new(cocoa::base::nil) };
    let device = Device::system_default();
    let command_queue = device.new_command_queue();

    let data = [
        1u32,  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,
    ];

    let buffer = device.new_buffer_with_data(
        unsafe { mem::transmute(data.as_ptr()) },
        (data.len() * mem::size_of::<u32>()) as u64,
        MTLResourceOptions::CPUCacheModeDefaultCache);

    let sum = {
        let data = [0u32];
        device.new_buffer_with_data(
            unsafe { mem::transmute(data.as_ptr()) },
            (data.len() * mem::size_of::<u32>()) as u64,
            MTLResourceOptions::CPUCacheModeDefaultCache)
    };

    let command_buffer = command_queue.new_command_buffer();
    let encoder = command_buffer.new_compute_command_encoder();

    let library = device.new_library_with_source(LIBRARY_SRC, &CompileOptions::new()).unwrap();
    let kernel = library.get_function("sum", None).unwrap();

    let argument_encoder = kernel.new_argument_encoder(0);
    let arg_buffer = device.new_buffer(
        argument_encoder.encoded_length(),
        MTLResourceOptions::empty());
    argument_encoder.set_argument_buffer(&arg_buffer, 0);
    argument_encoder.set_buffer(&buffer, 0, 0);
    argument_encoder.set_buffer(&sum, 0, 1);

    let pipeline_state_descriptor = ComputePipelineDescriptor::new();
    pipeline_state_descriptor.set_compute_function(Some(&kernel));

    let pipeline_state = device.new_compute_pipeline_state(&pipeline_state_descriptor).unwrap();

    encoder.set_compute_pipeline_state(&pipeline_state);
    encoder.set_buffer(0, 0, Some(&arg_buffer));

    encoder.use_resource(&buffer, MTLResourceUsage::Read);
    encoder.use_resource(&sum, MTLResourceUsage::Write);

    let width = 16;

    let thread_group_count = MTLSize {
        width,
        height: 1,
        depth: 1
    };

    let thread_group_size = MTLSize {
        width: (data.len() as u64 + width) / width,
        height: 1,
        depth: 1
    };

    encoder.dispatch_thread_groups(thread_group_count, thread_group_size);
    encoder.end_encoding();
    command_buffer.commit();
    command_buffer.wait_until_completed();

    let ptr = sum.contents() as *mut u32;

    unsafe {
        assert_eq!(465, *ptr);
        msg_send![pool, release];
    }
}