Sample Scene for Project Discussion with TA

Made it human-friendly
Improving parallelism in fragment shader leading to more than 20x the framerate
2019-12-29 23:39:33 +02:00 · 2019-12-27 22:53:14 +02:00 · 2019-12-26 23:07:03 +02:00 · 2019-12-25 00:09:14 +02:00 · 2019-12-18 19:09:00 +02:00 · 2019-12-17 21:24:03 +02:00
57 changed files with 978016 additions and 9449 deletions
@@ -13,9 +13,6 @@
 # User-specific files (MonoDevelop/Xamarin Studio)
 *.userprefs

-# Mono auto generated files
-mono_crash.*
-
 # Build results
 [Dd]ebug/
 [Dd]ebugPublic/
@@ -70,7 +67,6 @@ StyleCopReport.xml
 *_h.h
 *.ilk
 *.meta
-*.obj
 *.iobj
 *.pch
 *.pdb
@@ -1,9 +0,0 @@
-#include "COLLADA.h"
-
-COLLADA::COLLADA(const char* filename)
-{
-	tinyxml2::XMLDocument doc;
-	doc.LoadFile(filename);
-	tinyxml2::XMLElement* xml = doc.FirstChildElement("COLLADA");
-	m_GeometryLibrary = GeometryLibrary(xml->FirstChildElement("library_geometries"));
-}
@@ -1,18 +0,0 @@
-#pragma once
-#include "GeometryLibrary.h"
-class COLLADA
-{
-private:
-	GeometryLibrary     m_GeometryLibrary;
-	//ImagesLibrary       m_ImagesLibrary;
-	//EffectsLibrary      m_EffectsLibrary;
-	//MaterialLibrary     m_MaterialLibrary;
-	//ControllerLibrary   m_ControllerLibrary;
-	//AnimationLibrary    m_AnimationLibrary;
-	VisualSceneLibrary  m_VisualSceneLibrary;
-
-public:
-	COLLADA() = default;
-	COLLADA(const char* filename);
-};
-
@@ -1,178 +0,0 @@
-#include <iostream>
-#include <fstream>
-#include <sstream>
-#include "ColladaModel.h"
-
-#define PI 3.14159265358979323846
-#define DEG2RAD PI/180
-
-ColladaModel::ColladaModel(const char* filename) : faces_(), vertices_(), normals_()
-{
-	Transform = Matrix::identity();
-	Rotation = Matrix::identity();
-	Scale = Matrix::identity();
-	Translation = Matrix::identity();
-
-	tinyxml2::XMLDocument doc;
-	doc.LoadFile(filename);
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	////////////////////////////////////////////////faces//////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	face_count = 0;
-	tinyxml2::XMLElement* xml_face = doc.FirstChildElement("COLLADA")->FirstChildElement("library_geometries")->FirstChildElement("geometry")->FirstChildElement("mesh")->FirstChildElement("triangles");
-	xml_face->QueryIntAttribute("count", &face_count); //get count
-	std::stringstream str_triangle(xml_face->FirstChildElement("p")->GetText()); //get values as a string
-
-	for (int i = 0; i < face_count; i++)
-	{
-		faces_.push_back(std::vector<Vec3i>());
-		for (int j = 0; j < 3; j++)
-		{
-			Vec3i temp;
-			str_triangle >> temp.x;
-			str_triangle >> temp.y;
-			str_triangle >> temp.z;
-			faces_[i].push_back(temp);
-		}
-	}
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	////////////////////////////////////////////////vertex/////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	vertex_count = 0;
-	tinyxml2::XMLElement* xml_vertex = doc.FirstChildElement("COLLADA")->FirstChildElement("library_geometries")->FirstChildElement("geometry")->FirstChildElement("mesh")->FirstChildElement("source")->FirstChildElement("float_array");
-	xml_vertex->QueryIntAttribute("count", &vertex_count);
-	std::stringstream str_vertex(xml_vertex->GetText());
-
-	for (int i = 0; i < vertex_count / 3; i++)
-	{
-		Vec3f temp;
-		str_vertex >> temp.x;
-		str_vertex >> temp.y;
-		str_vertex >> temp.z;
-		vertices_.push_back(temp);
-	}
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	/////////////////////////////////////////////normal////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	normal_count = 0;
-	tinyxml2::XMLElement* xml_normal = doc.FirstChildElement("COLLADA")->FirstChildElement("library_geometries")->FirstChildElement("geometry")->FirstChildElement("mesh")->FirstChildElement("source")->NextSiblingElement()->FirstChildElement("float_array");
-	xml_normal->QueryIntAttribute("count", &normal_count);
-	std::stringstream str_normal(xml_normal->GetText());
-
-	for (int i = 0; i < normal_count / 3; i++)
-	{
-		Vec3f temp;
-		str_normal >> temp.x;
-		str_normal >> temp.y;
-		str_normal >> temp.z;
-		normals_.push_back(temp);
-	}
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	texcoord_count = 0;
-	tinyxml2::XMLElement* xml_texture = doc.FirstChildElement("COLLADA")->FirstChildElement("library_geometries")->FirstChildElement("geometry")->FirstChildElement("mesh")->FirstChildElement("source")->NextSiblingElement()->NextSiblingElement()->FirstChildElement("float_array");
-	xml_texture->QueryIntAttribute("count", &texcoord_count);
-	std::stringstream str_texcoord(xml_texture->GetText());
-
-	for (int i = 0; i < texcoord_count / 2; i++)
-	{
-		Vec2f temp;
-		str_texcoord >> temp.x;
-		str_texcoord >> temp.y;
-		textureco_.push_back(temp);
-	}
-
-	load_texture(filename, "_diffuse.tga", diffusemap_);
-	load_texture(filename, "_nm_tangent.tga", normalmap_);
-}
-
-ColladaModel::~ColladaModel() {
-}
-
-void ColladaModel::ApplyTransform() {
-	Transform = Translation * Scale * Rotation;
-}
-
-void ColladaModel::translate(Vec3f tr) {
-	Translation[0][3] += tr.x;
-	Translation[1][3] += tr.y;
-	Translation[2][3] += tr.z;
-}
-void ColladaModel::rotate(Vec3f rot) {
-	rot = rot * DEG2RAD;
-
-	Rotation[0][0] = cosf(rot.y) * cosf(rot.z);
-	Rotation[0][1] = -cosf(rot.y) * sinf(rot.z);
-	Rotation[0][2] = sinf(rot.y);
-	Rotation[1][0] = sinf(rot.x) * sinf(rot.y) * cosf(rot.z) + cosf(rot.x) * sinf(rot.z);
-	Rotation[1][1] = -sinf(rot.x) * sinf(rot.y) * sinf(rot.z) + cosf(rot.x) * cosf(rot.z);
-	Rotation[1][2] = -sinf(rot.x) * cosf(rot.y);
-	Rotation[2][0] = -cosf(rot.x) * sinf(rot.y) * cosf(rot.z) + sinf(rot.x) * sinf(rot.z);
-	Rotation[2][1] = cosf(rot.x) * sinf(rot.y) * sinf(rot.z) + sinf(rot.x) * cosf(rot.z);
-	Rotation[2][2] = cosf(rot.x) * cosf(rot.y);
-}
-void ColladaModel::scale(Vec3f scl) {
-	Scale[0][0] = scl.x;
-	Scale[1][1] = scl.y;
-	Scale[2][2] = scl.z;
-}
-
-int ColladaModel::nfaces() {
-	return face_count;
-}
-int ColladaModel::nvertices() {
-	return vertex_count;
-}
-
-std::vector<int> ColladaModel::face(int idx) {
-	std::vector<int> face;
-	for (int i = 0; i < (int)faces_[idx].size(); i++) face.push_back(faces_[idx][i][0]);
-	return face;
-}
-
-Vec3f ColladaModel::vertix(int i) {
-	return vertices_[i];
-}
-
-Vec3f ColladaModel::vertix(int iface, int nthvert) {
-	return vertices_[faces_[iface][nthvert][0]];
-}
-
-void ColladaModel::load_texture(std::string filename, const char* suffix, TGAImage& img) {
-	std::string texfile(filename);
-	size_t dot = texfile.find_last_of(".");
-	if (dot != std::string::npos) {
-		texfile = texfile.substr(0, dot) + std::string(suffix);
-		std::cerr << "texture file " << texfile << " loading " << (img.read_tga_file(texfile.c_str()) ? "ok" : "failed") << std::endl;
-		img.flip_vertically();
-	}
-}
-
-TGAColor ColladaModel::diffuse(Vec2f uvf) {
-	Vec2i uv(uvf[0] * diffusemap_.get_width(), uvf[1] * diffusemap_.get_height());
-	return diffusemap_.get(uv[0], uv[1]);
-}
-
-Vec3f ColladaModel::normal(Vec2f uvf) {
-	Vec2i uv(uvf[0] * normalmap_.get_width(), uvf[1] * normalmap_.get_height());
-	TGAColor c = normalmap_.get(uv[0], uv[1]);
-	Vec3f res;
-	for (int i = 0; i < 3; i++)
-		res[2 - i] = (float)c[i] / 255.f * 2.f - 1.f;
-	return res;
-}
-
-Vec2f ColladaModel::uv(int iface, int nthvert) {
-	return textureco_[faces_[iface][nthvert][2]];
-}
-
-float ColladaModel::specular(Vec2f uvf) {
-	Vec2i uv(uvf[0] * specularmap_.get_width(), uvf[1] * specularmap_.get_height());
-	return specularmap_.get(uv[0], uv[1])[0] / 1.f;
-}
-
-Vec3f ColladaModel::normal(int iface, int nthvert) {
-	int idx = faces_[iface][nthvert][1];
-	return normals_[idx].normalize();
-}
@@ -1,66 +0,0 @@
-#ifndef __MODEL_s__
-#define __MODEL_s__
-#pragma once
-
-#include "ColladaModel.h"
-#include "tinyxml2.h"
-#include "geometry.h"
-#include "tgaimage.h"
-
-#include <cerrno>
-#include <cstdlib>
-#include <cstring>
-#include <ctime>
-#include <fstream>
-#include <iostream>
-#include <string>
-#include <sstream>
-#include <vector>
-
-class ColladaModel {
-private:
-	int face_count;
-	int vertex_count;
-	int normal_count;
-	int texcoord_count;
-
-	std::vector<std::vector<Vec3i> > faces_; //vertex/normal/uv
-	std::vector<Vec3f> vertices_;
-	std::vector<Vec3f> normals_;
-	std::vector<Vec2f> textureco_;
-
-	TGAImage diffusemap_;
-	TGAImage normalmap_;
-	TGAImage specularmap_;
-
-	void load_texture(std::string filename, const char* suffix, TGAImage& img);
-public:
-	ColladaModel(const char* filename);
-	~ColladaModel();
-
-	int nfaces();
-	int nvertices();
-
-	Vec3f vertix(int iface, int nthvert);
-	Vec3f vertix(int i);
-
-	Vec3f normal(int iface, int nthvert);
-	Vec3f normal(Vec2f uv);
-
-	Vec2f uv(int iface, int nthvert);
-
-	Matrix Transform;
-	Matrix Rotation;
-	Matrix Scale;
-	Matrix Translation;
-
-	void translate(Vec3f tr);
-	void rotate(Vec3f rot);
-	void scale(Vec3f scl);
-	void ApplyTransform();
-
-	TGAColor diffuse(Vec2f uv);
-	float specular(Vec2f uv);
-	std::vector<int> face(int idx);
-};
-#endif
@@ -1,9 +0,0 @@
-#include "GeometryLibrary.h"
-
-GeometryLibrary::GeometryLibrary(tinyxml2::XMLElement* xml)
-{
-	auto xm = xml;
-	auto g = Geometry(xml->FirstChildElement("geometry"));
-	m_Geometries.push_back(g);
-}
-
@@ -1,22 +0,0 @@
-#pragma once
-#include <vector>
-#include "geometry.h"
-
-class GeometryLibrary
-{
-	std::vector<Geometry>    m_Geometries;
-
-public:
-	GeometryLibrary() = default;
-	GeometryLibrary(tinyxml2::XMLElement * xml);
-};
-
-class VisualSceneLibrary
-{
-	Visualscene    m_visual_scene;
-
-public:
-	VisualSceneLibrary() = default;
-	VisualSceneLibrary(tinyxml2::XMLElement* xml);
-};
-
@@ -1,4 +1,4 @@
-<?xml version="1.0" encoding="utf-8"?>
+<?xml version="1.0" encoding="utf-8"?>
 <Project DefaultTargets="Build" ToolsVersion="15.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
  <ItemGroup Label="ProjectConfigurations">
    <ProjectConfiguration Include="Debug|Win32">
@@ -22,32 +22,32 @@
    <VCProjectVersion>15.0</VCProjectVersion>
    <ProjectGuid>{42971E68-F861-4D45-9DA6-F5E163705584}</ProjectGuid>
    <RootNamespace>OpenWindow</RootNamespace>
-    <WindowsTargetPlatformVersion>10.0</WindowsTargetPlatformVersion>
+    <WindowsTargetPlatformVersion>10.0.18362.0</WindowsTargetPlatformVersion>
  </PropertyGroup>
  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'" Label="Configuration">
    <ConfigurationType>Application</ConfigurationType>
    <UseDebugLibraries>true</UseDebugLibraries>
-    <PlatformToolset>v142</PlatformToolset>
+    <PlatformToolset>v141</PlatformToolset>
    <CharacterSet>MultiByte</CharacterSet>
  </PropertyGroup>
  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration">
    <ConfigurationType>Application</ConfigurationType>
    <UseDebugLibraries>false</UseDebugLibraries>
-    <PlatformToolset>v142</PlatformToolset>
+    <PlatformToolset>v141</PlatformToolset>
    <WholeProgramOptimization>true</WholeProgramOptimization>
    <CharacterSet>MultiByte</CharacterSet>
  </PropertyGroup>
  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" Label="Configuration">
    <ConfigurationType>Application</ConfigurationType>
    <UseDebugLibraries>true</UseDebugLibraries>
-    <PlatformToolset>v142</PlatformToolset>
+    <PlatformToolset>v141</PlatformToolset>
    <CharacterSet>MultiByte</CharacterSet>
  </PropertyGroup>
  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="Configuration">
    <ConfigurationType>Application</ConfigurationType>
    <UseDebugLibraries>false</UseDebugLibraries>
-    <PlatformToolset>v142</PlatformToolset>
+    <PlatformToolset>v141</PlatformToolset>
    <WholeProgramOptimization>true</WholeProgramOptimization>
    <CharacterSet>MultiByte</CharacterSet>
  </PropertyGroup>
@@ -76,7 +76,14 @@
      <Optimization>Disabled</Optimization>
      <SDLCheck>true</SDLCheck>
      <ConformanceMode>true</ConformanceMode>
+      <AdditionalIncludeDirectories>$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
+      <PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
    </ClCompile>
+    <Link>
+      <AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
+      <SubSystem>NotSet</SubSystem>
+    </Link>
  </ItemDefinitionGroup>
  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">
    <ClCompile>
@@ -84,7 +91,14 @@
      <Optimization>Disabled</Optimization>
      <SDLCheck>true</SDLCheck>
      <ConformanceMode>true</ConformanceMode>
+      <AdditionalIncludeDirectories>$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
+      <PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
    </ClCompile>
+    <Link>
+      <AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
+      <SubSystem>NotSet</SubSystem>
+    </Link>
  </ItemDefinitionGroup>
  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
    <ClCompile>
@@ -94,10 +108,15 @@
      <IntrinsicFunctions>true</IntrinsicFunctions>
      <SDLCheck>true</SDLCheck>
      <ConformanceMode>true</ConformanceMode>
+      <AdditionalIncludeDirectories>$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
+      <PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
    </ClCompile>
    <Link>
      <EnableCOMDATFolding>true</EnableCOMDATFolding>
      <OptimizeReferences>true</OptimizeReferences>
+      <AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
+      <SubSystem>NotSet</SubSystem>
    </Link>
  </ItemDefinitionGroup>
  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'">
@@ -108,35 +127,44 @@
      <IntrinsicFunctions>true</IntrinsicFunctions>
      <SDLCheck>true</SDLCheck>
      <ConformanceMode>true</ConformanceMode>
+      <AdditionalIncludeDirectories>$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
+      <PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
    </ClCompile>
    <Link>
      <EnableCOMDATFolding>true</EnableCOMDATFolding>
      <OptimizeReferences>true</OptimizeReferences>
+      <AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
+      <SubSystem>NotSet</SubSystem>
    </Link>
  </ItemDefinitionGroup>
  <ItemGroup>
-    <ClCompile Include="COLLADA.cpp" />
-    <ClCompile Include="ColladaModel.cpp" />
-    <ClCompile Include="data.cpp" />
+    <ClCompile Include="camera.cpp" />
    <ClCompile Include="geometry.cpp" />
-    <ClCompile Include="GeometryLibrary.cpp" />
+    <ClCompile Include="improv_gfx.cpp" />
+    <ClCompile Include="kernels.cpp">
+      <AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Release|x64'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+    </ClCompile>
+    <ClCompile Include="kernel_sources.cpp" />
    <ClCompile Include="main.cpp" />
    <ClCompile Include="model.cpp" />
    <ClCompile Include="renderer.cpp" />
    <ClCompile Include="tgaimage.cpp" />
-    <ClCompile Include="tinyxml2.cpp" />
+    <ClCompile Include="util_renderer.cpp" />
    <ClCompile Include="util_window.cpp" />
  </ItemGroup>
  <ItemGroup>
