/*###########################################################################################################################
# Copyright (c) 1997-2012 Ufasoft   http://ufasoft.com   mailto:support@ufasoft.com                                         #
#                                                                                                                           #
# This program is free software; you can redistribute it and/or modify it under the terms of the GNU General Public License #
# as published by the Free Software Foundation; either version 3, or (at your option) any later version.                    #                                                          #
#                                                                                                                           #
# This program is distributed in the hope that it will be useful, but WITHOUT ANY WARRANTY; without even the implied        #
# warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License for more details.     #
#                                                                                                                           #
# You should have received a copy of the GNU General Public License along with this program;                                #
# If not, see <http://www.gnu.org/licenses/>                                                                                #
###########################################################################################################################*/


#include <el/ext.h>
#include "ext-opencl.h"

#pragma comment(lib, "opencl")

namespace Ext { namespace Cl {

void ClCheck(cl_int rc) {
	if (rc != CL_SUCCESS)
		throw OpenclException(MAKE_HRESULT(SEVERITY_ERROR, FACILITY_OPENCL, -rc));
}

vector<Platform> Platform::GetAll() {
	cl_uint n;
	int rc = ::clGetPlatformIDs(0, 0, &n);
	if (rc == CL_PLATFORM_NOT_FOUND_KHR)
		return vector<Platform>();			// OpenCL not installed
	ClCheck(rc);
	cl_platform_id *p = (cl_platform_id*)alloca(n*sizeof(cl_platform_id));
	ClCheck(::clGetPlatformIDs(n, p, 0));
	return vector<Platform>(p, p+n);
}

String Platform::get_Vendor() const {
	size_t size;
	ClCheck(::clGetPlatformInfo(Id, CL_PLATFORM_VENDOR, 0, 0, &size));
	char *p  = (char*)alloca(size);
	ClCheck(::clGetPlatformInfo(Id, CL_PLATFORM_VENDOR, size, p, 0));
	return String(p);
}

vector<Device> Platform::GetDevices(cl_device_type typ) const {
	cl_uint n;
	ClCheck(::clGetDeviceIDs(Id, typ, 0, 0, &n));
	cl_device_id *p = (cl_device_id*)alloca(n*sizeof(cl_device_id));
	ClCheck(::clGetDeviceIDs(Id, typ, n, p, 0));
	return vector<Device>(p, p+n);
}

Blob Device::GetInfo(cl_device_info pname) const {
	size_t size;
	ClCheck(::clGetDeviceInfo(Id, pname, 0, 0, &size));
	Blob r(0, size);
	ClCheck(::clGetDeviceInfo(Id, pname, r.Size, r.data(), 0));
	return r;
}

NDRange::NDRange(int n)
	:	m_globalSize(1)
	,	m_localSize(1)
{
	m_global[0] = n;
	m_local[0] = DEFAULT_WORK_GROUP_SIZE;
}

ExecutionStatus Event::get_Status() {
	cl_int r;
	ClCheck(::clGetEventInfo(_self, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof r, &r, 0));
	return ExecutionStatus(r);
}

int Event::get_RefCount() {
	cl_uint r;
	ClCheck(::clGetEventInfo(_self, CL_EVENT_REFERENCE_COUNT, sizeof r, &r, 0));
	return r;
}

void Event::SetCallback(void (CL_CALLBACK *pfn)(cl_event, cl_int, void *), void *userData) {
	ClCheck(::clSetEventCallback(_self, CL_COMPLETE, pfn, userData));
}

void Wait(const Event& ev) {
	ClCheck(::clWaitForEvents(1, &ev.m_h));
}

Event CommandQueue::Enqueue(const Kernel& kernel, const NDRange& range) {
	Event ev;
	ClCheck(::clEnqueueNDRangeKernel(_self, kernel, range.m_globalSize, 0, &range.m_global[0],  &range.m_local[0], 0, 0, &ev.m_h));
	return ev;
}

Event CommandQueue::Enqueue(const Kernel& kernel) {
	Event ev;
	ClCheck(::clEnqueueTask(_self, kernel, 0, 0, &ev.m_h));
	return ev;
}

Event CommandQueue::EnqueueRead(const Buffer& buffer, void *p, size_t cb) {
	Event ev;
	ClCheck(::clEnqueueReadBuffer(_self, buffer, false, 0, cb, p, 0, 0, &ev.m_h));
	return ev;
}

Event CommandQueue::EnqueueWrite(const Buffer& buffer, const void *p, size_t cb) {
	Event ev;
	ClCheck(::clEnqueueWriteBuffer(_self, buffer, false, 0, cb, p, 0, 0, &ev.m_h));
	return ev;
}

Event CommandQueue::EnqueueMarker() {
	Event ev;
	ClCheck(::clEnqueueMarker(_self, &ev.m_h));
	return ev;
}

void CommandQueue::EnqueueBarrier() {
	ClCheck(::clEnqueueBarrier(_self));
}

Context::~Context() {
}

void Context::Create(const CContextProps& props, const Device& dev) {
	cl_context_properties *p = (cl_context_properties*)alloca((props.size()*2+1)*sizeof(cl_context_properties));
	int i = 0;
	for (auto it=props.begin(); it!=props.end(); ++it, i+=2) {
		p[i] = it->first;
		p[i+1] = it->second;
	}
	p[i] = 0;
	cl_int rc;
	OutRef() = ::clCreateContext(p, 1, &dev.Id, 0, 0, &rc);
	if (!m_h)
		ClCheck(rc);
}

void Context::CreateFromType(const CContextProps& props, cl_device_type typ) {
	cl_context_properties *p = (cl_context_properties*)alloca((props.size()*2+1)*sizeof(cl_context_properties));
	int i = 0;
	for (auto it=props.begin(); it!=props.end(); ++it, i+=2) {
		p[i] = it->first;
		p[i+1] = it->second;
	}
	p[i] = 0;
	cl_int rc;
	OutRef() = ::clCreateContextFromType(p, typ, 0, 0, &rc);
	if (!m_h)
		ClCheck(rc);
}

CommandQueue Context::CreateCommandQueue(const Device& dev, cl_command_queue_properties prop) {
	CommandQueue r;
	cl_int rc;
	r.OutRef() = ::clCreateCommandQueue(_self, dev.Id, prop, &rc);
	if (!r.m_h)
		ClCheck(rc);
	return r;
}

Buffer Context::CreateBuffer(void *p, size_t size, cl_mem_flags flags) {
	Buffer r;
	cl_int rc;
	r.OutRef() = ::clCreateBuffer(_self, flags, size, p, &rc);
	if (!r.m_h)
		ClCheck(rc);
	return r;
}

vector<Device> Context::get_Devices() const {
	cl_uint n;
	ClCheck(::clGetContextInfo(_self, CL_CONTEXT_NUM_DEVICES, sizeof n, &n, 0));
	cl_device_id *p = (cl_device_id*)alloca(n*sizeof(cl_device_id));
	ClCheck(::clGetContextInfo(_self, CL_CONTEXT_DEVICES, n*sizeof(cl_device_id), p, 0));
	return vector<Device>(p, p+n);
}

void Program::CreateFromSources(Context& ctx, const vector<String>& ss) {
	Ctx = &ctx;
	vector<const char*> pp;
	for (int i=0; i<ss.size(); ++i)
		pp.push_back(ss[i]);
	cl_int rc;
	OutRef() = ::clCreateProgramWithSource(ctx, 1, &pp[0], 0, &rc);
	if (!m_h)
		ClCheck(rc);
}

Program Program::FromBinary(Context& ctx, const Device& dev, const ConstBuf& mb) {
	cl_int st, rc;
	Program r;
	const byte *p = mb.P;
	r.OutRef() = ::clCreateProgramWithBinary(ctx, 1, &dev.Id, &mb.Size, &p, &st, &rc);
	if (!r.m_h)
		ClCheck(rc);
	return r;
}

void Program::Build(Device& dev) {
	cl_int rc = ::clBuildProgram(_self, 1, &dev.Id, 0, 0, 0);
	if (rc == CL_SUCCESS) // && || != CL_BUILD_PROGRAM_FAILURE)
		ClCheck(rc);
	else {
		cl_build_status st;
		ClCheck(::clGetProgramBuildInfo(_self, dev.Id, CL_PROGRAM_BUILD_STATUS, sizeof st, &st, 0));
		if (st == CL_SUCCESS)
			Throw(E_FAIL); //!!!
		size_t size;
		ClCheck(::clGetProgramBuildInfo(_self, dev.Id, CL_PROGRAM_BUILD_LOG, 0, 0, &size));
		vector<char> buf(size);
		ClCheck(::clGetProgramBuildInfo(_self, dev.Id, CL_PROGRAM_BUILD_LOG, size, &buf[0], 0));
		throw BuildException(MAKE_HRESULT(SEVERITY_ERROR, FACILITY_OPENCL, -rc), &buf[0]);
	}
}

Kernel Program::CreateKernel(RCString name) {
	Kernel r;
	cl_int rc;
	r.OutRef() = ::clCreateKernel(_self, name, &rc);
	if (!r.m_h)
		ClCheck(rc);
	return r;
}

vector<ProgramBinary> Program::get_Binaries() {
	vector<ProgramBinary> r;

	cl_uint n;
	ClCheck(::clGetProgramInfo(_self, CL_PROGRAM_NUM_DEVICES, sizeof(n), &n, 0));
	if (n) {
		vector<cl_device_id> vId(n);
		vector<size_t> vSize(n);
		ClCheck(::clGetProgramInfo(_self, CL_PROGRAM_DEVICES, sizeof(cl_device_id)*n, &vId[0], 0));
		ClCheck(::clGetProgramInfo(_self, CL_PROGRAM_BINARY_SIZES, sizeof(size_t)*n, &vSize[0], 0));
		r.resize(n);
		vector<void*> vp(n);
		for (int i=0; i<n; ++i) {
			r[i].DeviceId = vId[i];
			r[i].Binary.Size = vSize[i];
			vp[i] = r[i].Binary.data();
		}
		ClCheck(::clGetProgramInfo(_self, CL_PROGRAM_BINARIES,	sizeof(void*)*n, &vp[0], 0));
	}
	return r;
}


}} // Ext::Cl::