-    <ClInclude Include="COLLADA.h" />
-    <ClInclude Include="ColladaModel.h" />
-    <ClInclude Include="data.h" />
    <ClInclude Include="geometry.h" />
-    <ClInclude Include="GeometryLibrary.h" />
+    <ClInclude Include="camera.h" />
+    <ClInclude Include="improv_gfx.h" />
+    <ClInclude Include="kernels.h" />
    <ClInclude Include="model.h" />
    <ClInclude Include="renderer.h" />
    <ClInclude Include="tgaimage.h" />
-    <ClInclude Include="tinyxml2.h" />
+    <ClInclude Include="util_renderer.h" />
    <ClInclude Include="util_window.h" />
  </ItemGroup>
  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
@@ -33,19 +33,19 @@
    <ClCompile Include="renderer.cpp">
      <Filter>Source Files</Filter>
    </ClCompile>
-    <ClCompile Include="ColladaModel.cpp">
+    <ClCompile Include="camera.cpp">
      <Filter>Source Files</Filter>
    </ClCompile>
-    <ClCompile Include="tinyxml2.cpp">
+    <ClCompile Include="util_renderer.cpp">
      <Filter>Source Files</Filter>
    </ClCompile>
-    <ClCompile Include="COLLADA.cpp">
+    <ClCompile Include="kernels.cpp">
      <Filter>Source Files</Filter>
    </ClCompile>
-    <ClCompile Include="GeometryLibrary.cpp">
+    <ClCompile Include="kernel_sources.cpp">
      <Filter>Source Files</Filter>
    </ClCompile>
-    <ClCompile Include="data.cpp">
+    <ClCompile Include="improv_gfx.cpp">
      <Filter>Source Files</Filter>
    </ClCompile>
  </ItemGroup>
@@ -65,19 +65,16 @@
    <ClInclude Include="renderer.h">
      <Filter>Header Files</Filter>
    </ClInclude>
-    <ClInclude Include="ColladaModel.h">
+    <ClInclude Include="camera.h">
      <Filter>Header Files</Filter>
    </ClInclude>
-    <ClInclude Include="tinyxml2.h">
+    <ClInclude Include="util_renderer.h">
      <Filter>Header Files</Filter>
    </ClInclude>
-    <ClInclude Include="COLLADA.h">
+    <ClInclude Include="kernels.h">
      <Filter>Header Files</Filter>
    </ClInclude>
-    <ClInclude Include="GeometryLibrary.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
-    <ClInclude Include="data.h">
+    <ClInclude Include="improv_gfx.h">
      <Filter>Header Files</Filter>
    </ClInclude>
  </ItemGroup>
@@ -0,0 +1,124 @@
+#include "camera.h"
+#include "util_window.h"
+
+Camera::Camera()
+{
+	fov = 30;
+	position = Vec3f(0, 0, 0);
+	rotation = Vec3f(0, 0, 0);
+	up = Vec3f(0, 0, 0);
+	right = Vec3f(0, 0, 0);
+	forward = Vec3f(0, 0, 0);
+
+	near_plane = 0;
+	far_plane = 15;
+	horizontal_camera_speed = 0.5;
+	vertical_camera_speed = 0.5;
+	vertical_camera_clamp_up = 90;
+	vertical_camera_clamp_down = -90;
+	movement_speed = 1.f;
+}
+
+Vec3f Camera::GetForward() {
+	return forward;
+}
+
+void Camera::Move(Vec3f move_vec) { 
+	position = position + move_vec; 
+}
+
+void Camera::SetPosition(Vec3f pos) { 
+	position = pos; 
+}
+void Camera::Rotate(Vec3f rot_vec) { 
+	rotation = rotation + rot_vec; 
+}
+void Camera::SetRotation(Vec3f rot) { 
+	rotation = rot; 
+}
+void Camera::SetFOV(int angle) { 
+	fov = angle; 
+}
+void Camera::SetVerticalRotSpeed(float speed) {
+	vertical_camera_speed = speed;
+}
+void Camera::SetHorizontalRotSpeed(float speed) {
+	horizontal_camera_speed = speed;
+}
+void Camera::SetClampRotUp(float angle) {
+	vertical_camera_clamp_up = angle;
+}
+void Camera::SetClampRotDown(float angle) {
+	vertical_camera_clamp_down = angle;
+}
+void Camera::SetNearPlane(float near_val) { 
+	near_plane = near_val; 
+}
+void Camera::SetFarPlane(float far_val) { 
+	far_plane = far_val; 
+}
+
+void Camera::rotate_hor(float d_angle) {
+	rotation.y += d_angle * horizontal_camera_speed;
+}
+void Camera::rotate_ver(float d_angle) {
+	rotation.x += d_angle * vertical_camera_speed;
+	rotation.x = std::fmin(rotation.x, vertical_camera_clamp_up);
+	rotation.x = std::fmax(rotation.x, vertical_camera_clamp_down);
+}
+
+
+void Camera::move_camera_left() { 
+	position = position - right * movement_speed;
+}
+void Camera::move_camera_right() { 
+	position = position + right * movement_speed;
+}
+void Camera::move_camera_forward() {
+	position = position + forward * movement_speed;
+}
+void Camera::move_camera_backward() {
+	position = position - forward * movement_speed;
+}
+void Camera::rise() {
+	position = position + Vec3f(0, movement_speed, 0);
+}
+void Camera::fall() {
+	position = position - Vec3f(0, movement_speed, 0);
+}
+
+void Camera::SetMovementSpeed(float speed) {
+	movement_speed = speed;
+}
+
+void Camera::ApplyChanges() {
+	forward = Vec3f(sin(rotation.y * DEG2RAD), -sin(rotation.x * DEG2RAD), -cosf(rotation.y*DEG2RAD) * cosf(rotation.x*DEG2RAD)).normalize();
+	right = Vec3f(cos(rotation.y*DEG2RAD), 0, sin(rotation.y * DEG2RAD)).normalize();
+	up = cross(right, forward).normalize();
+}
+
+Matrix Camera::GetModelViewMatrix() {
+	Vec3f center = position + forward;
+	Vec3f z = forward * -1;
+	Matrix Minv = Matrix::identity();
+	for (int i = 0; i < 3; i++) {
+		Minv[0][i] = right[i];
+		Minv[1][i] = up[i];
+		Minv[2][i] = z[i];
+	}
+	Minv[0][3] = -(right * position);
+	Minv[1][3] = -(up * position);
+	Minv[2][3] = -(z * position);
+	return Minv;
+}
+
+Matrix Camera::GetProjectionMatrix() {
+	Matrix Projection = Matrix::identity();
+	Projection[0][0] = 1 / tan(fov * DEG2RAD / 2);
+	Projection[1][1] = 1 / tan(fov * DEG2RAD / 2);
+	Projection[2][2] = (far_plane + near_plane) / (far_plane - near_plane);
+	Projection[2][3] = (2 * far_plane * near_plane) / (far_plane - near_plane);
+	Projection[3][2] = -1;
+	Projection[3][3] = 0;
+	return Projection;
+}
@@ -0,0 +1,50 @@
+#ifndef __CAMERA_HEADER__
+#define __CAMERA_HEADER__
+#include "geometry.h"
+
+class Camera {
+private:
+	Vec3f position;
+	Vec3f rotation;
+	Vec3f up;
+	Vec3f right;
+	Vec3f forward;
+	float fov;
+	float near_plane;
+	float far_plane;
+	float horizontal_camera_speed;
+	float vertical_camera_speed;
+	float vertical_camera_clamp_up;
+	float vertical_camera_clamp_down;
+	float movement_speed;
+
+public:
+	Camera();
+	void SetPosition(Vec3f pos);
+	void SetRotation(Vec3f rot);
+	void Move(Vec3f move_vec);
+	void Rotate(Vec3f rot_vec);
+	void SetFOV(int angle);
+	void SetNearPlane(float near_val);
+	void SetFarPlane(float far_val);
+	void SetVerticalRotSpeed(float speed);
+	void SetHorizontalRotSpeed(float speed);
+	void SetClampRotUp(float angle);
+	void SetClampRotDown(float angle);
+	void SetMovementSpeed(float speed);
+	Vec3f GetForward();
+
+	void rotate_hor(float d_angle);
+	void rotate_ver(float d_angle);
+	void move_camera_left();
+	void move_camera_right();
+	void move_camera_forward();
+	void move_camera_backward();
+	void rise();
+	void fall();
+
+	void ApplyChanges();
+	Matrix GetModelViewMatrix();
+	Matrix GetProjectionMatrix();
+};
+#endif
@@ -1,187 +0,0 @@
-#include "data.h"
-
-Float_Array::Float_Array(tinyxml2::XMLElement* xml)
-{
-	xml->QueryIntAttribute("count", (int*)&m_Count);
-
-	const char* ch;
-	xml->QueryStringAttribute("count", &ch);
-	std::string s(ch);
-	m_ID = s;
-
-	std::stringstream str(xml->GetText());
-	m_Floats = new float[m_Count];
-
-	for (int i = 0; i < m_Count; i++)
-	{
-		 str >> m_Floats[i];
-	}
-	float f = m_Floats[55358];
-}
-
-Name_Array::Name_Array(tinyxml2::XMLElement* xml) 
-{
-	xml->QueryIntAttribute("count", (int*)&m_Count);
-
-	const char* ch;
-	xml->QueryStringAttribute("count", &ch);
-	std::string s(ch);
-	m_ID = s;
-
-	std::stringstream str(xml->GetText());
-	for (int i = 0; i < m_Count; i++)
-	{
-		str >> m_Names[i];
-	}
-}
-
-Param::Param(tinyxml2::XMLElement* xml) 
-{
-	const char* ch;
-	xml->QueryStringAttribute("name", &ch);
-	std::string s(ch);
-	m_Name = s;
-
-	const char* ch1;
-	xml->QueryStringAttribute("type", &ch1);
-	std::string s1(ch1);
-	m_Type = s1;
-}
-
-Accessor::Accessor(tinyxml2::XMLElement* xml) 
-{
-	const char* ch;
-	xml->QueryStringAttribute("source", &ch);
-	std::string s(ch);
-	m_Source = s;
-
-	xml->QueryIntAttribute("count", (int*)&m_Count);
-
-	xml->QueryIntAttribute("stride", (int*)&m_Stride);
-
-	for (tinyxml2::XMLElement* child = xml->FirstChildElement("param"); child != NULL; child = child->NextSiblingElement("param"))
-	{
-		Param p(child);
-		m_Params.push_back(p);
-	}
-}
-
-Technique_Common::Technique_Common(tinyxml2::XMLElement* xml)
-{
-	m_Accessor = Accessor(xml->FirstChildElement("accessor"));
-}
-
-Source::Source(tinyxml2::XMLElement* xml)
-{
-	const char* ch;
-	xml->QueryStringAttribute("id", &ch);
-	std::string s(ch);
-	m_ID = s;
-
-	tinyxml2::XMLElement* fa = xml->FirstChildElement("float_array");
-	if (fa != NULL)
-	{
-		m_Float_Array = Float_Array(fa);
-	}
-
-	tinyxml2::XMLElement* na = xml->FirstChildElement("Name_array");
-	if (na != NULL)
-	{
-		m_Name_Array = Name_Array(na);
-	}
-
-	m_TechniqueCommon = Technique_Common(xml->FirstChildElement("technique_common"));
-}
-
-Input::Input(tinyxml2::XMLElement* xml)
-{
-	const char* ch;
-	xml->QueryStringAttribute("semantic", &ch);
-	std::string s(ch);
-	m_Semantic = s;
-
-	const char* ch1;
-	xml->QueryStringAttribute("source", &ch1);
-	std::string s1(ch1);
-	m_Source = s1;
-	
-	if (xml->FindAttribute("offset"))
-	{
-		const char* ch2;
-		xml->QueryStringAttribute("offset", &ch2);
-		std::string s2(ch2);
-		m_Offset = s2;
-	}
-
-	if (xml->FindAttribute("set"))
-	{
-		const char* ch3;
-		xml->QueryStringAttribute("set", &ch3);
-		std::string s3(ch3);
-		m_Set = s3;
-	}
-}
-
-Vertices::Vertices(tinyxml2::XMLElement* xml)
-{
-	const char* ch;
-	xml->QueryStringAttribute("id", &ch);
-	std::string s(ch);
-	m_ID = s;
-
-	for (tinyxml2::XMLElement* child = xml->FirstChildElement("input"); child != NULL; child = child->NextSiblingElement("input"))
-	{
-		Input in(child);
-		m_Inputs.push_back(in);
-	}
-}
-
-P::P(tinyxml2::XMLElement* xml, unsigned int count)
-{
-	m_Count = count;
-	std::stringstream str(xml->GetText());
-	m_Indices = new unsigned short[count * 9];
-	for (int i = 0; i < count * 9; i++)
-	{
-		str >> m_Indices[i];
-	}
-}
-
-Triangles::Triangles(tinyxml2::XMLElement* xml)
-{
-	xml->QueryIntAttribute("count", (int*)&m_Count);
-
-	const char* ch;
-	xml->QueryStringAttribute("count", &ch);
-	std::string s(ch);
-	m_Material = s;
-
-	for (tinyxml2::XMLElement* child = xml->FirstChildElement("input"); child != NULL; child = child->NextSiblingElement("input"))
-	{
-		Input in(child);
-		m_Inputs.push_back(in);
-	}
-
-	tinyxml2::XMLElement* mp = xml->FirstChildElement("p");
-	if (mp != NULL)
-	{
-		m_P = P(mp, m_Count);
-	}
-}
-
-Mesh::Mesh(tinyxml2::XMLElement* xml)
-{
-	for (tinyxml2::XMLElement* child = xml->FirstChildElement("source"); child != NULL; child = child->NextSiblingElement("source"))
-	{
-		Source sc(child);
-		m_Sources.push_back(sc);
-	}
-
-	m_Vertices = Vertices(xml->FirstChildElement("vertices"));
-
-	for (tinyxml2::XMLElement* child = xml->FirstChildElement("triangles"); child != NULL; child = child->NextSiblingElement("triangles"))
-	{
-		Triangles tr(child);
-		m_Triangles.push_back(tr);
-	}
-}
@@ -1,137 +0,0 @@
-#pragma once
-#include <iostream>
-#include <fstream>
-#include <sstream>
-#include <string>
-#include <vector>
-#include "tinyxml2.h"
-//----------------------------------------------//
-class Float_Array
-{
-	unsigned int        m_Count;
-	std::string              m_ID;
-	float* m_Floats;
-
-public:
-	Float_Array() = default;
-	Float_Array(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Name_Array
-{
-	unsigned int        m_Count;
-	std::string              m_ID;
-	std::string* m_Names;
-public:
-	Name_Array() = default;
-	Name_Array(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Param
-{
-	std::string              m_Name;
-	std::string              m_Type;
-public:
-	Param() = default;
-	Param(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Accessor
-{
-	std::string              m_Source;
-	unsigned int        m_Count;
-	unsigned int        m_Stride;
-	std::vector<Param>       m_Params;
-public:
-	Accessor () = default;
-	Accessor(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Technique_Common
-{
-	Accessor            m_Accessor;
-public:
-	Technique_Common() = default;
-	Technique_Common(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Source
-{
-	std::string              m_ID;
-	Float_Array         m_Float_Array;
-	Name_Array          m_Name_Array;
-	Technique_Common    m_TechniqueCommon;
-public:
-	Source() = default;
-	Source(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Input
-{
-	std::string              m_Semantic;
-	std::string              m_Source;
-	std::string              m_Offset;
-	std::string              m_Set;
-public:
-	Input() = default;
-	Input(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Vertices
-{
-	std::string              m_ID;
-	std::vector<Input>       m_Inputs;
-public:
-	Vertices() = default;
-	Vertices(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class P
-{
-	unsigned int         m_Count;
-	unsigned short* m_Indices;
-public:
-	P() = default;
-	P(tinyxml2::XMLElement* xml, unsigned int count);
-};
-//----------------------------------------------//
-class Triangles
-{
-	unsigned int        m_Count;
-	std::string              m_Material;
-	std::vector<Input>       m_Inputs;
-	P                   m_P;
-public:
-	Triangles() = default;
-	Triangles(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Mesh
-{
-	std::vector<Source>      m_Sources;
-	Vertices            m_Vertices;
-	std::vector<Triangles>   m_Triangles;
-public:
-	Mesh() = default;
-	Mesh(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Node
-{
-	std::string      m_ID;
-	std::string      m_name;
-	std::string      m_type;
-	std::string      m_sid;
-public:
-	Node() = default;
-	Node(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
-class Matrix
-{	
-	std::string      m_sid;
-public:
-	Node() = default;
-	Node(tinyxml2::XMLElement* xml);
-};
-//----------------------------------------------//
@@ -5,7 +5,3 @@ template <> template <> vec<3,float>::vec(const vec<3,int> &v)   : x(v.x),y(v.y)
 template <> template <> vec<2,int>  ::vec(const vec<2,float> &v) : x(int(v.x+.5f)),y(int(v.y+.5f)) {}
 template <> template <> vec<2,float>::vec(const vec<2,int> &v)   : x(v.x),y(v.y) {}

-Geometry::Geometry(tinyxml2::XMLElement* xml)
-{
-	m_Mesh = Mesh(xml->FirstChildElement("mesh"));
-}
@@ -1,10 +1,12 @@
 #ifndef __GEOMETRY_H__
 #define __GEOMETRY_H__
+#define _USE_MATH_DEFINES
 #include <cmath>
 #include <vector>
 #include <cassert>
 #include <iostream>
-#include "data.h"
+
+#define DEG2RAD M_PI/180

 template<size_t DimCols,size_t DimRows,typename T> class mat;

@@ -234,46 +236,5 @@ typedef vec<3,  float> Vec3f;
 typedef vec<3,  int>   Vec3i;
 typedef vec<4,  float> Vec4f;
 typedef mat<4,4,float> Matrix;
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-class Geometry
-{
-	Mesh                m_Mesh;
-public:
-	Geometry() = default;
-	Geometry(tinyxml2::XMLElement * xml);
-};
-
-class Visualscene
-{
-	node                m_Mesh;
-public:
-	Visualscene() = default;
-	Visualscene(tinyxml2::XMLElement* xml);
-};
 #endif //__GEOMETRY_H__

@@ -0,0 +1,27 @@
+#include "improv_gfx.h"
+
+Model *octopus, *ocean, *sharky;
+
+void Setup() {
+	octopus = new Model("octojet.obj", 1);
+	octopus->scale(Vec3f(4, 4, 4));
+	//octopus->translate(Vec3f(0, -0.15, 0));
+	octopus->ApplyTransform();
+	AddModel(octopus);
+	ocean = new Model("ocean.obj", 1);
+	ocean->scale(Vec3f(1, 1, 1));
+	ocean->ApplyTransform();
+	AddModel(ocean);
+	sharky = new Model("sharky.obj", 0);
+	sharky->scale(Vec3f(0.02, 0.02, 0.02));
+	sharky->translate(Vec3f(0, 0.05, 0));
+	sharky->ApplyTransform();
+	AddModel(sharky);
+}
+
+void Draw() {
+	sharky->translate(Vec3f(sharky->Translation[0][3] + 0.01*cos(TIME), sharky->Translation[1][3], sharky->Translation[2][3] + 0.01*sin(TIME)));
+	sharky->rotate(Vec3f(0, (cos(TIME) + sin(TIME)) * -90, 0));
+	sharky->ApplyTransform();
+
+}
@@ -0,0 +1,11 @@
+#ifndef __IMPROV_GFX__
+#define __IMPROV_GFX__
+#include "model.h"
+#include "geometry.h"
+#include "renderer.h"
+
+void Setup();
+void Draw();
+void AddModel(Model* model);
+
+#endif
@@ -0,0 +1,126 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+#ifndef __OPENCL_CL_D3D10_H
+#define __OPENCL_CL_D3D10_H
+
+#include <d3d10.h>
+#include <CL/cl.h>
+#include <CL/cl_platform.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/******************************************************************************
+ * cl_khr_d3d10_sharing                                                       */
+#define cl_khr_d3d10_sharing 1
+
+typedef cl_uint cl_d3d10_device_source_khr;
+typedef cl_uint cl_d3d10_device_set_khr;
+
+/******************************************************************************/
+
+/* Error Codes */
+#define CL_INVALID_D3D10_DEVICE_KHR                  -1002
+#define CL_INVALID_D3D10_RESOURCE_KHR                -1003
+#define CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR       -1004
+#define CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR           -1005
+
+/* cl_d3d10_device_source_nv */
+#define CL_D3D10_DEVICE_KHR                          0x4010
+#define CL_D3D10_DXGI_ADAPTER_KHR                    0x4011
+
+/* cl_d3d10_device_set_nv */
+#define CL_PREFERRED_DEVICES_FOR_D3D10_KHR           0x4012
+#define CL_ALL_DEVICES_FOR_D3D10_KHR                 0x4013
+
+/* cl_context_info */
+#define CL_CONTEXT_D3D10_DEVICE_KHR                  0x4014
+#define CL_CONTEXT_D3D10_PREFER_SHARED_RESOURCES_KHR 0x402C
+
+/* cl_mem_info */
+#define CL_MEM_D3D10_RESOURCE_KHR                    0x4015
+
+/* cl_image_info */
+#define CL_IMAGE_D3D10_SUBRESOURCE_KHR               0x4016
+
+/* cl_command_type */
+#define CL_COMMAND_ACQUIRE_D3D10_OBJECTS_KHR         0x4017
+#define CL_COMMAND_RELEASE_D3D10_OBJECTS_KHR         0x4018
+
+/******************************************************************************/
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetDeviceIDsFromD3D10KHR_fn)(
+    cl_platform_id             platform,
+    cl_d3d10_device_source_khr d3d_device_source,
+    void *                     d3d_object,
+    cl_d3d10_device_set_khr    d3d_device_set,
+    cl_uint                    num_entries,
+    cl_device_id *             devices,
+    cl_uint *                  num_devices) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D10BufferKHR_fn)(
+    cl_context     context,
+    cl_mem_flags   flags,
+    ID3D10Buffer * resource,
+    cl_int *       errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D10Texture2DKHR_fn)(
+    cl_context        context,
+    cl_mem_flags      flags,
+    ID3D10Texture2D * resource,
+    UINT              subresource,
+    cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D10Texture3DKHR_fn)(
+    cl_context        context,
+    cl_mem_flags      flags,
+    ID3D10Texture3D * resource,
+    UINT              subresource,
+    cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireD3D10ObjectsKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseD3D10ObjectsKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_0;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_D3D10_H */
+
@@ -0,0 +1,126 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+#ifndef __OPENCL_CL_D3D11_H
+#define __OPENCL_CL_D3D11_H
+
+#include <d3d11.h>
+#include <CL/cl.h>
+#include <CL/cl_platform.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/******************************************************************************
+ * cl_khr_d3d11_sharing                                                       */
+#define cl_khr_d3d11_sharing 1
+
+typedef cl_uint cl_d3d11_device_source_khr;
+typedef cl_uint cl_d3d11_device_set_khr;
+
+/******************************************************************************/
+
+/* Error Codes */
+#define CL_INVALID_D3D11_DEVICE_KHR                  -1006
+#define CL_INVALID_D3D11_RESOURCE_KHR                -1007
+#define CL_D3D11_RESOURCE_ALREADY_ACQUIRED_KHR       -1008
+#define CL_D3D11_RESOURCE_NOT_ACQUIRED_KHR           -1009
+
+/* cl_d3d11_device_source */
+#define CL_D3D11_DEVICE_KHR                          0x4019
+#define CL_D3D11_DXGI_ADAPTER_KHR                    0x401A
+
+/* cl_d3d11_device_set */
+#define CL_PREFERRED_DEVICES_FOR_D3D11_KHR           0x401B
+#define CL_ALL_DEVICES_FOR_D3D11_KHR                 0x401C
+
+/* cl_context_info */
+#define CL_CONTEXT_D3D11_DEVICE_KHR                  0x401D
+#define CL_CONTEXT_D3D11_PREFER_SHARED_RESOURCES_KHR 0x402D
+
+/* cl_mem_info */
+#define CL_MEM_D3D11_RESOURCE_KHR                    0x401E
+
+/* cl_image_info */
+#define CL_IMAGE_D3D11_SUBRESOURCE_KHR               0x401F
+
+/* cl_command_type */
+#define CL_COMMAND_ACQUIRE_D3D11_OBJECTS_KHR         0x4020
+#define CL_COMMAND_RELEASE_D3D11_OBJECTS_KHR         0x4021
+
+/******************************************************************************/
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetDeviceIDsFromD3D11KHR_fn)(
+    cl_platform_id             platform,
+    cl_d3d11_device_source_khr d3d_device_source,
+    void *                     d3d_object,
+    cl_d3d11_device_set_khr    d3d_device_set,
+    cl_uint                    num_entries,
+    cl_device_id *             devices,
+    cl_uint *                  num_devices) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D11BufferKHR_fn)(
+    cl_context     context,
+    cl_mem_flags   flags,
+    ID3D11Buffer * resource,
+    cl_int *       errcode_ret) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D11Texture2DKHR_fn)(
+    cl_context        context,
+    cl_mem_flags      flags,
+    ID3D11Texture2D * resource,
+    UINT              subresource,
+    cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D11Texture3DKHR_fn)(
+    cl_context        context,
+    cl_mem_flags      flags,
+    ID3D11Texture3D * resource,
+    UINT              subresource,
+    cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireD3D11ObjectsKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseD3D11ObjectsKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_2;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_D3D11_H */
+
@@ -0,0 +1,127 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+#ifndef __OPENCL_CL_DX9_MEDIA_SHARING_H
+#define __OPENCL_CL_DX9_MEDIA_SHARING_H
+
+#include <CL/cl.h>
+#include <CL/cl_platform.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/******************************************************************************
+/* cl_khr_dx9_media_sharing                                                   */
+#define cl_khr_dx9_media_sharing 1
+
+typedef cl_uint             cl_dx9_media_adapter_type_khr;
+typedef cl_uint             cl_dx9_media_adapter_set_khr;
+    
+#if defined(_WIN32)
+#include <d3d9.h>
+typedef struct _cl_dx9_surface_info_khr
+{
+    IDirect3DSurface9 *resource;
+    HANDLE shared_handle;
+} cl_dx9_surface_info_khr;
+#endif
+
+
+/******************************************************************************/
+
+/* Error Codes */
+#define CL_INVALID_DX9_MEDIA_ADAPTER_KHR                -1010
+#define CL_INVALID_DX9_MEDIA_SURFACE_KHR                -1011
+#define CL_DX9_MEDIA_SURFACE_ALREADY_ACQUIRED_KHR       -1012
+#define CL_DX9_MEDIA_SURFACE_NOT_ACQUIRED_KHR           -1013
+
+/* cl_media_adapter_type_khr */
+#define CL_ADAPTER_D3D9_KHR                              0x2020
+#define CL_ADAPTER_D3D9EX_KHR                            0x2021
+#define CL_ADAPTER_DXVA_KHR                              0x2022
+
+/* cl_media_adapter_set_khr */
+#define CL_PREFERRED_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR   0x2023
+#define CL_ALL_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR         0x2024
+
+/* cl_context_info */
+#define CL_CONTEXT_ADAPTER_D3D9_KHR                      0x2025
+#define CL_CONTEXT_ADAPTER_D3D9EX_KHR                    0x2026
+#define CL_CONTEXT_ADAPTER_DXVA_KHR                      0x2027
+
+/* cl_mem_info */
+#define CL_MEM_DX9_MEDIA_ADAPTER_TYPE_KHR                0x2028
+#define CL_MEM_DX9_MEDIA_SURFACE_INFO_KHR                0x2029
+
+/* cl_image_info */
+#define CL_IMAGE_DX9_MEDIA_PLANE_KHR                     0x202A
+
+/* cl_command_type */
+#define CL_COMMAND_ACQUIRE_DX9_MEDIA_SURFACES_KHR        0x202B
+#define CL_COMMAND_RELEASE_DX9_MEDIA_SURFACES_KHR        0x202C
+
+/******************************************************************************/
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetDeviceIDsFromDX9MediaAdapterKHR_fn)(
+    cl_platform_id                   platform,
+    cl_uint                          num_media_adapters,
+    cl_dx9_media_adapter_type_khr *  media_adapter_type,
+    void *                           media_adapters,
+    cl_dx9_media_adapter_set_khr     media_adapter_set,
+    cl_uint                          num_entries,
+    cl_device_id *                   devices,
+    cl_uint *                        num_devices) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromDX9MediaSurfaceKHR_fn)(
+    cl_context                    context,
+    cl_mem_flags                  flags,
+    cl_dx9_media_adapter_type_khr adapter_type,
+    void *                        surface_info,
+    cl_uint                       plane,                                                                          
+    cl_int *                      errcode_ret) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireDX9MediaSurfacesKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseDX9MediaSurfacesKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_2;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_DX9_MEDIA_SHARING_H */
+
@@ -0,0 +1,443 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+/* $Revision: 14835 $ on $Date: 2011-05-26 11:32:00 -0700 (Thu, 26 May 2011) $ */
+
+/* cl_ext.h contains OpenCL extensions which don't have external */
+/* (OpenGL, D3D) dependencies.                                   */
+
+#ifndef __CL_EXT_H
+#define __CL_EXT_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifdef __APPLE__
+        #include <OpenCL/cl.h>
+    #include <AvailabilityMacros.h>
+#else
+        #include <CL/cl.h>
+#endif
+
+/* cl_khr_fp16 extension - no extension #define since it has no functions  */
+#define CL_DEVICE_HALF_FP_CONFIG                    0x1033
+
+/* Memory object destruction
+ *
+ * Apple extension for use to manage externally allocated buffers used with cl_mem objects with CL_MEM_USE_HOST_PTR
+ *
+ * Registers a user callback function that will be called when the memory object is deleted and its resources 
+ * freed. Each call to clSetMemObjectCallbackFn registers the specified user callback function on a callback 
+ * stack associated with memobj. The registered user callback functions are called in the reverse order in 
+ * which they were registered. The user callback functions are called and then the memory object is deleted 
+ * and its resources freed. This provides a mechanism for the application (and libraries) using memobj to be 
+ * notified when the memory referenced by host_ptr, specified when the memory object is created and used as 
+ * the storage bits for the memory object, can be reused or freed.
+ *
+ * The application may not call CL api's with the cl_mem object passed to the pfn_notify.
+ *
+ * Please check for the "cl_APPLE_SetMemObjectDestructor" extension using clGetDeviceInfo(CL_DEVICE_EXTENSIONS)
+ * before using.
+ */
+#define cl_APPLE_SetMemObjectDestructor 1
+cl_int  CL_API_ENTRY clSetMemObjectDestructorAPPLE(  cl_mem /* memobj */, 
+                                        void (* /*pfn_notify*/)( cl_mem /* memobj */, void* /*user_data*/), 
+                                        void * /*user_data */ )             CL_EXT_SUFFIX__VERSION_1_0;  
+
+
+/* Context Logging Functions
+ *
+ * The next three convenience functions are intended to be used as the pfn_notify parameter to clCreateContext().
+ * Please check for the "cl_APPLE_ContextLoggingFunctions" extension using clGetDeviceInfo(CL_DEVICE_EXTENSIONS)
+ * before using.
+ *
+ * clLogMessagesToSystemLog fowards on all log messages to the Apple System Logger 
+ */
+#define cl_APPLE_ContextLoggingFunctions 1
+extern void CL_API_ENTRY clLogMessagesToSystemLogAPPLE(  const char * /* errstr */, 
+                                            const void * /* private_info */, 
+                                            size_t       /* cb */, 
+                                            void *       /* user_data */ )  CL_EXT_SUFFIX__VERSION_1_0;
+
+/* clLogMessagesToStdout sends all log messages to the file descriptor stdout */
+extern void CL_API_ENTRY clLogMessagesToStdoutAPPLE(   const char * /* errstr */, 
+                                          const void * /* private_info */, 
+                                          size_t       /* cb */, 
+                                          void *       /* user_data */ )    CL_EXT_SUFFIX__VERSION_1_0;
+
+/* clLogMessagesToStderr sends all log messages to the file descriptor stderr */
+extern void CL_API_ENTRY clLogMessagesToStderrAPPLE(   const char * /* errstr */, 
+                                          const void * /* private_info */, 
+                                          size_t       /* cb */, 
+                                          void *       /* user_data */ )    CL_EXT_SUFFIX__VERSION_1_0;
+
+
+/************************ 
+* cl_khr_icd extension *                                                  
+************************/
+#define cl_khr_icd 1
+
+/* cl_platform_info                                                        */
+#define CL_PLATFORM_ICD_SUFFIX_KHR                  0x0920
+
+/* Additional Error Codes                                                  */
+#define CL_PLATFORM_NOT_FOUND_KHR                   -1001
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clIcdGetPlatformIDsKHR(cl_uint          /* num_entries */,
+                       cl_platform_id * /* platforms */,
+                       cl_uint *        /* num_platforms */);
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clIcdGetPlatformIDsKHR_fn)(
+    cl_uint          /* num_entries */,
+    cl_platform_id * /* platforms */,
+    cl_uint *        /* num_platforms */);
+
+
+/* Extension: cl_khr_image2D_buffer
+ *
+ * This extension allows a 2D image to be created from a cl_mem buffer without a copy.
+ * The type associated with a 2D image created from a buffer in an OpenCL program is image2d_t.
+ * Both the sampler and sampler-less read_image built-in functions are supported for 2D images
+ * and 2D images created from a buffer.  Similarly, the write_image built-ins are also supported
+ * for 2D images created from a buffer.
+ *
+ * When the 2D image from buffer is created, the client must specify the width,
+ * height, image format (i.e. channel order and channel data type) and optionally the row pitch
+ *
+ * The pitch specified must be a multiple of CL_DEVICE_IMAGE_PITCH_ALIGNMENT pixels.
+ * The base address of the buffer must be aligned to CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT pixels.
+ */
+    
+/*************************************
+ * cl_khr_initalize_memory extension *
+ *************************************/
+    
+#define CL_CONTEXT_MEMORY_INITIALIZE_KHR            0x2030
+    
+    
+/**************************************
+ * cl_khr_terminate_context extension *
+ **************************************/
+    
+#define CL_DEVICE_TERMINATE_CAPABILITY_KHR          0x2031
+#define CL_CONTEXT_TERMINATE_KHR                    0x2032
+
+#define cl_khr_terminate_context 1
+extern CL_API_ENTRY cl_int CL_API_CALL clTerminateContextKHR(cl_context /* context */) CL_EXT_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clTerminateContextKHR_fn)(cl_context /* context */) CL_EXT_SUFFIX__VERSION_1_2;
+    
+/*
+ * Extension: cl_khr_spir
+ *
+ * This extension adds support to create an OpenCL program object from a 
+ * Standard Portable Intermediate Representation (SPIR) instance
+ */
+
+#define CL_DEVICE_SPIR_VERSIONS                     0x40E0
+#define CL_PROGRAM_BINARY_TYPE_INTERMEDIATE         0x40E1
+
+
+/******************************************
+* cl_nv_device_attribute_query extension *
+******************************************/
+/* cl_nv_device_attribute_query extension - no extension #define since it has no functions */
+#define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV       0x4000
+#define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV       0x4001
+#define CL_DEVICE_REGISTERS_PER_BLOCK_NV            0x4002
+#define CL_DEVICE_WARP_SIZE_NV                      0x4003
+#define CL_DEVICE_GPU_OVERLAP_NV                    0x4004
+#define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV            0x4005
+#define CL_DEVICE_INTEGRATED_MEMORY_NV              0x4006
+
+/*********************************
+* cl_amd_device_memory_flags *
+*********************************/
+#define cl_amd_device_memory_flags 1
+#define CL_MEM_USE_PERSISTENT_MEM_AMD       (1 << 6)        // Alloc from GPU's CPU visible heap
+
+/* cl_device_info */
+#define CL_DEVICE_MAX_ATOMIC_COUNTERS_EXT           0x4032
+
+/*********************************
+* cl_amd_device_attribute_query *
+*********************************/
+#define CL_DEVICE_PROFILING_TIMER_OFFSET_AMD        0x4036
+#define CL_DEVICE_TOPOLOGY_AMD                      0x4037
+#define CL_DEVICE_BOARD_NAME_AMD                    0x4038
+#define CL_DEVICE_GLOBAL_FREE_MEMORY_AMD            0x4039
+#define CL_DEVICE_SIMD_PER_COMPUTE_UNIT_AMD         0x4040
+#define CL_DEVICE_SIMD_WIDTH_AMD                    0x4041
+#define CL_DEVICE_SIMD_INSTRUCTION_WIDTH_AMD        0x4042
+#define CL_DEVICE_WAVEFRONT_WIDTH_AMD               0x4043
+#define CL_DEVICE_GLOBAL_MEM_CHANNELS_AMD           0x4044
+#define CL_DEVICE_GLOBAL_MEM_CHANNEL_BANKS_AMD      0x4045
+#define CL_DEVICE_GLOBAL_MEM_CHANNEL_BANK_WIDTH_AMD 0x4046
+#define CL_DEVICE_LOCAL_MEM_SIZE_PER_COMPUTE_UNIT_AMD   0x4047
+#define CL_DEVICE_LOCAL_MEM_BANKS_AMD               0x4048
+#define CL_DEVICE_THREAD_TRACE_SUPPORTED_AMD        0x4049
+#define CL_DEVICE_GFXIP_MAJOR_AMD                   0x404A
+#define CL_DEVICE_GFXIP_MINOR_AMD                   0x404B
+#define CL_DEVICE_AVAILABLE_ASYNC_QUEUES_AMD        0x404C
+
+typedef union
+{
+    struct { cl_uint type; cl_uint data[5]; } raw;
+    struct { cl_uint type; cl_char unused[17]; cl_char bus; cl_char device; cl_char function; } pcie;
+} cl_device_topology_amd;
+
+#define CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD            1
+
+/*************
+* cl_amd_hsa *
+**************/
+#define CL_HSA_ENABLED_AMD                          (1ull << 62)
+#define CL_HSA_DISABLED_AMD                         (1ull << 63)
+
+
+/**************************
+* cl_amd_offline_devices *
+**************************/
+#define CL_CONTEXT_OFFLINE_DEVICES_AMD              0x403F
+
+/*********************************
+* cl_arm_printf extension
+*********************************/
+#define CL_PRINTF_CALLBACK_ARM                      0x40B0
+#define CL_PRINTF_BUFFERSIZE_ARM                    0x40B1
+
+#ifdef CL_VERSION_1_1
+   /***********************************
+    * cl_ext_device_fission extension *
+    ***********************************/
+    #define cl_ext_device_fission   1
+    
+    extern CL_API_ENTRY cl_int CL_API_CALL
+    clReleaseDeviceEXT( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1; 
+    
+    typedef CL_API_ENTRY cl_int 
+    (CL_API_CALL *clReleaseDeviceEXT_fn)( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    extern CL_API_ENTRY cl_int CL_API_CALL
+    clRetainDeviceEXT( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1; 
+    
+    typedef CL_API_ENTRY cl_int 
+    (CL_API_CALL *clRetainDeviceEXT_fn)( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    typedef cl_ulong  cl_device_partition_property_ext;
+    extern CL_API_ENTRY cl_int CL_API_CALL
+    clCreateSubDevicesEXT(  cl_device_id /*in_device*/,
+                            const cl_device_partition_property_ext * /* properties */,
+                            cl_uint /*num_entries*/,
+                            cl_device_id * /*out_devices*/,
+                            cl_uint * /*num_devices*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    typedef CL_API_ENTRY cl_int 
+    ( CL_API_CALL * clCreateSubDevicesEXT_fn)(  cl_device_id /*in_device*/,
+                                                const cl_device_partition_property_ext * /* properties */,
+                                                cl_uint /*num_entries*/,
+                                                cl_device_id * /*out_devices*/,
+                                                cl_uint * /*num_devices*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    /* cl_device_partition_property_ext */
+    #define CL_DEVICE_PARTITION_EQUALLY_EXT             0x4050
+    #define CL_DEVICE_PARTITION_BY_COUNTS_EXT           0x4051
+    #define CL_DEVICE_PARTITION_BY_NAMES_EXT            0x4052
+    #define CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN_EXT  0x4053
+    
+    /* clDeviceGetInfo selectors */
+    #define CL_DEVICE_PARENT_DEVICE_EXT                 0x4054
+    #define CL_DEVICE_PARTITION_TYPES_EXT               0x4055
+    #define CL_DEVICE_AFFINITY_DOMAINS_EXT              0x4056
+    #define CL_DEVICE_REFERENCE_COUNT_EXT               0x4057
+    #define CL_DEVICE_PARTITION_STYLE_EXT               0x4058
+    
+    /* clGetImageInfo enum */
+    #define CL_IMAGE_BYTE_PITCH_AMD                     0x4059
+
+    /* error codes */
+    #define CL_DEVICE_PARTITION_FAILED_EXT              -1057
+    #define CL_INVALID_PARTITION_COUNT_EXT              -1058
+    #define CL_INVALID_PARTITION_NAME_EXT               -1059
+    
+    /* CL_AFFINITY_DOMAINs */
+    #define CL_AFFINITY_DOMAIN_L1_CACHE_EXT             0x1
+    #define CL_AFFINITY_DOMAIN_L2_CACHE_EXT             0x2
+    #define CL_AFFINITY_DOMAIN_L3_CACHE_EXT             0x3
+    #define CL_AFFINITY_DOMAIN_L4_CACHE_EXT             0x4
+    #define CL_AFFINITY_DOMAIN_NUMA_EXT                 0x10
+    #define CL_AFFINITY_DOMAIN_NEXT_FISSIONABLE_EXT     0x100
+    /* cl_device_partition_property_ext list terminators */
+    #define CL_PROPERTIES_LIST_END_EXT                  ((cl_device_partition_property_ext) 0)
+    #define CL_PARTITION_BY_COUNTS_LIST_END_EXT         ((cl_device_partition_property_ext) 0)
+    #define CL_PARTITION_BY_NAMES_LIST_END_EXT          ((cl_device_partition_property_ext) 0 - 1)
+
+/*********************************
+* cl_qcom_ext_host_ptr extension
+*********************************/
+
+#define CL_MEM_EXT_HOST_PTR_QCOM                  (1 << 29)
+
+#define CL_DEVICE_EXT_MEM_PADDING_IN_BYTES_QCOM   0x40A0      
+#define CL_DEVICE_PAGE_SIZE_QCOM                  0x40A1
+#define CL_IMAGE_ROW_ALIGNMENT_QCOM               0x40A2
+#define CL_IMAGE_SLICE_ALIGNMENT_QCOM             0x40A3
+#define CL_MEM_HOST_UNCACHED_QCOM                 0x40A4
+#define CL_MEM_HOST_WRITEBACK_QCOM                0x40A5
+#define CL_MEM_HOST_WRITETHROUGH_QCOM             0x40A6
+#define CL_MEM_HOST_WRITE_COMBINING_QCOM          0x40A7
+
+typedef cl_uint                                   cl_image_pitch_info_qcom;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetDeviceImageInfoQCOM(cl_device_id             device,
+                         size_t                   image_width,
+                         size_t                   image_height,
+                         const cl_image_format   *image_format,
+                         cl_image_pitch_info_qcom param_name,
+                         size_t                   param_value_size,
+                         void                    *param_value,
+                         size_t                  *param_value_size_ret);
+
+typedef struct _cl_mem_ext_host_ptr
+{
+    /* Type of external memory allocation. */
+    /* Legal values will be defined in layered extensions. */
+    cl_uint  allocation_type;
+            
+    /* Host cache policy for this external memory allocation. */
+    cl_uint  host_cache_policy;
+
+} cl_mem_ext_host_ptr;
+
+/*********************************
+* cl_qcom_ion_host_ptr extension
+*********************************/
+
+#define CL_MEM_ION_HOST_PTR_QCOM                  0x40A8
+
+typedef struct _cl_mem_ion_host_ptr
+{
+    /* Type of external memory allocation. */
+    /* Must be CL_MEM_ION_HOST_PTR_QCOM for ION allocations. */
+    cl_mem_ext_host_ptr  ext_host_ptr;
+
+    /* ION file descriptor */
+    int                  ion_filedesc;
+            
+    /* Host pointer to the ION allocated memory */
+    void*                ion_hostptr;
+
+} cl_mem_ion_host_ptr;
+
+
+#endif /* CL_VERSION_1_1 */
+
+#ifdef CL_VERSION_1_2
+/********************************
+* cl_amd_bus_addressable_memory *
+********************************/
+
+/* cl_mem flag - bitfield */
+#define CL_MEM_BUS_ADDRESSABLE_AMD               (1<<30)
+#define CL_MEM_EXTERNAL_PHYSICAL_AMD             (1<<31)
+
+#define CL_COMMAND_WAIT_SIGNAL_AMD                0x4080
+#define CL_COMMAND_WRITE_SIGNAL_AMD               0x4081
+#define CL_COMMAND_MAKE_BUFFERS_RESIDENT_AMD      0x4082
+
+typedef struct _cl_bus_address_amd
+{
+    cl_ulong surface_bus_address;
+    cl_ulong marker_bus_address;
+} cl_bus_address_amd;
+
+typedef CL_API_ENTRY cl_int
+(CL_API_CALL * clEnqueueWaitSignalAMD_fn)( cl_command_queue /*command_queue*/,
+                                           cl_mem /*mem_object*/,
+                                           cl_uint /*value*/,
+                                           cl_uint /*num_events*/,
+                                           const cl_event * /*event_wait_list*/,
+                                           cl_event * /*event*/) CL_EXT_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int
+(CL_API_CALL * clEnqueueWriteSignalAMD_fn)( cl_command_queue /*command_queue*/,
+                                            cl_mem /*mem_object*/,
+                                            cl_uint /*value*/,
+                                            cl_ulong /*offset*/,
+                                            cl_uint /*num_events*/,
+                                            const cl_event * /*event_list*/,
+                                            cl_event * /*event*/) CL_EXT_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int
+(CL_API_CALL * clEnqueueMakeBuffersResidentAMD_fn)( cl_command_queue /*command_queue*/,
+                                                    cl_uint /*num_mem_objs*/,
+                                                    cl_mem * /*mem_objects*/,
+                                                    cl_bool /*blocking_make_resident*/,
+                                                    cl_bus_address_amd * /*bus_addresses*/,
+                                                    cl_uint /*num_events*/,
+                                                    const cl_event * /*event_list*/,
+                                                    cl_event * /*event*/) CL_EXT_SUFFIX__VERSION_1_2;
+#endif /* CL_VERSION_1_2 */
+
+#ifdef CL_VERSION_2_0
+/*********************************
+* cl_khr_sub_groups extension
+*********************************/
+#define cl_khr_sub_groups 1
+
+typedef cl_uint  cl_kernel_sub_group_info;
+
+/* cl_khr_sub_group_info */
+#define CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE_KHR	0x2033
+#define CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR		0x2034
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetKernelSubGroupInfoKHR(cl_kernel /* in_kernel */,
+						   cl_device_id /*in_device*/,
+						   cl_kernel_sub_group_info /* param_name */,
+						   size_t /*input_value_size*/,
+						   const void * /*input_value*/,
+						   size_t /*param_value_size*/,
+						   void* /*param_value*/,
+						   size_t* /*param_value_size_ret*/ ) CL_EXT_SUFFIX__VERSION_2_0;
+						   
+typedef CL_API_ENTRY cl_int
+     ( CL_API_CALL * clGetKernelSubGroupInfoKHR_fn)(cl_kernel /* in_kernel */,
+						      cl_device_id /*in_device*/,
+						      cl_kernel_sub_group_info /* param_name */,
+						      size_t /*input_value_size*/,
+						      const void * /*input_value*/,
+						      size_t /*param_value_size*/,
+						      void* /*param_value*/,
+						      size_t* /*param_value_size_ret*/ ) CL_EXT_SUFFIX__VERSION_2_0;
+#endif /* CL_VERSION_2_0 */
+
+#ifdef __cplusplus
+}
+#endif
+
+
+#endif /* __CL_EXT_H */
@@ -0,0 +1,162 @@
+/**********************************************************************************
+ * Copyright (c) 2008 - 2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+#ifndef __OPENCL_CL_GL_H
+#define __OPENCL_CL_GL_H
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif	
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef cl_uint     cl_gl_object_type;
+typedef cl_uint     cl_gl_texture_info;
+typedef cl_uint     cl_gl_platform_info;
+typedef struct __GLsync *cl_GLsync;
+
+/* cl_gl_object_type = 0x2000 - 0x200F enum values are currently taken           */
+#define CL_GL_OBJECT_BUFFER                     0x2000
+#define CL_GL_OBJECT_TEXTURE2D                  0x2001
+#define CL_GL_OBJECT_TEXTURE3D                  0x2002
+#define CL_GL_OBJECT_RENDERBUFFER               0x2003
+#define CL_GL_OBJECT_TEXTURE2D_ARRAY            0x200E
+#define CL_GL_OBJECT_TEXTURE1D                  0x200F
+#define CL_GL_OBJECT_TEXTURE1D_ARRAY            0x2010
+#define CL_GL_OBJECT_TEXTURE_BUFFER             0x2011
+
+/* cl_gl_texture_info           */
+#define CL_GL_TEXTURE_TARGET                    0x2004
+#define CL_GL_MIPMAP_LEVEL                      0x2005
+#define CL_GL_NUM_SAMPLES                       0x2012
+
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLBuffer(cl_context     /* context */,
+                     cl_mem_flags   /* flags */,
+                     cl_GLuint      /* bufobj */,
+                     int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLTexture(cl_context      /* context */,
+                      cl_mem_flags    /* flags */,
+                      cl_GLenum       /* target */,
+                      cl_GLint        /* miplevel */,
+                      cl_GLuint       /* texture */,
+                      cl_int *        /* errcode_ret */) CL_API_SUFFIX__VERSION_1_2;
+    
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLRenderbuffer(cl_context   /* context */,
+                           cl_mem_flags /* flags */,
+                           cl_GLuint    /* renderbuffer */,
+                           cl_int *     /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLObjectInfo(cl_mem                /* memobj */,
+                  cl_gl_object_type *   /* gl_object_type */,
+                  cl_GLuint *           /* gl_object_name */) CL_API_SUFFIX__VERSION_1_0;
+                  
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLTextureInfo(cl_mem               /* memobj */,
+                   cl_gl_texture_info   /* param_name */,
+                   size_t               /* param_value_size */,
+                   void *               /* param_value */,
+                   size_t *             /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueAcquireGLObjects(cl_command_queue      /* command_queue */,
+                          cl_uint               /* num_objects */,
+                          const cl_mem *        /* mem_objects */,
+                          cl_uint               /* num_events_in_wait_list */,
+                          const cl_event *      /* event_wait_list */,
+                          cl_event *            /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueReleaseGLObjects(cl_command_queue      /* command_queue */,
+                          cl_uint               /* num_objects */,
+                          const cl_mem *        /* mem_objects */,
+                          cl_uint               /* num_events_in_wait_list */,
+                          const cl_event *      /* event_wait_list */,
+                          cl_event *            /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+
+/* Deprecated OpenCL 1.1 APIs */
+extern CL_API_ENTRY CL_EXT_PREFIX__VERSION_1_1_DEPRECATED cl_mem CL_API_CALL
+clCreateFromGLTexture2D(cl_context      /* context */,
+                        cl_mem_flags    /* flags */,
+                        cl_GLenum       /* target */,
+                        cl_GLint        /* miplevel */,
+                        cl_GLuint       /* texture */,
+                        cl_int *        /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1_DEPRECATED;
+    
+extern CL_API_ENTRY CL_EXT_PREFIX__VERSION_1_1_DEPRECATED cl_mem CL_API_CALL
+clCreateFromGLTexture3D(cl_context      /* context */,
+                        cl_mem_flags    /* flags */,
+                        cl_GLenum       /* target */,
+                        cl_GLint        /* miplevel */,
+                        cl_GLuint       /* texture */,
+                        cl_int *        /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1_DEPRECATED;
+    
+/* cl_khr_gl_sharing extension  */
+    
+#define cl_khr_gl_sharing 1
+    
+typedef cl_uint     cl_gl_context_info;
+    
+/* Additional Error Codes  */
+#define CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR  -1000
+    
+/* cl_gl_context_info  */
+#define CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR    0x2006
+#define CL_DEVICES_FOR_GL_CONTEXT_KHR           0x2007
+    
+/* Additional cl_context_properties  */
+#define CL_GL_CONTEXT_KHR                       0x2008
+#define CL_EGL_DISPLAY_KHR                      0x2009
+#define CL_GLX_DISPLAY_KHR                      0x200A
+#define CL_WGL_HDC_KHR                          0x200B
+#define CL_CGL_SHAREGROUP_KHR                   0x200C
+    
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLContextInfoKHR(const cl_context_properties * /* properties */,
+                      cl_gl_context_info            /* param_name */,
+                      size_t                        /* param_value_size */,
+                      void *                        /* param_value */,
+                      size_t *                      /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+    
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetGLContextInfoKHR_fn)(
+    const cl_context_properties * properties,
+    cl_gl_context_info            param_name,
+    size_t                        param_value_size,
+    void *                        param_value,
+    size_t *                      param_value_size_ret);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_GL_H */
@@ -0,0 +1,69 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+/* cl_gl_ext.h contains vendor (non-KHR) OpenCL extensions which have           */
+/* OpenGL dependencies.                                                         */
+
+#ifndef __OPENCL_CL_GL_EXT_H
+#define __OPENCL_CL_GL_EXT_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifdef __APPLE__
+    #include <OpenCL/cl_gl.h>
+#else
+    #include <CL/cl_gl.h>
+#endif
+
+/*
+ * For each extension, follow this template
+ *  cl_VEN_extname extension  */
+/* #define cl_VEN_extname 1
+ * ... define new types, if any
+ * ... define new tokens, if any
+ * ... define new APIs, if any
+ *
+ *  If you need GLtypes here, mirror them with a cl_GLtype, rather than including a GL header
+ *  This allows us to avoid having to decide whether to include GL headers or GLES here.
+ */
+
+/* 
+ *  cl_khr_gl_event  extension
+ *  See section 9.9 in the OpenCL 1.1 spec for more information
+ */
+#define CL_COMMAND_GL_FENCE_SYNC_OBJECT_KHR     0x200D
+
+extern CL_API_ENTRY cl_event CL_API_CALL
+clCreateEventFromGLsyncKHR(cl_context           /* context */,
+                           cl_GLsync            /* cl_GLsync */,
+                           cl_int *             /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif	/* __OPENCL_CL_GL_EXT_H  */
@@ -0,0 +1,54 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+#ifndef __OPENCL_H
+#define __OPENCL_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifdef __APPLE__
+
+#include <OpenCL/cl.h>
+#include <OpenCL/cl_gl.h>
+#include <OpenCL/cl_gl_ext.h>
+#include <OpenCL/cl_ext.h>
+
+#else
+
+#include <CL/cl.h>
+#include <CL/cl_gl.h>
+#include <CL/cl_gl_ext.h>
+#include <CL/cl_ext.h>
+
+#endif
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_H   */
+
@@ -0,0 +1,145 @@
+//#include "kernels.h"
+//
+//const char* vertex_shader_kernel_source =
+//"__kernel                                                                                                                    \n"
+//"void vertex_shader( __global float* m,                                                                                      \n"
+//"                    __global float* VertexBuffer,                                                                           \n"
+//"                    __global float* NewVertexBuffer)                                                                        \n"
+//"{                                                                                                                           \n"
+//"    int local_index = get_local_id(0);                                                                                      \n"
+//"    int global_index = get_group_id(0);                                                                                     \n"
+//"    NewVertexBuffer[4*global_index+local_index] =                                                                           \n"
+//"						  m[local_index*4]*VertexBuffer[3*global_index]                                                      \n"
+//"	                    + m[local_index*4 + 1]*VertexBuffer[3*global_index+1]                                                \n"
+//"	                    + m[local_index*4 + 2]*VertexBuffer[3*global_index+2]                                                \n"
+//"	                    + m[local_index*4 + 3];                                                                              \n"
+//"}                                                                                                                           \n";
+//
+//const char* fragment_shader_kernel_source =
+//"float3 barycentric(float3* pts, float3 P)                                                                                                                                                                    \n"
+//"{                                                                                                                                                                                                          \n"
+//"    float3 u = cross(                                                                                                                                                                                      \n"
+//"        (float3){pts[0][2] - pts[0][0], pts[0][1] - pts[0][0], pts[0][0] - P[0]}, // AC_x, AB_x, distance_x                                                                                                \n"
+//"        (float3){pts[1][2] - pts[1][0], pts[1][1] - pts[1][0], pts[1][0] - P[1]}  // AC_y, AB_y, distance_y                                                                                                \n"
+//"    );                                                                                                                                                                                                     \n"
+//"	 if (fabs(u[2]) < 1) return (float3){-1, 1, 1};                                                                                                                                                         \n"
+//"	 return (float3){1.f - (u[0] + u[1]) / u[2], u[1] / u[2], u[0] / u[2]};                                                                                                                                 \n"
+//"}                                                                                                                                                                                                          \n"
+//"                                                                                                                                                                                                           \n"
+//"__kernel void fragment_shader (                                                                                                                                                                            \n"
+//"    __global int3* faces,                                                                                                                                                                                   \n"
+//"    __global float* vertices,                                                                                                                                                                             \n"
+//"    __global int* pixels,                                                                                                                                                                                  \n"
+//"    __global int* screen_width,                                                                                                                                                                            \n"
+//"    __global float* z_buffer,                                                                                                                                                                             \n"
+//"    __global int* nfaces,                                                                                                                                                                                 \n"
+//"    __global float* uv_buffer,                                                                                                                                                                           \n"
+//"    __global int* map_size,                                                                                                                                                                              \n"
+//"    __global float* light_dir,                                                                                                                                                                            \n"
+//"    __global float* norms_buff,                                                                                                                                                                           \n"
+//"    __global uchar* diffuse_map                                                                                                                                                                              \n"
+//") {                                                                                                                                                                                                        \n"
+//"    int GROUP_ID = get_group_id(0);                                                                                                                                                                          \n"
+//"    int GROUP_SIZE = get_local_size(0);                                                                                                                                                                          \n"
+//"    int LOCAL_ID = get_local_id(0);                                                                                                                                                                          \n"
+//"                                                                                                                                                                                                           \n"
+//"    bool out = true;                                                                                                                                                                                       \n"
+//"    float3 vertices3[3];                                                                                                                                                                                   \n"
+//"    float2 uv_coords[3];                                 \n"
+//"    float3 norms[3];                                                                                                                                                                                   \n"
+//"                                                                                                                                                                                                           \n"
+//"    for(int i = 0; i < 3; i++) {                                                                                                                                                                           \n"
+//"        float4 vertex;// =  vertices[faces[GROUP_ID * 3 + i ][0]];                                                                                                                                                \n"
+//"        for(int j = 0; j < 4; j ++) {                                                                                                                                                                                                   \n"
+//"            vertex[j] = vertices[4 * faces[GROUP_ID*3 + i][0] + j];                                                                                                                                          \n"
+//"        }                                                                                                                                                                                                  \n"
+//"                                                                                                                                                                                                           \n"
+//"        for(int j = 0; j < 2; j++) {                                                                                                                                                                       \n"
+//"            uv_coords[i][j]  = uv_buffer [2 * faces[GROUP_ID * 3 + i][1] + j];                                                                                                                                                \n"
+//"        }                                                                                                                                                                                                  \n"
+//"                                                                                                                                                                                                           \n"
+//"        for( int j = 0; j < 3; j++ ) {																		\n"
+//"            vertices3[j][i] = (vertex[j]/vertex[3]);                                                                                                                                                       \n"
+//"		     norms[j][i] = norms_buff[3 * (faces[GROUP_ID * 3 + i][2]) + j];                                                                                                                                          \n"
+//"        }                                                                                                                                                                                                  \n"
+//"                                                                                                                                                                                                           \n"
+//"        if (   vertices3[0][i] > 0 && vertices3[0][i] < *screen_width                                                                                                                                      \n"
+//"            && vertices3[1][i] > 0 && vertices3[1][i] < *screen_width   )                                                                                                                                  \n"
+//"                out = false;                                                                                                                                                                               \n"
+//"    }                                                                                                                                                                                                      \n"
+//"                                                                                                                                                                                                           \n"
+//"    if(out) return;                                                                                                                                                                                        \n"
+//"                                                                                                                                                                                                           \n"
+//"    //if(vertices3[1][0] == vertices3[1][1] && vertices3[1][2] == vertices3[1][1]) return;                                                                                                                   \n"
+//"                                                                                                                                                                                                           \n"
+//"    int2 bounding_box_min = (int2) { *screen_width - 1, *screen_width - 1 };                                                                                                                                    \n"
+//"    int2 bounding_box_max = (int2) { 0, 0 };                                                                                                                                                                    \n"
+//"    int2 clamper = (int2) { *screen_width - 1, *screen_width - 1 };                                                                                                                                               \n"
+//"                                                                                                                                                                                                           \n"
+//"    for(int i = 0; i < 3; i++) {                                                                                                                                                                           \n"
+//"        for(int j = 0; j < 2; j++) {                                                                                                                                                                       \n"
+//"            bounding_box_min[j] = max(0, min(bounding_box_min[j], (int)vertices3[j][i]));                                                                                                                  \n"
+//"            bounding_box_max[j] = min(clamper[j], max(bounding_box_max[j], (int)vertices3[j][i]));                                                                                                           \n"
+//"        }                                                                                                                                                                                                  \n"
+//"    }                                                                                                                                                                                                      \n"
+//"                                                                                                                                                                                                           \n"
+//"    if(bounding_box_min[0] > *screen_width || bounding_box_max[0] < 0 || bounding_box_min[1] > *screen_width || bounding_box_max[1] < 0) return;                                                           \n"
+//"                                                                                                                                                                                                           \n"
+//"                                                                                                                                                                                                           \n"
+//"    int X_PER_ITEM = (int)(ceil((float)(bounding_box_max[0] - bounding_box_min[0]) / 16.f));                                                                                                                                                                                            \n"
+//"    int Y_PER_ITEM = (int)(ceil((float)(bounding_box_max[1] - bounding_box_min[1]) / 16.f));                                                                                                                                                                                            \n"
+//"    float STARTING_X = bounding_box_min[0] + X_PER_ITEM * (LOCAL_ID % 16);                                                                                                                                                                                            \n"
+//"    float ENDING_X = STARTING_X + X_PER_ITEM;                                                                                                                                                                                            \n"
+//"    float STARTING_Y = bounding_box_min[1] + Y_PER_ITEM * (LOCAL_ID / 16);                                                                                                                                                                                            \n"
+//"    float ENDING_Y = STARTING_Y + Y_PER_ITEM;                                                                                                                                                                                            \n"
+//"                                                                                                                                                                                                           \n"
+//"                                                                                                                                                                                                           \n"
+//"    float3 point;                                                                                                                                                                                            \n"
+//"    for(point[0] = STARTING_X; point[0] <= ENDING_X; point[0]++) {                                                                                                                     \n"
+//"        for(point[1] = STARTING_Y; point[1] <= ENDING_Y; point[1]++) {                                                                                                                 \n"
+//"           if(point[1] >= *screen_width || point[1] >= *screen_width) break;                                                                                                                                                                                                \n"
+//"			float3 bc_coord = barycentric(vertices3, point);                                                                                                                                                \n"
+//"           if (bc_coord[0] < 0 || bc_coord[1] < 0 || bc_coord[2] < 0) continue;                                                                                                                                                                                               \n"
+//"                                                                                                                                                                                                           \n"
+//"           float2 uv_vec = (float2){0, 0};                                                                                                                                                                                  \n"
+//"           float3 normal = (float3){0, 0, 0};                                                                                                                                                                \n"
+//"                                                                                                                                                                                                           \n"
+//"           point[2] = dot(vertices3[2], bc_coord);                                                                                                                                                                                   \n"
+//"			if (z_buffer[(int)(point[0] + point[1] * *screen_width)] > point[2]) {                                                                        \n"
+//"               continue;                                                                                                                                                                                   \n"
+//"           }                                                                                                                                                                                               \n"
+//"                                                                                                                                                                                                           \n"
+//"                                                                                                                                                                                                           \n"
+//"			for (int i = 0; i < 3; i++) {                                                                                                                                                                   \n"
+//"               uv_vec[0]  += uv_coords[i][0] * bc_coord[i];                                                                                                                                                 \n"
+//"               uv_vec[1]  += uv_coords[i][1] * bc_coord[i];                                                                                                                                                 \n"
+//"               normal[i]  = dot(norms[i], bc_coord);                                                                                                                                                 \n"
+//"           }                                                                                                                                                                                               \n"
+//"           int2 uv_point = (int2) { (int)(uv_vec[0] * map_size[0]), (int)(uv_vec[1] * map_size[1]) };                                                                                                                    \n"
+//"                                                                                                                                                                                                           \n"
+//"                                                                                                                                                                                                           \n"
+//"           int col_index = 3 * (uv_point[0] + uv_point[1] * map_size[0]);                                                                                                                                       \n"
+//"                                                                                                                                                                                                           \n"
+//"                                                                                                                                                                                                           \n"
+//"                                                                                                                                                                                                           \n"
+//"                                                                                                                                                                                                           \n"
+//"           float3 normalized_norm = normalize(normal);                                                                                                                                                                                                \n"
+//"                                                                                                                                                                                                           \n"
+//"                                                                                                                                                                                                           \n"
+//"           float intensity = clamp((dot(normalized_norm , (float3){light_dir[0], light_dir[1], light_dir[2]})), 0.f, 1.f) + 0.2;                                                                                                                                                                                                \n"
+//"                                                                                                                                                                                                           \n"
+//"           int color = 0;                                                                                                                                                                                  \n"
+//"           color |= ((int)fmin((float)(diffuse_map[col_index + 0]) * intensity, (float) 0xff)) << 16;                                                                                                                                                               \n"
+//"           color |= ((int)fmin((float)(diffuse_map[col_index + 1]) * intensity, (float) 0xff)) <<  8;                                                                                                                                                               \n"
+//"           color |= ((int)fmin((float)(diffuse_map[col_index + 2]) * intensity, (float) 0xff)) <<  0;                                                                                                                                                               \n"
+//"                                                                                                                                                                                                          \n"
+//"           //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 16;                                                                                                                                                               \n"
+//"           //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) <<  8;                                                                                                                                                               \n"
+//"           //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) <<  0;                                                                                                                                                               \n"
+//"                                                                                                                                                                                                          \n"
+//"                                                                                                                                                                                                          \n"
+//"           z_buffer[(int) (point[0] + point[1] * *screen_width)] = point[2];                                                                                                                                 \n"
+//"			pixels  [(int) (point[0] + point[1] * *screen_width)] = color;// & 0x00ffffff;                                                                                                                                     \n"
+//"        }                                                                                                                                                                                                  \n"
+//"    }                                                                                                                                                                                                      \n"
+//"}                                                                                                                                                                                                          \n";
@@ -0,0 +1,78 @@
+#include "kernels.h"
+#include "renderer.h"
+#include <debugapi.h>
+#include <ctime>
+#include <climits>
+
+cl_uint numPlatforms;
+cl_uint numDevices;
+
+cl_platform_id platform_id;
+cl_device_id device_id;
+cl_context context;
+
+int* hidden_pixel_buffer;
+cl_mem z_buffer_mem;
+cl_mem pixel_data_buffer;
+cl_mem time_buffer;
+//==================
+
+cl_platform_id* platforms = NULL;
+cl_device_id* devices = NULL;
+
+cl_command_queue commands;
+
+int err;
+
+
+void init_kernels() {
+	err = clGetPlatformIDs(0, NULL, &numPlatforms);
+	platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * numPlatforms);
+	err = clGetPlatformIDs(numPlatforms, platforms, NULL);
+
+	err = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices);
+	devices = (cl_device_id*)malloc(sizeof(cl_device_id)*numDevices);
+	err = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
+
+	context = clCreateContext(NULL, numDevices, devices, NULL, NULL, &err);
+
+	commands = clCreateCommandQueue(context, devices[0], 0, &err);
+
+	hidden_pixel_buffer = new int[screen_width*screen_height];
+	pixel_data_buffer  = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * screen_height * screen_width  , NULL, &err);
+	z_buffer_mem       = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * screen_height * screen_width, NULL, &err);
+	time_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float), NULL, &err);
+}
+
+
+void clear(cl_mem* buffer, size_t size, const int pattern) {
+	clEnqueueFillBuffer(commands, *buffer, &pattern, sizeof(int), 0, size, 0, NULL, NULL);
+}
+
+void new_frame() {
+	clear(&pixel_data_buffer, sizeof(int) * screen_width * screen_height, 0);
+	clear(&z_buffer_mem, sizeof(float) * screen_width * screen_height, 0);
+	clEnqueueWriteBuffer(commands, time_buffer, CL_TRUE, 0, sizeof(float), &TIME, 0, NULL, NULL);
+	clFinish(commands);
+}
+
+void end_frame() {
+	clFinish(commands);
+	clEnqueueReadBuffer(commands, pixel_data_buffer, CL_TRUE, 0, sizeof(int) * screen_height * screen_width, hidden_pixel_buffer, 0, NULL, NULL);
+
+	for (int j = 0; j < screen_height; j++)
+		for (int i = 0; i < screen_width; i++)
+			set_pixel(i, j, hidden_pixel_buffer[j * screen_width + i]);
+}
+
+void destroy_kernels()
+{
+	clReleaseMemObject(z_buffer_mem);
+	clReleaseMemObject(pixel_data_buffer);
+	clReleaseMemObject(time_buffer);
+	clReleaseCommandQueue(commands);
+
+	free(hidden_pixel_buffer);
+
+	clReleaseContext(context);
+}
@@ -0,0 +1,62 @@
+#ifndef _KERNEL_HEADER_
+#define _KERNEL_HEADER_
+#include <Windows.h>
+#include <CL/cl.h>
+#include "geometry.h"
+#include "util_window.h"
+
+
+extern cl_uint numPlatforms;
+extern cl_platform_id firstPlatformId;
+extern cl_device_id device_id;
+extern cl_context context;
+
+
+
+// Matrix Multiplication Variables
+#define ORDER 4
+#define MATRIX_SIZE 16
+extern int Mdim, Pdim, Ndim;
+extern cl_mem mat_a;
+extern cl_mem mat_b;
+extern cl_mem mat_c;
+
+extern cl_program mat_mul_prog;
+extern cl_kernel mat_mul_kernel;
+extern cl_command_queue commands;
+
+extern int err;
+
+extern const char* matrix_mul_kernel_source;
+// End of: Matrix Multiplication Variables
+
+extern const char* vertex_shader_kernel_source;
+extern const char* fragment_shader_kernel_source;
+extern const char* fluid_vertex_shader_kernel_source;
+
+extern cl_command_queue commands;
+extern cl_platform_id* platforms;
+extern cl_device_id* devices;
+extern cl_context context;
+
+extern cl_mem z_buffer_mem;
+extern cl_mem pixel_data_buffer;
+extern cl_mem time_buffer;
+
+void init_kernels();
+void destroy_kernels();
+void new_frame();
+void end_frame();
+
+
+
+
+
+
+
+
+
+
+
+
+#endif // !_KERNEL_HEADER_
@@ -2,28 +2,43 @@
 #include "util_window.h"
 #include "renderer.h"
 #include "ctime"
+#include "kernels.h"

-const int width = 800;
-const int height = 800;
+const int screen_width = 1000;
+const int screen_height = 1000;
+
+#define TARGET_FRAMERATE 30
+
+int prev_mouse_x = screen_width/2;
+int prev_mouse_y = screen_height/2;
+
+float TIME = 0;

 int WINAPI WinMain(HINSTANCE hInstance, HINSTANCE hPrevInstance, LPSTR lpCmdLine, int nCmdShow)
 {
   HWND                hwnd;
   MSG                 Msg;

-   hwnd = create_window(width, height, hInstance);
+   hwnd = create_window(hInstance);
+   ShowCursor(false);
+
+    AllocConsole();
+	freopen("CONIN$", "r",stdin);
+	freopen("CONOUT$", "w",stdout);
+	freopen("CONOUT$", "w",stderr);
+
   ShowWindow(hwnd, nCmdShow);

-   char debug_str[100];
+   init_camera();
+   init_kernels();

-   clock_t first_time = clock();
-   render();
-   clock_t end_time = clock();
-   Update();

-   sprintf_s(debug_str, "%f\n", (double)(end_time - first_time) / CLOCKS_PER_SEC);
-   OutputDebugString(debug_str);
-   
+	//camera.ApplyChanges();
+	//render();
+	//Update();
+
+   SetTimer(hwnd, NULL, 1000 / TARGET_FRAMERATE, (TIMERPROC)FixedUpdate);
+
   while (GetMessage(&Msg, NULL, 0, 0))
   {
      TranslateMessage(&Msg);
@@ -32,17 +47,59 @@ int WINAPI WinMain(HINSTANCE hInstance, HINSTANCE hPrevInstance, LPSTR lpCmdLine
   return Msg.wParam;
 }

+void HandleMouseMovement() {
+		POINT point;
+		GetCursorPos(&point);
+
+		camera.rotate_hor(point.x - prev_mouse_x);
+		camera.rotate_ver(point.y - prev_mouse_y);
+
+		SetCursorPos(prev_mouse_x, prev_mouse_y);
+}
+
+bool HandleButtonPressed() {
+	if (GetAsyncKeyState(0x57) & 0x8000)
+		camera.move_camera_forward();
+	if (GetAsyncKeyState(0x53) & 0x8000)
+		camera.move_camera_backward();
+	if (GetAsyncKeyState(0x44) & 0x8000)
+		camera.move_camera_right();
+	if (GetAsyncKeyState(0x41) & 0x8000)
+		camera.move_camera_left();
+	if (GetAsyncKeyState(VK_SPACE) & 0x8000)
+		camera.rise();
+	if (GetAsyncKeyState(VK_CONTROL) & 0x8000)
+		camera.fall();
+	return true;
+}
+
+void CALLBACK FixedUpdate(HWND hwnd, UINT message, UINT uInt, DWORD dWord)
+{
+	TIME += 1.f / (float)TARGET_FRAMERATE;
+	HandleButtonPressed();
+	camera.ApplyChanges();
+	clock_t start = clock();
+	render();
+	Update();
+	clock_t end = clock();
+	printf("FPS: %f\n", 1/((float)(end-start)/CLOCKS_PER_SEC));
+}
+
 LRESULT CALLBACK WndProc(HWND hwnd, UINT message, WPARAM wParam, LPARAM lParam)
 {
 	switch (message)
 	{
+	case WM_MOUSEMOVE:
+		HandleMouseMovement();
+		break;
 	case WM_RBUTTONDOWN:
 		break;
 	case WM_LBUTTONDOWN:
-		Update();
 		break;
 	case WM_CLOSE:
 		DestroyWindow(hwnd);
+		free_renderer();
+		destroy_kernels();
 		break;
 	case WM_DESTROY:
 		PostQuitMessage(0);
@@ -6,7 +6,8 @@
 #define PI 3.14159265358979323846
 #define DEG2RAD PI/180

-Model::Model(const char *filename) : verts_(), faces_(), norms_(), uv_(), diffusemap_(), normalmap_(), specularmap_() {
+Model::Model(const char *filename, int watery) : verts_(), faces_(), norms_(), uv_(), diffusemap_(), normalmap_(), specularmap_() {
+	fluid = watery;
    std::ifstream in;
 	Transform = Matrix::identity();
 	Rotation = Matrix::identity();
@@ -45,22 +46,22 @@ Model::Model(const char *filename) : verts_(), faces_(), norms_(), uv_(), diffus
            faces_.push_back(f);
        }
    }
-    std::cerr << "# v# " << verts_.size() << " f# "  << faces_.size() << " vt# " << uv_.size() << " vn# " << norms_.size() << std::endl;
    load_texture(filename, "_diffuse.tga", diffusemap_);
-    load_texture(filename, "_nm_tangent.tga",      normalmap_);
-    //load_texture(filename, "_spec.tga",    specularmap_);
+	init_kernels();
 }

-Model::~Model() {}
+Model::~Model() {
+	release_kernels();
+}

 void Model::ApplyTransform() {
 	Transform = Translation * Scale * Rotation;
 }

 void Model::translate(Vec3f tr) {
-	Translation[0][3] += tr.x;
-	Translation[1][3] += tr.y;
-	Translation[2][3] += tr.z;
+	Translation[0][3] = tr.x;
+	Translation[1][3] = tr.y;
+	Translation[2][3] = tr.z;
 }
 void Model::rotate(Vec3f rot) {
 	rot = rot * DEG2RAD;
@@ -109,7 +110,7 @@ void Model::load_texture(std::string filename, const char *suffix, TGAImage &img
    if (dot!=std::string::npos) {
        texfile = texfile.substr(0,dot) + std::string(suffix);
        std::cerr << "texture file " << texfile << " loading " << (img.read_tga_file(texfile.c_str()) ? "ok" : "failed") << std::endl;
-        img.flip_vertically();
+        //img.flip_vertically();
    }
 }

@@ -141,3 +142,275 @@ Vec3f Model::normal(int iface, int nthvert) {
    return norms_[idx].normalize();
 }

+void Model::init_kernels() {
+	if(fluid)
+		vertex_shader_prog = clCreateProgramWithSource(context, 1, (const char **)&fluid_vertex_shader_kernel_source, NULL, &err);
+	else
+		vertex_shader_prog = clCreateProgramWithSource(context, 1, (const char **)&vertex_shader_kernel_source, NULL, &err);
+	err = clBuildProgram(vertex_shader_prog, 1, devices, NULL, NULL, NULL);
+	vertex_shader_kernel = clCreateKernel(vertex_shader_prog, "vertex_shader", &err);
+	vertex_shader_matz = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * 16, NULL, &err);
+	vertex_shader_vertices = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * nverts() * 3, NULL, &err);
+	new_vertices_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float4) * nverts(), NULL, &err);
+
+	clSetKernelArg(vertex_shader_kernel, 0, sizeof(cl_mem), &vertex_shader_matz);
+	clSetKernelArg(vertex_shader_kernel, 1, sizeof(cl_mem), &vertex_shader_vertices);
+	clSetKernelArg(vertex_shader_kernel, 2, sizeof(cl_mem), &new_vertices_mem);
+	clSetKernelArg(vertex_shader_kernel, 3, sizeof(cl_mem), &time_buffer);
+	clEnqueueWriteBuffer(commands, vertex_shader_vertices, CL_TRUE, 0, sizeof(float) * nverts() * 3, *(float**)((Vec3f*) &verts_), 0, NULL, NULL);
+
+	int map_size[] = { diffusemap_.get_width(), diffusemap_.get_height() };
+	faces = (cl_int3*)malloc(3 * sizeof(cl_int3) * nfaces());
+	for (int i = 0; i < nfaces(); i++) {
+		for (int j = 0; j < 3; j++) {
+			faces[i * 3 + j].x = faces_[i][j][0];
+			faces[i * 3 + j].y = faces_[i][j][1];
+			faces[i * 3 + j].z = faces_[i][j][2];
+		}
+	}
+
+	fragment_shader_prog = clCreateProgramWithSource(context, 1, (const char**)&fragment_shader_kernel_source, NULL, &err);
+	err = clBuildProgram(fragment_shader_prog, 1, devices, NULL, NULL, NULL);
+	fragment_shader_kernel = clCreateKernel(fragment_shader_prog, "fragment_shader", &err);
+
+	fragment_shader_faces        = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(cl_int3) * 3 * nfaces()               , NULL, &err);
+	fragment_shader_screen_width = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int)                                  , NULL, &err);
+	fragment_shader_uv           = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * uv_.size() * 2               , NULL, &err);
+	fragment_shader_map_size     = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int) * 2                              , NULL, &err);
+	fragment_shader_norms        = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * norms_.size() * 3            , NULL, &err);
+	fragment_shader_light_dir    = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * 3                            , NULL, &err);
+	fragment_shader_diffuse_map  = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(char) * 3 * map_size[0] * map_size[1] , NULL, &err);
+
+	clSetKernelArg(fragment_shader_kernel, 0, sizeof(cl_mem), &fragment_shader_faces);
+	clSetKernelArg(fragment_shader_kernel, 1, sizeof(cl_mem), &new_vertices_mem);
+	clSetKernelArg(fragment_shader_kernel, 2, sizeof(cl_mem), &pixel_data_buffer);
+	clSetKernelArg(fragment_shader_kernel, 3, sizeof(cl_mem), &fragment_shader_screen_width);
+	clSetKernelArg(fragment_shader_kernel, 4, sizeof(cl_mem), &z_buffer_mem);
+	clSetKernelArg(fragment_shader_kernel, 5, sizeof(cl_mem), &fragment_shader_uv);
+	clSetKernelArg(fragment_shader_kernel, 6, sizeof(cl_mem), &fragment_shader_map_size);
+	clSetKernelArg(fragment_shader_kernel, 7, sizeof(cl_mem), &fragment_shader_light_dir);
+	clSetKernelArg(fragment_shader_kernel, 8, sizeof(cl_mem), &fragment_shader_norms);
+	clSetKernelArg(fragment_shader_kernel, 9, sizeof(cl_mem), &fragment_shader_diffuse_map);
+
+	clEnqueueWriteBuffer(commands, fragment_shader_screen_width, CL_FALSE, 0, sizeof(int)                                 , &screen_width              , 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_faces       , CL_FALSE, 0, sizeof(cl_int3) * 3 * nfaces()              , faces                      , 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_uv          , CL_FALSE, 0, sizeof(float) * 2 * uv_.size()              , *(float**)((Vec2f*)&uv_)   , 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_map_size    , CL_FALSE, 0, sizeof(int) * 2                             , map_size                   , 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_norms       , CL_FALSE, 0, sizeof(float) * norms_.size() * 3           , *(float**)((Vec3f*)&norms_), 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_diffuse_map , CL_TRUE, 0, sizeof(char) * 3 * map_size[0] * map_size[1], diffusemap_.data           , 0, NULL, NULL);
+}
+
+void Model::vertex(float* z) {
+	size_t vertex_shader_global[] = { nverts() * 4 };
+	size_t vertex_shader_local[] = { 4 };
+
+	clEnqueueWriteBuffer(commands, vertex_shader_matz, CL_TRUE, 0, sizeof(float) * 16, z, 0, NULL, NULL);
+
+	clEnqueueNDRangeKernel(commands, vertex_shader_kernel, 1, NULL, vertex_shader_global, vertex_shader_local, 0, NULL, NULL);
+}
+void Model::fragment(float* light_dir) {
+	clEnqueueWriteBuffer(commands, fragment_shader_light_dir, CL_FALSE, 0, sizeof(float) * 3, light_dir, 0, NULL, NULL);
+	size_t fragment_shader_global[] = { nfaces() * 256 };
+	size_t fragment_shader_local[] = { 256 };
+
+	clEnqueueNDRangeKernel(commands, fragment_shader_kernel, 1, NULL, fragment_shader_global, fragment_shader_local, 0, NULL, NULL);
+
+}
+void Model::render(Matrix* z, float* light_dir) {
+	Matrix m = *z * Transform;
+	vertex((float*)&m);
+	clFinish(commands);
+	fragment(light_dir);
+}
+void Model::release_kernels() {
+
+	clReleaseProgram(vertex_shader_prog);
+	clReleaseKernel(vertex_shader_kernel);
+	clReleaseMemObject(vertex_shader_matz);
+	clReleaseMemObject(vertex_shader_vertices);
+	clReleaseMemObject(new_vertices_mem);
+
+	clReleaseProgram(fragment_shader_prog);
+	clReleaseKernel(fragment_shader_kernel);
+	clReleaseMemObject(fragment_shader_faces);
+	clReleaseMemObject(fragment_shader_screen_width);
+	clReleaseMemObject(fragment_shader_uv);
+	clReleaseMemObject(fragment_shader_map_size);
+	clReleaseMemObject(fragment_shader_norms);
+	clReleaseMemObject(fragment_shader_light_dir);
+	clReleaseMemObject(fragment_shader_diffuse_map);
+
+	free(faces);
+}
+
+
+const char* fluid_vertex_shader_kernel_source =
+"__kernel                                                                                                                    \n"
+"void vertex_shader( __global float* m,                                                                                      \n"
+"                    __global float* VertexBuffer,                                                                           \n"
+"                    __global float* NewVertexBuffer,                                                                        \n"
+"                    __global float* TIME           )                                                                        \n"
+"{                                                                                                                           \n"
+"    int local_index = get_local_id(0);                                                                                      \n"
+"    int global_index = get_group_id(0);                                                                                     \n"
+"    float AMPLITUDE = 5;                                                                                     \n"
+"    float FREQ = 2.0f;                                                                                     \n"
+"                                                                                                                            \n"
+"    float offset = 0.f;                                                                                                     \n"
+"    if(local_index == 1) {                                                                                                  \n"
+"    offset = (sin(TIME[0]*(fmod(10-VertexBuffer[3*global_index],5.f) * FREQ)) + sin(TIME[0]*(fmod(10-VertexBuffer[3*global_index+2], 5.f))*FREQ)) * AMPLITUDE;                                                                                                \n"
+"    }                                                                                                                       \n"
+"                                                                                                                            \n"
+"    NewVertexBuffer[4*global_index+local_index] =                                                                           \n"
+"						  m[local_index*4]*VertexBuffer[3*global_index]                                                      \n"
+"	                    + m[local_index*4 + 1]*VertexBuffer[3*global_index+1]                                                \n"
+"	                    + m[local_index*4 + 2]*VertexBuffer[3*global_index+2]                                                \n"
+"	                    + m[local_index*4 + 3]                                                                               \n"
+"	                    + offset;                                                                                            \n"
+"}                                                                                                                           \n";
+
+
+const char* vertex_shader_kernel_source =
+"__kernel                                                                                                                    \n"
+"void vertex_shader( __global float* m,                                                                                      \n"
+"                    __global float* VertexBuffer,                                                                           \n"
+"                    __global float* NewVertexBuffer,                                                                        \n"
+"                    __global float* TIME           )                                                                        \n"
+"{                                                                                                                           \n"
+"    int local_index = get_local_id(0);                                                                                      \n"
+"    int global_index = get_group_id(0);                                                                                     \n"
+"                                                                                                                            \n"
+"    NewVertexBuffer[4*global_index+local_index] =                                                                           \n"
+"						  m[local_index*4]*VertexBuffer[3*global_index]                                                      \n"
+"	                    + m[local_index*4 + 1]*VertexBuffer[3*global_index+1]                                                \n"
+"	                    + m[local_index*4 + 2]*VertexBuffer[3*global_index+2]                                                \n"
+"	                    + m[local_index*4 + 3];                                                                              \n"
+"}                                                                                                                           \n";
+
+const char* fragment_shader_kernel_source =
+"float3 barycentric(float3* pts, float3 P)                                                                                                                                                                    \n"
+"{                                                                                                                                                                                                          \n"
+"    float3 u = cross(                                                                                                                                                                                      \n"
+"        (float3){pts[0][2] - pts[0][0], pts[0][1] - pts[0][0], pts[0][0] - P[0]}, // AC_x, AB_x, distance_x                                                                                                \n"
+"        (float3){pts[1][2] - pts[1][0], pts[1][1] - pts[1][0], pts[1][0] - P[1]}  // AC_y, AB_y, distance_y                                                                                                \n"
+"    );                                                                                                                                                                                                     \n"
+"	 if (fabs(u[2]) < 1) return (float3){-1, 1, 1};                                                                                                                                                         \n"
+"	 return (float3){1.f - (u[0] + u[1]) / u[2], u[1] / u[2], u[0] / u[2]};                                                                                                                                 \n"
+"}                                                                                                                                                                                                          \n"
+"                                                                                                                                                                                                    \n"
+"__kernel void fragment_shader (                                                                                                                                                                \n"
+"    __global int3* faces,                                                                                                                                                                        \n"
+"    __global float* vertices,                                                                                                                                                                             \n"
+"    __global int* pixels,                                                                                                                                                                  \n"
+"    __global int* screen_width,                                                                                                                                                                  \n"
+"    __global float* z_buffer,                                                                                                                                                                         \n"
+"    __global float* uv_buffer,                                                                                                                                                                        \n"
+"    __global int* map_size,                                                                                                                                                                       \n"
+"    __global float* light_dir,                                                                                                                                                                           \n"
+"    __global float* norms_buff,                                                                                                                                                                           \n"
+"    __global uchar* diffuse_map                                                                                                                                                                              \n"
+") {                                                                                                                                                                                                        \n"
+"    int GROUP_ID = get_group_id(0);                                                                                                                                                                          \n"
+"    int GROUP_SIZE = get_local_size(0);                                                                                                                                                                          \n"
+"    int LOCAL_ID = get_local_id(0);                                                                                                                                                                          \n"
+"                                                                                                                                                                                                           \n"
+"    bool out = true;                                                                                                                                                                                       \n"
+"    float3 vertices3[3];                                                                                                                                                                                   \n"
+"    float2 uv_coords[3];                                 \n"
+"    float3 norms[3];                                                                                                                                                                                   \n"
+"                                                                                                                                                                                                           \n"
+"    for(int i = 0; i < 3; i++) {                                                                                                                                                                           \n"
+"        float4 vertex;// =  vertices[faces[GROUP_ID * 3 + i ][0]];                                                                                                                                                \n"
+"        for(int j = 0; j < 4; j ++) {                                                                                                                                                                                                   \n"
+"            vertex[j] = vertices[4 * faces[GROUP_ID*3 + i][0] + j];                                                                                                                                          \n"
+"        }                                                                                                                                                                                                  \n"
+"                                                                                                                                                                                                           \n"
+"        for(int j = 0; j < 2; j++) {                                                                                                                                                                       \n"
+"            uv_coords[i][j]  = uv_buffer [2 * faces[GROUP_ID * 3 + i][1] + j];                                                                                                                                                \n"
+"        }                                                                                                                                                                                                  \n"
+"                                                                                                                                                                                                           \n"
+"        for( int j = 0; j < 3; j++ ) {																		\n"
+"            vertices3[j][i] = (vertex[j]/vertex[3]);                                                                                                                                                       \n"
+"		     norms[j][i] = norms_buff[3 * (faces[GROUP_ID * 3 + i][2]) + j];                                                                                                                                          \n"
+"        }                                                                                                                                                                                                  \n"
+"                                                                                                                                                                                                           \n"
+"        if (   vertices3[0][i] > 0 && vertices3[0][i] < *screen_width                                                                                                                                      \n"
+"            && vertices3[1][i] > 0 && vertices3[1][i] < *screen_width   )                                                                                                                                  \n"
+"                out = false;                                                                                                                                                                               \n"
+"    }                                                                                                                                                                                                      \n"
+"                                                                                                                                                                                                           \n"
+"    if(out) return;                                                                                                                                                                                        \n"
+"                                                                                                                                                                                                           \n"
+"    //if(vertices3[1][0] == vertices3[1][1] && vertices3[1][2] == vertices3[1][1]) return;                                                                                                                   \n"
+"                                                                                                                                                                                                           \n"
+"    int2 bounding_box_min = (int2) { *screen_width - 1, *screen_width - 1 };                                                                                                                                    \n"
+"    int2 bounding_box_max = (int2) { 0, 0 };                                                                                                                                                                    \n"
+"    int2 clamper = (int2) { *screen_width - 1, *screen_width - 1 };                                                                                                                                               \n"
+"                                                                                                                                                                                                           \n"
+"    for(int i = 0; i < 3; i++) {                                                                                                                                                                           \n"
+"        for(int j = 0; j < 2; j++) {                                                                                                                                                                       \n"
+"            bounding_box_min[j] = max(0, min(bounding_box_min[j], (int)vertices3[j][i]));                                                                                                                  \n"
+"            bounding_box_max[j] = min(clamper[j], max(bounding_box_max[j], (int)vertices3[j][i]));                                                                                                           \n"
+"        }                                                                                                                                                                                                  \n"
+"    }                                                                                                                                                                                                      \n"
+"                                                                                                                                                                                                           \n"
+"    if(bounding_box_min[0] > *screen_width || bounding_box_max[0] < 0 || bounding_box_min[1] > *screen_width || bounding_box_max[1] < 0) return;                                                           \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"    int X_PER_ITEM = (int)(ceil((float)(bounding_box_max[0] - bounding_box_min[0]) / 16.f));                                                                                                                                                                                            \n"
+"    int Y_PER_ITEM = (int)(ceil((float)(bounding_box_max[1] - bounding_box_min[1]) / 16.f));                                                                                                                                                                                            \n"
+"    float STARTING_X = bounding_box_min[0] + X_PER_ITEM * (LOCAL_ID % 16);                                                                                                                                                                                            \n"
+"    float ENDING_X = STARTING_X + X_PER_ITEM;                                                                                                                                                                                            \n"
+"    float STARTING_Y = bounding_box_min[1] + Y_PER_ITEM * (LOCAL_ID / 16);                                                                                                                                                                                            \n"
+"    float ENDING_Y = STARTING_Y + Y_PER_ITEM;                                                                                                                                                                                            \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"    float3 point;                                                                                                                                                                                            \n"
+"    for(point[0] = STARTING_X; point[0] <= ENDING_X; point[0]++) {                                                                                                                     \n"
+"        for(point[1] = STARTING_Y; point[1] <= ENDING_Y; point[1]++) {                                                                                                                 \n"
+"           if(point[1] >= *screen_width || point[1] >= *screen_width) break;                                                                                                                                                                                                \n"
+"			float3 bc_coord = barycentric(vertices3, point);                                                                                                                                                \n"
+"           if (bc_coord[0] < 0 || bc_coord[1] < 0 || bc_coord[2] < 0) continue;                                                                                                                                                                                               \n"
+"                                                                                                                                                                                                           \n"
+"           float2 uv_vec = (float2){0, 0};                                                                                                                                                                                  \n"
+"           float3 normal = (float3){0, 0, 0};                                                                                                                                                                \n"
+"                                                                                                                                                                                                           \n"
+"           point[2] = dot(vertices3[2], bc_coord);                                                                                                                                                                                   \n"
+"			if (z_buffer[(int)(point[0] + point[1] * *screen_width)] > point[2]) {                                                                        \n"
+"               continue;                                                                                                                                                                                   \n"
+"           }                                                                                                                                                                                               \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"			for (int i = 0; i < 3; i++) {                                                                                                                                                                   \n"
+"               uv_vec[0]  += uv_coords[i][0] * bc_coord[i];                                                                                                                                                 \n"
+"               uv_vec[1]  += uv_coords[i][1] * bc_coord[i];                                                                                                                                                 \n"
+"               normal[i]  = dot(norms[i], bc_coord);                                                                                                                                                 \n"
+"           }                                                                                                                                                                                               \n"
+"           int2 uv_point = (int2) { (int)(uv_vec[0] * map_size[0]), (int)(uv_vec[1] * map_size[1]) };                                                                                                                    \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"           int col_index = 3 * (uv_point[0] + uv_point[1] * map_size[0]);                                                                                                                                       \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"           float3 normalized_norm = normalize(normal);                                                                                                                                                                                                \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"           float intensity = clamp((dot(normalized_norm , (float3){light_dir[0], light_dir[1], light_dir[2]})), 0.f, 1.f) + 0.2;                                                                                                                                                                                                \n"
+"                                                                                                                                                                                                           \n"
+"           int color = 0;                                                                                                                                                                                  \n"
+"           color |= ((int)fmin((float)(diffuse_map[col_index + 0]) * intensity, (float) 0xff)) << 16;                                                                                                                                                               \n"
+"           color |= ((int)fmin((float)(diffuse_map[col_index + 1]) * intensity, (float) 0xff)) <<  8;                                                                                                                                                               \n"
+"           color |= ((int)fmin((float)(diffuse_map[col_index + 2]) * intensity, (float) 0xff)) <<  0;                                                                                                                                                               \n"
+"                                                                                                                                                                                                          \n"
+"           //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 16;                                                                                                                                                               \n"
+"           //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) <<  8;                                                                                                                                                               \n"
+"           //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) <<  0;                                                                                                                                                               \n"
+"                                                                                                                                                                                                          \n"
+"                                                                                                                                                                                                          \n"
+"           z_buffer[(int) (point[0] + point[1] * *screen_width)] = point[2];                                                                                                                                 \n"
+"			pixels  [(int) (point[0] + point[1] * *screen_width)] = color;// & 0x00ffffff;                                                                                                                                     \n"
+"        }                                                                                                                                                                                                  \n"
+"    }                                                                                                                                                                                                      \n"
+"}                                                                                                                                                                                                          \n";
@@ -4,19 +4,38 @@
 #include <string>
 #include "geometry.h"
 #include "tgaimage.h"
+#include "kernels.h"
+#include "util_window.h"

 class Model {
 private:
+    void load_texture(std::string filename, const char *suffix, TGAImage &img);
+	int fluid;
+public:
+    TGAImage diffusemap_;
+    TGAImage normalmap_;
+    TGAImage specularmap_;
    std::vector<Vec3f> verts_;
    std::vector<std::vector<Vec3i> > faces_; // attention, this Vec3i means vertex/uv/normal
    std::vector<Vec3f> norms_;
    std::vector<Vec2f> uv_;
-    TGAImage diffusemap_;
-    TGAImage normalmap_;
-    TGAImage specularmap_;
-    void load_texture(std::string filename, const char *suffix, TGAImage &img);
-public:
-    Model(const char *filename);
+	cl_program vertex_shader_prog;
+	cl_program fragment_shader_prog;
+	cl_kernel vertex_shader_kernel;
+	cl_kernel fragment_shader_kernel;
+	cl_mem vertex_shader_matz;
+	cl_mem vertex_shader_vertices;
+	cl_mem new_vertices_mem;
+	cl_mem fragment_shader_faces;
+	cl_mem fragment_shader_screen_width;
+	cl_mem fragment_shader_uv;
+	cl_mem fragment_shader_map_size;
+	cl_mem fragment_shader_norms;
+	cl_mem fragment_shader_light_dir;
+	cl_mem fragment_shader_diffuse_map;
+	cl_int3* faces;
+
+    Model(const char *filename, int watery);
    ~Model();
    int nverts();
    int nfaces();
@@ -33,6 +52,11 @@ public:
 	void rotate(Vec3f rot);
 	void scale(Vec3f scl);
 	void ApplyTransform();
+	void init_kernels();
+	void vertex(float* z);
+	void fragment(float* light_dir);
+	void render(Matrix* z, float* light_dir);
+	void release_kernels();
    TGAColor diffuse(Vec2f uv);
    float specular(Vec2f uv);
    std::vector<int> face(int idx);
@@ -1,225 +1,83 @@
-#define _USE_MATH_DEFINES
-#include <vector>
-#include <algorithm>
-#include <limits>
-#include "tgaimage.h"
-#include "model.h"
-#include "geometry.h"
+#include "model.h"
 #include "renderer.h"
 #include "util_window.h"
+#include "camera.h"
+#include "util_renderer.h"
+#include "CL/cl.h"
+#include "kernels.h"
+#include "improv_gfx.h"
 #include <ctime>
-#include "ColladaModel.h"
-#include "COLLADA.h"

-const TGAColor white = TGAColor(255, 255, 255, 255);
-const TGAColor red = TGAColor(255, 0, 0, 255);
-const TGAColor green = TGAColor(0, 255, 0, 255);
-const TGAColor blue = TGAColor(0, 0, 255, 255);
+#pragma comment (lib, "x86_64/opencl.lib")

-const int depth = 255;
+#define HORIZONTAL_CAMERA_SPEED           0.1
+#define VERTICAL_CAMERA_SPEED             0.1
+#define VERTICAL_CAMERA_CLAMP_UP           90
+#define VERTICAL_CAMERA_CLAMP_DOWN        -90
+#define NEAR_CLIP_PLANE                    1.f 
+#define FAR_CLIP_PLANE                     2000.0f
+#define FOV                                50
+#define CAMERA_MOVEMENT_SPEED              .1f
+#define DEFAULT_CAMERA_POS     Vec3f(0, 0, 5)
+#define DEFAULT_CAMERA_ROT     Vec3f(0, 0, 0)
+#define LIGHT_INTENSITY                   1.5

-float* z_buffer;
-Vec3f light_dir = Vec3f(0, 0, 1).normalize();
-Vec3f eye(0, 0, 3);
-Vec3f center(0, 0, 0);
+Matrix ViewPort = Matrix::identity();
+Matrix ModelView = Matrix::identity();
+Matrix Projection = Matrix::identity();

-Matrix viewport(int x, int y, int w, int h) {
-	Matrix m = Matrix::identity();
-	m[0][3] = x + w / 2.f;
-	m[1][3] = y + h / 2.f;
-	m[2][3] = depth / 2.f;
+Model* model, *model2;
+Camera camera;

-	m[0][0] = w / 2.f;
-	m[1][1] = h / 2.f;
-	m[2][2] = depth / 2.f;
-	return m;
-}
+Vec3f light_dir = Vec3f(0, 1, 1).normalize();
+bool init_flag = false;

-void line(Vec2i p0, Vec2i p1, TGAImage &image, TGAColor color)
-{
-	bool steep = false;
-
-	if (std::abs(p0[0] - p1[0]) < std::abs(p0[1] - p1[1])) {
-		std::swap(p0[0], p0[1]);
-		std::swap(p1[0], p1[1]);
-		steep = true;
-
-	}
-
-	if (p0[0] > p1[0]) {
-		std::swap(p0[0], p1[0]);
-		std::swap(p0[1], p1[1]);
-	}
-
-	int dx = p1[0] - p0[0];
-	int dy = p1[1] - p0[1];
-	int derror2 = std::abs(dy) * 2;
-	int error2 = 0;
-	int y = p0[1];
-	int y_step = p1[1] > p0[1] ? 1 : -1;
-	int dx_2 = 2 * dx;
-
-	for (int x = p0[0]; x <= p1[0]; x++) {
-		if (steep) {
-			image.set(y, x, color);
-		}
-		else {
-			image.set(x, y, color);
-		}
-		error2 += derror2;
-		if (error2 > dx) {
-			y += (y_step);
-			error2 -= dx_2;
-		}
-	}
-}
-
-Vec3f barycentric(Vec3f* pts, Vec3f P)
-{
-	Vec3f u = cross(
-		Vec3f(pts[2][0] - pts[0][0], pts[1][0] - pts[0][0], pts[0][0] - P[0]), // AC_x, AB_x, distance_x
-		Vec3f(pts[2][1] - pts[0][1], pts[1][1] - pts[0][1], pts[0][1] - P[1])  // AC_y, AB_y, distance_y
-	);
-
-	if (std::abs(u[2]) < 1) return Vec3f(-1, 1, 1);
-	return Vec3f(1.f - (u.x + u.y) / u.z, u.y / u.z, u.x / u.z);
+void init_camera() {
+	camera.SetPosition(DEFAULT_CAMERA_POS);
+	camera.SetRotation(DEFAULT_CAMERA_ROT);
+	camera.SetFOV(FOV);
+	camera.SetNearPlane(NEAR_CLIP_PLANE);
+	camera.SetFarPlane(FAR_CLIP_PLANE);
+	camera.SetClampRotDown(VERTICAL_CAMERA_CLAMP_DOWN);
+	camera.SetClampRotUp(VERTICAL_CAMERA_CLAMP_UP);
+	camera.SetHorizontalRotSpeed(HORIZONTAL_CAMERA_SPEED);
+	camera.SetVerticalRotSpeed(VERTICAL_CAMERA_SPEED);
+	camera.SetMovementSpeed(CAMERA_MOVEMENT_SPEED);
+	camera.ApplyChanges();
 }


-void triangle(
-	Vec3f* pts,         // Needed
-	ColladaModel* model,       // Should be removed
-	Vec2f* diff_pts,    // Should be removed
-	float* intensities,
-	Vec3f camera_pos)   // Not really sure yet
-{
+std::vector<Model*> models_in_scene;

-	if (pts[0].y == pts[1].y && pts[0].y == pts[2].y) return; // i dont care about degenerate triangles
-	if (pts[0].y > pts[1].y) {
-		std::swap(pts[0], pts[1]);
-		std::swap(diff_pts[0], diff_pts[1]);
-		std::swap(intensities[0], intensities[1]);
-	}
-	if (pts[0].y > pts[2].y) {
-		std::swap(pts[0], pts[2]);
-		std::swap(diff_pts[0], diff_pts[2]);
-		std::swap(intensities[0], intensities[2]);
-	}
-	if (pts[1].y > pts[2].y) {
-		std::swap(pts[1], pts[2]);
-		std::swap(diff_pts[1], diff_pts[2]);
-		std::swap(intensities[1], intensities[2]);
-	}
-	      Vec2i bounding_box_min(screen_width - 1, screen_height - 1);
-        Vec2i bounding_box_max(0, 0);
-        Vec2i clamp(screen_width - 1, screen_height - 1);
-        TGAColor color = white;
-
-        for(int i = 0; i < 3; i++) {
-                for(int j =0; j < 2; j++) {
-                        bounding_box_min[j] = std::fmax(0, std::fmin(bounding_box_min[j], (int)pts[i][j]));
-                        bounding_box_max[j] = std::fmin(clamp[j], std::fmax(bounding_box_max[j], (int)pts[i][j]));
-                }
-        }
-
-        Vec3f P;
-        for(P.x = bounding_box_min.x; P.x <= bounding_box_max.x; P.x++) {
-                for(P.y = bounding_box_min.y; P.y <= bounding_box_max.y; P.y++) {
-                        Vec3f bc_coord = barycentric(pts, P);
-                        if(bc_coord.x < 0 || bc_coord.y < 0 || bc_coord.z < 0) continue;
-
-
-                        float intensity =
-                                intensities[0]
-                                + (intensities[1] - intensities[0]) * bc_coord[1]
-                                + (intensities[2] - intensities[0]) * bc_coord[2];
-
-
-                        // Interpolating Z using the barycentric coordinates
-                        P.z = 0;
-                        for(int i = 0; i < 3; i++) P.z += pts[i][2] * bc_coord[i];
-
-                        // Coloring according to the Z-Buffer
-                        if (P.z > z_buffer[(int)(P.x + P.y * screen_width)])
-                        {
-                                z_buffer[(int)(P.x + P.y * screen_width)] = P.z;
-
-                                // If diff_pts (Diffusemap Points) were passed, then find the
-								                     // color of the current pixel
-                                if(diff_pts) {
-                                        Vec2f diff_pt =
-                                                diff_pts[0]
-                                                + (diff_pts[1] - diff_pts[0]) * bc_coord[1]
-                                                + (diff_pts[2] - diff_pts[0]) * bc_coord[2];
-
-                                        color = model->diffuse(diff_pt);
-                                }
-								color = color * intensity;
-                                set_pixel(P.x, P.y, color_to_int(color));
-                        }
-                }
-        }
-}
-
-int color_to_int(TGAColor col) {
-	return (col[2] << 16) | (col[1] << 8) | col[0];
-}
-
-void init_zbuffer()
-{
-	z_buffer = new float[screen_width*screen_height];
-	for (int i = 0; i < screen_width * screen_height; i++)
-		z_buffer[i] = INT_MIN;
-}
-
-Matrix lookat(Vec3f eye, Vec3f center, Vec3f up) {
-	Vec3f z = (eye - center).normalize();
-	Vec3f x = cross(up, z).normalize();
-	Vec3f y = cross(z, x).normalize();
-	Matrix Minv = Matrix::identity();
-	Matrix Tr = Matrix::identity();
-	for (int i = 0; i < 3; i++) {
-		Minv[0][i] = x[i];
-		Minv[1][i] = y[i];
-		Minv[2][i] = z[i];
-		Tr[i][3] = -center[i];
-	}
-	return Minv * Tr;
-}

 void render()
 {
-	COLLADA C("sssssssssssss.dae");
-	ColladaModel* model = new ColladaModel("african_head.dae");
-	Matrix ViewPort = viewport(screen_width / 8, screen_height / 8, screen_width * 3 / 4, screen_height * 3 / 4);
-	Matrix Projection = Matrix::identity();
-	Matrix ModelView = lookat(eye, center, Vec3f(0, 1, 0));
-
-	Projection[3][2] = -1.f / (eye - center).norm();
-
-	Matrix z = ViewPort * Projection * ModelView * model->Transform;
-
-	init_zbuffer();
-	for (int i = 0; i < model->nfaces(); i++)
-	{
-		std::vector<int> face = model->face(i);
-		Vec3f screen_coords[3];
-		Vec3f world_coords[3];
-		Vec2f diffuse_coords[3];
-		float intensities[3];
-
-		for (int j = 0; j < 3; j++)
-		{
-			Vec3f v = model->vertix(face[j]);
-			Vec4f v4(v);
-			Vec3f coord(z * v4);
-
-			screen_coords[j] = coord;
-			world_coords[j] = v;
-			diffuse_coords[j] = model->uv(i, j);
-			intensities[j] = model->normal(i, j) * light_dir;
-		}
-
-		triangle(screen_coords, model, diffuse_coords, intensities, Vec3f(0, 0, 5));
+	if (!init_flag) {
+		viewport(0, 0, screen_width, screen_height, FAR_CLIP_PLANE, NEAR_CLIP_PLANE);
+		Setup();
+		init_flag = true;
 	}
-}
+
+	{
+		Projection = camera.GetProjectionMatrix();
+		ModelView = camera.GetModelViewMatrix();
+	}
+
+	Matrix z = ViewPort * Projection * ModelView;
+
+	new_frame();
+	Draw();
+	for(Model* model : models_in_scene)
+		model->render(&z, (float*)&light_dir);
+	end_frame();
+
+}
+
+void AddModel(Model* model) {
+	models_in_scene.push_back(model);
+}
+
+void free_renderer() {
+	for(Model* model : models_in_scene)
+		delete model;
+}
@@ -1,7 +1,16 @@
 #ifndef RENDERER_HEADER
 #define RENDERER_HEADER
 #include "tgaimage.h"
+#include "camera.h"

+
+extern float* z_buffer;
+extern Camera camera;
+extern float TIME;
+
+
+void init_camera();
 void render();
+void free_renderer();
 int color_to_int(TGAColor col);
 #endif
@@ -62,14 +62,13 @@ struct TGAColor {

 class TGAImage {
 protected:
-    unsigned char* data;
    int width;
    int height;
    int bytespp;
-
    bool   load_rle_data(std::ifstream &in);
    bool unload_rle_data(std::ofstream &out);
 public:
+    unsigned char* data;
    enum Format {
        GRAYSCALE=1, RGB=3, RGBA=4
    };
@@ -0,0 +1,90 @@
+#include "util_renderer.h"
+#include "util_window.h"
+#include "kernels.h"
+
+IShader::~IShader() {}
+
+float* z_buffer = new float[screen_width * screen_height];
+
+void viewport(int x, int y, int w, int h, int far_plane, int near_plane) {
+	ViewPort[0][3] = x + w / 2.f;
+	ViewPort[1][3] = y + h / 2.f;
+	ViewPort[2][3] = (far_plane-near_plane) / 2.f;
+
+	ViewPort[0][0] = w / 2.f;
+	ViewPort[1][1] = h / 2.f;
+	ViewPort[2][2] = (far_plane+near_plane) / 2.f;
+}
+
+int color_to_int(TGAColor col) {
+	return (col[0] << 24) | (col[1] << 16) | (col[2] << 8) | col[3];
+}
+
+Vec3f barycentric(Vec3f* pts, Vec3f P)
+{
+	Vec3f u = cross(
+		Vec3f(pts[2][0] - pts[0][0], pts[1][0] - pts[0][0], pts[0][0] - P[0]), // AC_x, AB_x, distance_x
+		Vec3f(pts[2][1] - pts[0][1], pts[1][1] - pts[0][1], pts[0][1] - P[1])  // AC_y, AB_y, distance_y
+	);
+
+	if (std::abs(u[2]) < 1) return Vec3f(-1, 1, 1);
+	return Vec3f(1.f - (u.x + u.y) / u.z, u.y / u.z, u.x / u.z);
+}
+
+void triangle( Vec4f* pts, IShader &shader)
+{
+	Vec3f pts3[3];
+	for (int i = 0; i < 3; i++)
+		pts3[i] = Vec3f(pts[i]);
+
+	if (pts3[0].y == pts3[1].y && pts3[0].y == pts3[2].y) return; // i dont care about degenerate triangles
+	//if (pts3[0].y > pts3[1].y) { std::swap(pts3[0], pts3[1]); }
+	//if (pts3[0].y > pts3[2].y) { std::swap(pts3[0], pts3[2]); }
+	//if (pts3[1].y > pts3[2].y) { std::swap(pts3[1], pts3[2]); }
+
+	Vec2i bounding_box_min(screen_width - 1, screen_height - 1);
+	Vec2i bounding_box_max(0, 0);
+	Vec2i clamp(screen_width - 1, screen_height - 1);
+
+	#pragma omp parallel for
+	for (int i = 0; i < 3; i++)
+		for (int j = 0; j < 2; j++) {
+			bounding_box_min[j] = std::fmax(0, std::fmin(bounding_box_min[j], (int)pts3[i][j]));
+			bounding_box_max[j] = std::fmin(clamp[j], std::fmax(bounding_box_max[j], (int)pts3[i][j]));
+		}
+
+	//printf("%f, %f, %f\n", pts3[0][0], pts3[0][1], pts3[0][2]);
+	//float* points_test = (float*)&pts3;
+
+	//printf("%f, %f, %f\n", points_test[0], points_test[1], points_test[2]);
+	//printf("%d, %d\n", bounding_box_max[0], bounding_box_max[1]);
+
+	//fragment_shader_pixelwise(
+	//	*(float**)((Vec3f*)&pts3),
+	//	*(int**)((Vec2i*)&bounding_box_min),
+	//	*(int**)((Vec2i*)&bounding_box_max));
+
+	//fragment_shader_pixelwise(
+	//	(float*)(&pts3),
+	//	(int*)(&bounding_box_min),
+	//	(int*)(&bounding_box_max));
+
+	Vec3i P;
+	#pragma omp parallel for
+	for (P.x = bounding_box_min.x; P.x <= bounding_box_max.x; P.x++) {
+		for (P.y = bounding_box_min.y; P.y <= bounding_box_max.y; P.y++) {
+			//Vec3f bc_coord = barycentric(pts3, P);
+			//float frag_depth = 0;
+			//for (int i = 0; i < 3; i++)
+			//	frag_depth += pts3[i][2] * bc_coord[i];
+			//if (bc_coord.x < 0 || bc_coord.y < 0 || bc_coord.z < 0 || z_buffer[ P.x + P.y * screen_width ]>frag_depth) continue;
+			//TGAColor color;
+			//bool discard = shader.fragment(bc_coord, color);
+			//if (!discard) {
+			//	z_buffer[P.x + P.y * screen_width] = frag_depth;
+			//	set_pixel(P.x, P.y, color_to_int(color));
+			    set_pixel(P.x, P.y, 0x00ff00);
+			//}
+		}
+	}
+}
@@ -0,0 +1,22 @@
+#pragma once
+#include "tgaimage.h"
+#include "model.h"
+#include "geometry.h"
+
+extern Matrix ModelView;
+extern Matrix ViewPort;
+extern Matrix Projection;
+
+void viewport(int x, int y, int w, int h, int far, int near);
+
+struct IShader {
+    virtual ~IShader();
+    virtual Vec4f vertex(int iface, int nthvert) = 0;
+    virtual bool fragment(Vec3f bar, TGAColor &color) = 0;
+};
+
+//void triangle(Vec4f *pts, IShader &shader, TGAImage &image, TGAImage &zbuffer);
+
+//void triangle( Vec3f* pts, Vec2f* diff_pts, Model* model, float* intensities)
+
+void triangle(Vec4f* pts, IShader &shader);
@@ -8,8 +8,6 @@ BITMAPINFO info;
 HBITMAP hbm;
 const int BITCOUNT_PER_PIXEL = 24;
 const int title_height = 39;
-int screen_width = 0;
-int screen_height = 0;
 long long bytes_per_row;
 bool screen_changed = false;
 HDC hdc;
@@ -17,15 +15,13 @@ HDRAWDIB hdd;
 HDC bitmap_dc;
 HGDIOBJ old_obj;

-HWND create_window(int width, int height, HINSTANCE &hInstance) {
+HWND create_window(HINSTANCE &hInstance) {
 	HWND hWnd;
 	WNDCLASSEX wnd_class = { 0 };
 	init_wnd_class(wnd_class, hInstance);

 	RegisterClassEx(&wnd_class);

-	screen_width = width;
-	screen_height = height;
 	create_hwnd(hWnd, hInstance);
 	init(hWnd);

@@ -99,14 +95,24 @@ void destroy_window() {
 void set_pixel(unsigned int x, unsigned int y, unsigned int color) {
 	unsigned long pixel_index = y * bytes_per_row + x * 3;

-	pixel_data[pixel_index + 0] = (char)(color >>  0);
+	pixel_data[pixel_index + 2] = (char)(color >>  0);
 	pixel_data[pixel_index + 1] = (char)(color >>  8);
-	pixel_data[pixel_index + 2] = (char)(color >> 16);
+	pixel_data[pixel_index + 0] = (char)(color >> 16);

 	if (!screen_changed)
 		screen_changed = true;
 }

-void Update() {
-   DrawDibDraw(hdd, hdc, 0, 0, screen_width, screen_height, &info.bmiHeader, pixels, 0, 0, screen_width, screen_height, 0);
+void clear_screen() {
+	for (int x = 0; x < screen_width; x++)
+		for (int y = 0; y < screen_height; y++)
+			set_pixel(x, y, 0);
+}
+
+void Update() {
+	if (screen_changed) {
+		DrawDibDraw(hdd, hdc, 0, 0, screen_width, screen_height, &info.bmiHeader, pixels, 0, 0, screen_width, screen_height, 0);
+		clear_screen();
+		screen_changed = false;
+	}
 }
@@ -18,18 +18,19 @@ extern HDC hdc;
 extern HDRAWDIB hdd;
 extern HDC bitmap_dc;
 extern HGDIOBJ old_obj;
-extern int screen_width;
-extern int screen_height;
+extern const int screen_width;
+extern const int screen_height;


 void init(HWND &hWnd);
 void destroy_window(); 
-HWND create_window(int width, int height, HINSTANCE &hInstance);
+HWND create_window(HINSTANCE &hInstance);
 void create_hwnd(HWND &hwnd, HINSTANCE &hInstance);
 void init_wnd_class(WNDCLASSEX &wndClass, HINSTANCE &hInstance);
 void set_pixel(unsigned int x, unsigned int y, unsigned int color);
 void Update();

 LRESULT CALLBACK WndProc(HWND hwnd, UINT message, WPARAM wParam, LPARAM lParam);
+void CALLBACK FixedUpdate(HWND hwnd, UINT message, UINT uInt, DWORD dWord);

 #endif
Author	SHA1	Message	Date
mo7sen	9a2f6129bc	Sample Scene for Project Discussion with TA	2019-12-29 23:39:33 +02:00
mo7sen	5dd2de5d77	Made it human-friendly	2019-12-27 22:53:14 +02:00
mo7sen	6cef4df543	Improving parallelism in fragment shader leading to more than 20x the framerate	2019-12-26 23:07:03 +02:00
mo7sen	2e6ddd2ff2	Adding basic texture fragment shader kernel	2019-12-25 00:09:14 +02:00
mo7sen	e57702936e	Vertex Shader is now GPU-Accelerated	2019-12-18 19:09:00 +02:00
mo7sen	a47ccbc3cd	Adding some parallelism, Not working yet	2019-12-17 21:24:03 +02:00
mo7sen	d156fb4044	Changing controls and probably some matrices	2019-12-16 23:19:46 +02:00
mo7sen	8b1546cb54	Added Specular map to the shader	2019-12-10 13:47:08 +02:00
mo7sen	26a7e48784	Adding Shaders and a 'Time' Variable	2019-12-10 01:25:24 +02:00
mo7sen	b936db0ce9	Fixing Camera bugs	2019-12-09 12:24:34 +02:00
mo7sen	b51c63ea0c	Adding a buggy FPS Camera with position and rotation	2019-12-09 05:46:30 +02:00
mo7sen	ad56096eb3	Editing the Projection Matrix	2019-12-08 20:57:18 +02:00
mo7sen	2f5437396c	Added messed up low performance camera rotation	2019-12-08 18:47:16 +02:00