Improving parallelism in fragment shader leading to more than 20x the framerate

OpenWindow/ColladaModel.cpp

@@ -1,297 +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_(), rootjoint()
-{
-	Transform = Matrix::identity();
-	Rotation = Matrix::identity();
-	Scale = Matrix::identity();
-	Translation = Matrix::identity();
-
-	tinyxml2::XMLDocument doc;
-	doc.LoadFile(filename);
-
-	load_texture(filename, "_diffuse.tga", diffusemap_);
-	load_texture(filename, "_nm_tangent.tga", normalmap_);
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	////////////////////////////////////////////////faces//////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	face_count = 0;
-	tinyxml2::XMLElement* xml_face = doc.FirstChildElement("COLLADA")->FirstChildElement("library_geometries")->FirstChildElement()->FirstChildElement()->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////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	int 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);
-	}
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	/////////////////////////////////////////////textcoord/////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	int 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;
-		texturecos_.push_back(temp);
-	}
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	int vertexweights_count = 0;
-	tinyxml2::XMLElement* xml_vertexweights = doc.FirstChildElement("COLLADA")->FirstChildElement("library_controllers")->FirstChildElement()->FirstChildElement()->FirstChildElement("vertex_weights");
-	xml_vertexweights->QueryIntAttribute("count", &vertexweights_count);
-	std::stringstream str_vertexweights(xml_vertexweights->FirstChildElement("vcount")->GetText());
-	std::stringstream str_vertexweights1(xml_vertexweights->FirstChildElement("v")->GetText());
-
-	for (int i = 0; i < vertexweights_count; i++)
-	{
-		int temp;
-		Vec2i temp1;
-		vertexweights_.push_back(std::vector<Vec2i>());
-		str_vertexweights >> temp;
-		
-		for (int j = 0; j < temp; j++)
-		{
-			str_vertexweights1 >> temp1.x;
-			str_vertexweights1 >> temp1.y;
-			vertexweights_[i].push_back(temp1);
-		}
-	}
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	int weight_count = 0;
-	tinyxml2::XMLElement* xml_weight = doc.FirstChildElement("COLLADA")->FirstChildElement("library_controllers")->FirstChildElement()->FirstChildElement()->FirstChildElement("source")->NextSiblingElement()->NextSiblingElement()->FirstChildElement();
-	xml_weight->QueryIntAttribute("count", &weight_count);
-	std::stringstream str_weight(xml_weight->GetText());
-
-	for (int i = 0; i < weight_count; i++)
-	{
-		float temp;
-		str_weight >> temp;
-		weights_.push_back(temp);
-	}
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	/*joint_count = 0;
-	tinyxml2::XMLElement* xml_joint = doc.FirstChildElement("COLLADA")->FirstChildElement("library_controllers")->FirstChildElement()->FirstChildElement()->FirstChildElement("source")->FirstChildElement();
-	xml_joint->QueryIntAttribute("count", &joint_count);
-	std::stringstream str_joint(xml_joint->GetText());
-	 */
-	/*for (int i = 0; i < joint_count; i++)
-	{
-		Vec2f temp;
-		str_texcoord >> temp.x;
-		str_texcoord >> temp.y;
-		texturecos_.push_back(temp);
-	}*/
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	///////////////////////////////////////////////////////////////////////////////////////////////////////
-	//tinyxml2::XMLElement* xml_hierachy = doc.FirstChildElement("COLLADA")->FirstChildElement("library_visual_scenes")->FirstChildElement("visual_scene")->FirstChildElement("node")->FirstChildElement("node");
-	//std::stringstream str_hierachy(xml_hierachy->GetText());
-
-	//ColladaModel::rootjoint = buildjoint(xml_hierachy);
-	joint_count = 0;
-	tinyxml2::XMLElement* xml_joint = doc.FirstChildElement("COLLADA")->FirstChildElement("library_controllers")->FirstChildElement("controller")->FirstChildElement("skin")->FirstChildElement("source")->FirstChildElement("Name_array");
-	xml_joint->QueryIntAttribute("count", &joint_count);
-	std::stringstream str_joint(xml_joint->GetText());
-
-	for (int i = 0; i < vertex_count / 3; i++)
-	{
-		std::string temp;
-		str_vertex >> temp;
-
-		jointIDs_.push_back(temp);
-	}
-
-}
-
-//Joint buildjoint(tinyxml2::XMLElement* xml_hierachy)
-//{
-//	const char* name;
-//	xml_hierachy->QueryStringAttribute("name", &name);
-//	std::stringstream str_transform(xml_hierachy->FirstChildElement("matrix")->GetText());
-//	Matrix transform;
-//	
-//	if (xml_hierachy->FirstChildElement("node") != NULL)
-//	{
-//		xml_hierachy = xml_hierachy->FirstChildElement("node");
-//		do
-//		{
-////			buildjoint
-//		} while (xml_hierachy->NextSiblingElement() != NULL);
-//	}
-//}
-
-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 texturecos_[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();
-}
-
-Joint ColladaModel::getrootjoint()
-{
-	return rootjoint;
-}
-
-void ColladaModel::doanimation(Animation animation)
-{
-	//animator.doanimation(animation);
-}
-
-void ColladaModel::updateanimator()
-{
-	//animator.update();
-}
-
-std::vector<Matrix> ColladaModel::getjointtransforms()
-{
-	std::vector<Matrix> jointMat(joint_count);
-	addjointtoarray(rootjoint, jointMat);
-	return jointMat;
-}
-
-void ColladaModel::addjointtoarray(Joint parent, std::vector<Matrix> jointMat)
-{
-	jointMat[parent.index] = parent.getTransform();
-
-	for (auto& child : parent.children)
-	{
-		addjointtoarray(child, jointMat);
-	}
-}
-

OpenWindow/ColladaModel.h

@@ -1,110 +0,0 @@
-#ifndef __MODEL_s__
-#define __MODEL_s__
-#pragma once
-
-#include "ColladaModel.h"
-#include "tinyxml2.h"
-#include "geometry.h"
-#include "tgaimage.h"
-#include "joint.h"
-#include "animator.h"
-#include "animation.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 joint_count;
-
-	Joint rootjoint;
-	int rootindex=0;
-	char* rootname;
-	Matrix root_transform;
-
-	//Animator animator;
-
-	std::vector<std::vector<Vec3i> > faces_; //vertex/normal/uv
-	std::vector<Vec3f> vertices_;
-	std::vector<Vec3f> normals_;
-	std::vector<Vec2f> texturecos_;
-
-	// #include <map>
-	// map<const char*, int> joint_names;
-	// joints.add(pair<const char*, int>("TORSO", 0));
-
-	// {TORSO, 0}
-	// {LEG,   1}
-	// {HAND,  2}
-
-	//joint_names.get("TORSO");
-
-	// 1. Get jointnames and map them to ids
-	// 2. get each joint and fill in the Joint class
-	// 3. get inversebindtransforms
-
-	// Vector -> Vectors -> (joint , weight)
-	std::vector<std::vector<Vec2i>> vertexweights_; //joint/weight
-	std::vector<std::string> jointIDs_;
-	// { TORSO , LEG, TEZ}
-	// { 0, 1, 2}
-	// jointIDs[0]
-	std::vector<float> weights_;
-
-	TGAImage diffusemap_;
-	TGAImage normalmap_;
-	TGAImage specularmap_;
-
-	void load_texture(std::string filename, const char* suffix, TGAImage& img);
-public:
-	ColladaModel() = default;
-	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);
-
-	/////////////////////////////
-	Joint getrootjoint();
-
-	void doanimation(Animation animation);
-
-	void updateanimator();
-
-	std::vector<Matrix> getjointtransforms();
-
-	void addjointtoarray(Joint parent, std::vector<Matrix>);
-};
-#endif

OpenWindow/Joint.cpp

@@ -1,36 +0,0 @@
-#include "Joint.h"
-
-Joint::Joint(int index, Matrix transform, char* name) {
-	Joint::index = index;
-	_transform = transform;
-	Joint::name = name;
-}
-
-Joint::~Joint() {
-}
-
-void Joint::addChild(Joint child) {
-	children.push_back(child);
-}
-
-Matrix Joint::getTransform() {
-	return _transform;
-}
-
-void Joint::setTransform(Matrix transform) {
-	_transform = transform;
-}
-
-Matrix Joint::getInverseBindTransform() {
-	return _inverseBindtransform;
-}
-
-void Joint::calculateInverseBindTransform(Matrix parentBindTransform) {
-	Matrix bindTransform = parentBindTransform * _localBindTransform;
-	_inverseBindtransform = bindTransform.invert();
-
-	for (auto &child : children)
-	{
-		child.calculateInverseBindTransform(bindTransform);
-	}
-}

OpenWindow/Joint.h

@@ -1,28 +0,0 @@
-#pragma once
-#include <vector>
-#include "geometry.h"
-class Joint
-{
-private:
-	Matrix _transform;	//model space
-	Matrix _localBindTransform;
-	Matrix _inverseBindtransform;
-
-public:
-	int index;
-	char* name;
-	std::vector<Joint> children;
-
-	Joint() = default;
-	Joint(int index,Matrix transform,char* name);
-	~Joint();
-
-	void addChild(Joint child);
-
-	Matrix getTransform();
-	void setTransform(Matrix transform);
-
-	Matrix getInverseBindTransform();
-	void calculateInverseBindTransform(Matrix parentBindTransform);
-};
-

OpenWindow/OpenWindow.vcxproj

@@ -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,44 +127,41 @@
       <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="animation.cpp" />
-    <ClCompile Include="animator.cpp" />
     <ClCompile Include="camera.cpp" />
-    <ClCompile Include="colladamodel.cpp" />
     <ClCompile Include="geometry.cpp" />
-    <ClCompile Include="joint.cpp" />
+    <ClCompile Include="kernels.cpp">
-    <ClCompile Include="jointtransform.cpp" />
+      <AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
-    <ClCompile Include="keyframe.cpp" />
+      <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="quaternion.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="animation.h" />
-    <ClInclude Include="animator.h" />
-    <ClInclude Include="colladamodel.h" />
     <ClInclude Include="geometry.h" />
     <ClInclude Include="camera.h" />
-    <ClInclude Include="joint.h" />
+    <ClInclude Include="kernels.h" />
-    <ClInclude Include="jointtransform.h" />
-    <ClInclude Include="keyframe.h" />
     <ClInclude Include="model.h" />
-    <ClInclude Include="quaternion.h" />
     <ClInclude Include="renderer.h" />
     <ClInclude Include="tgaimage.h" />
-    <ClInclude Include="tinyxml2.h" />
     <ClInclude Include="util_renderer.h" />
     <ClInclude Include="util_window.h" />
   </ItemGroup>

OpenWindow/OpenWindow.vcxproj.filters

@@ -39,28 +39,10 @@
     <ClCompile Include="util_renderer.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
-    <ClCompile Include="colladamodel.cpp">
+    <ClCompile Include="kernels.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
-    <ClCompile Include="tinyxml2.cpp">
+    <ClCompile Include="kernel_sources.cpp">
-      <Filter>Source Files</Filter>
-    </ClCompile>
-    <ClCompile Include="joint.cpp">
-      <Filter>Source Files</Filter>
-    </ClCompile>
-    <ClCompile Include="animator.cpp">
-      <Filter>Source Files</Filter>
-    </ClCompile>
-    <ClCompile Include="animation.cpp">
-      <Filter>Source Files</Filter>
-    </ClCompile>
-    <ClCompile Include="keyframe.cpp">
-      <Filter>Source Files</Filter>
-    </ClCompile>
-    <ClCompile Include="jointtransform.cpp">
-      <Filter>Source Files</Filter>
-    </ClCompile>
-    <ClCompile Include="quaternion.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
   </ItemGroup>
@@ -86,28 +68,7 @@
     <ClInclude Include="util_renderer.h">
       <Filter>Header Files</Filter>
     </ClInclude>
-    <ClInclude Include="colladamodel.h">
+    <ClInclude Include="kernels.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
-    <ClInclude Include="tinyxml2.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
-    <ClInclude Include="joint.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
-    <ClInclude Include="animator.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
-    <ClInclude Include="animation.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
-    <ClInclude Include="keyframe.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
-    <ClInclude Include="jointtransform.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
-    <ClInclude Include="quaternion.h">
       <Filter>Header Files</Filter>
     </ClInclude>
   </ItemGroup>

OpenWindow/animation.cpp

@@ -1,17 +0,0 @@
-#include "animation.h"
-
-Animation::Animation(float seconds, std::vector<Keyframe> keyframes)
-{
-	length = seconds;
-	keyframes_ = keyframes;
-}
-
-float Animation::getlength()
-{
-	return length;
-}
-
-std::vector<Keyframe> Animation::getkeyframe()
-{
-	return keyframes_;
-}

OpenWindow/animation.h

@@ -1,16 +0,0 @@
-#pragma once
-
-#include "keyframe.h"
-#include <vector>
-
-class Animation {
-private:
-	float length;
-	std::vector<Keyframe> keyframes_;
-
-public:
-	Animation() = default;
-	Animation(float seconds, std::vector<Keyframe> keyframes);
-	float getlength();
-	std::vector<Keyframe>getkeyframe();
-};

OpenWindow/animator.cpp

@@ -1,108 +0,0 @@
-#include "animator.h"
-#include "colladamodel.h"
-
-void Animator::increaseAnimationTime()
-{
-	animation_time_ += (1/60.f);
-	if (animation_time_ > current_animation_.getlength()) 
-	{
-		animation_time_ = (int)(animation_time_) % (int)(current_animation_.getlength());
-	}
-}
-
-
-
-void Animator::applyPoseToJoints(std::map<std::string, Matrix> currentPose, Joint joint, Matrix parentTransform)
-{
-	Matrix currentLocalTransform = currentPose.at(joint.name);
-	Matrix currentTransform = parentTransform * currentLocalTransform;
-	for (Joint childJoint : joint.children) {
-		applyPoseToJoints(currentPose, childJoint, currentTransform);
-	}
-	currentTransform = currentTransform * joint.getInverseBindTransform();
-	joint.setTransform(currentTransform);
-}
-
-std::vector<Keyframe> Animator::getPreviousAndNextFrames()
-{
-	std::vector<Keyframe> allFrames = current_animation_.getkeyframe();
-	Keyframe previousFrame = allFrames[0];
-	Keyframe nextFrame = allFrames[0];
-	for (int i = 1; i < allFrames.capacity(); i++) {
-		nextFrame = allFrames[i];
-		if (nextFrame.gettimestamp() > animation_time_) {
-			break;
-		}
-		previousFrame = allFrames[i];
-	}
-	return std::vector<Keyframe>{ previousFrame, nextFrame };
-}
-
-float Animator::calculateProgression(Keyframe previousFrame, Keyframe nextFrame)
-{
-	float totalTime = nextFrame.gettimestamp() - previousFrame.gettimestamp();
-	float currentTime = animation_time_ - previousFrame.gettimestamp();
-	return currentTime / totalTime;
-}
-
-std::map<std::string, Matrix> Animator::interpolatePoses(Keyframe previousFrame, Keyframe nextFrame, float progression)
-{
-	std::map<std::string, Matrix> currentPose ;
-	std::set<std::string> key_set;
-	make_key_set(previousFrame.getJointKeyFrames(), key_set);
-	for (std::string jointName : key_set)
-	{
-		JointTransform previousTransform = previousFrame.getJointKeyFrames().at(jointName);
-		JointTransform nextTransform = nextFrame.getJointKeyFrames().at(jointName);
-		JointTransform currentTransform = JointTransform::interpolate(previousTransform, nextTransform, progression);
-		currentPose.insert_or_assign(jointName, currentTransform.getlocationtransform());
-		make_key_set(previousFrame.getJointKeyFrames(), key_set);
-	}
-	return currentPose;
-}
-
-Animator::Animator(ColladaModel entity)
-{
-	entity_ = &entity;
-}
-
-void Animator::doanimation(Animation animation)
-{
-	animation_time_ = 0;
-	current_animation_ = animation;
-}
-
-void Animator::update()
-{
-	if (&current_animation_ == NULL) {
-		return;
-	}
-	increaseAnimationTime();
-	std::map<std::string, Matrix> currentPose = calculateCurrentAnimationPose();
-	applyPoseToJoints(currentPose, entity_->getrootjoint(), Matrix());
-}
-
-std::map<std::string, Matrix> Animator::calculateCurrentAnimationPose()
-{
-	std::vector<Keyframe> frames = getPreviousAndNextFrames();
-	float progression = calculateProgression(frames[0], frames[1]);
-	return interpolatePoses(frames[0], frames[1], progression);
-}
-
-template<typename TK, typename TV>
-std::vector<TK> extract_keys(std::map<TK, TV> const& input_map) {
-	std::vector<TK> retval;
-	for (auto const& element : input_map) {
-		retval.push_back(element.first);
-	}
-	return retval;
-}
-
-template<typename TK, typename TV>
-std::vector<TV> extract_values(std::map<TK, TV> const& input_map) {
-	std::vector<TV> retval;
-	for (auto const& element : input_map) {
-		retval.push_back(element.second);
-	}
-	return retval;
-}

OpenWindow/animator.h

@@ -1,46 +0,0 @@
-#pragma once
-#include "animation.h"
-#include "geometry.h"
-#include "keyframe.h"
-#include "jointtransform.h"
-#include "joint.h"
-#include <map> 
-#include <set>
-class ColladaModel;
-
-class Animator {
-private:
-	ColladaModel* entity_;
-	Animation current_animation_;
-	float animation_time_;
-
-	void increaseAnimationTime();
-	std::map<std::string,Matrix> calculateCurrentAnimationPose();
-	void applyPoseToJoints(std::map<std::string, Matrix> currentPose, Joint joint, Matrix parentTransform);
-	std::vector<Keyframe> getPreviousAndNextFrames();
-	float calculateProgression(Keyframe previousFrame, Keyframe nextFrame);
-	std::map<std::string, Matrix> interpolatePoses(Keyframe previousFrame, Keyframe nextFrame, float progression);
-
-public:
-	Animator(ColladaModel entity);
-	void doanimation(Animation animation);
-	void update();	
-};
-
-template< class Key,
-	class T,
-	class Comparator,
-	class MapAllocator,
-	class SetAllocator>
-	void make_key_set(const std::map<Key, T, Comparator, MapAllocator>& map,
-		std::set<Key, Comparator, SetAllocator>& set)
-{
-	set.clear();
-	typedef typename std::map<Key, T, Comparator, MapAllocator> map_type;
-	typename map_type::const_iterator itr = map.begin();
-	while (map.end() != itr)
-	{
-		set.insert((itr++)->first);
-	}
-}
-

OpenWindow/camera.cpp

@@ -1,4 +1,5 @@
 #include "camera.h"
+#include "util_window.h"
 
 Camera::Camera()
 {
@@ -100,22 +101,24 @@ Matrix Camera::GetModelViewMatrix() {
 	Vec3f center = position + forward;
 	Vec3f z = forward * -1;
 	Matrix Minv = Matrix::identity();
-	Matrix Tr = Matrix::identity();
 	for (int i = 0; i < 3; i++) {
 		Minv[0][i] = right[i];
 		Minv[1][i] = up[i];
 		Minv[2][i] = z[i];
-		Tr[i][3] = -center[i];
 	}
-	return Minv * Tr;
+	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);
+	Projection[0][0] = 1 / tan(fov * DEG2RAD / 2);
-	Projection[1][1] = 1 / tan(fov * DEG2RAD);
+	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[2][3] = (2 * far_plane * near_plane) / (far_plane - near_plane);
 	Projection[3][2] = -1;
+	Projection[3][3] = 0;
 	return Projection;
 }

OpenWindow/include/CL/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_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 */
+

OpenWindow/include/CL/cl_d3d11.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 */
+

OpenWindow/include/CL/cl_dx9_media_sharing.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 */
+

OpenWindow/include/CL/cl_ext.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 */

OpenWindow/include/CL/cl_gl.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 */

OpenWindow/include/CL/cl_gl_ext.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  */

OpenWindow/include/CL/opencl.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   */
+

OpenWindow/jointtransform.cpp

@@ -1,34 +0,0 @@
-#include "jointtransform.h"
-
-
-
-JointTransform::JointTransform(Vec3f position, Quaternion rotation)
-{
-	position_ = position;
-	rotation_ = rotation;
-}
-
-Matrix JointTransform::getlocationtransform()
-{
-	Matrix matrix;
-	Vec4f column = Vec4f(position_.x, position_.y, position_.z, 1);
-	matrix.set_col(3, column);
-	matrix = matrix * rotation_.toRotationMatrix();
-	return matrix;
-}
-
-JointTransform JointTransform::interpolate(JointTransform frame1, JointTransform frame2, float progression)
-{
-	Vec3f pos = interpolate(frame1.position_, frame2.position_, progression);
-	Quaternion rot = Quaternion::interpolate(frame1.rotation_, frame2.rotation_, progression);
-	return JointTransform(pos, rot);
-}
-
-Vec3f JointTransform::interpolate(Vec3f start, Vec3f end, float progression)
-{
-	float x = start.x + (end.x - start.x) * progression;
-	float y = start.y + (end.y - start.y) * progression;
-	float z = start.z + (end.z - start.z) * progression;
-
-	return Vec3f(x,y,z);
-}

OpenWindow/jointtransform.h

@@ -1,24 +0,0 @@
-#pragma once
-
-#include "quaternion.h"
-#include "geometry.h"
-
-class JointTransform
-{
-private:
-
-	Vec3f position_;
-
-	Quaternion rotation_;
-
-public:
-	JointTransform() = default;
-	JointTransform(Vec3f position,Quaternion rotation);
-
-	Matrix getlocationtransform();
-
-	static JointTransform interpolate(JointTransform frame1, JointTransform frame2, float progression);
-
-	static Vec3f interpolate(Vec3f start,Vec3f end, float progression);
-};
-

OpenWindow/kernel_sources.cpp

@@ -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";

OpenWindow/kernels.cpp

@@ -0,0 +1,237 @@
+#include "kernels.h"
+#include "util_window.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 triangles_drawn = 0;
+
+// Matrix Multiplication Variables
+#define MATRIX_SIZE 16
+
+// vertex shader
+cl_mem mat_z;
+cl_mem vertices_mem;
+cl_mem new_vertices_mem;
+bool vertex_shader_buffers_initialized = false;
+//==================
+// fragment shader
+//	int* faces,
+//	int nfaces,
+//	float* new_verts,
+//	int nverts,
+//	int screen_width,
+//	int screen_height,
+//	float* z_buffer,
+//	float* uniform_m,
+//	float* uniform_mit,
+//	float* light_dir,
+//	const char* diffuse_map,
+//	const char* normal_map,
+//	const char* spec_map,
+//	char* pixel_data
+cl_mem faces_buffer;
+cl_mem uniform_m;
+cl_mem uniform_mit;
+cl_mem light_dir_buffer;
+cl_mem diffuse_map_buffer;
+cl_mem norms_mem;
+cl_mem spec_map_buffer;
+int* hidden_pixel_buffer;
+bool fragment_shader_buffers_initialized = false;
+cl_mem triangles_verts_mem;
+cl_mem bounding_box_min_mem;
+cl_mem z_buffer_mem;
+cl_mem pixel_data_buffer;
+cl_mem screen_width_mem;
+cl_mem nfaces_mem;
+cl_mem uv_buffer;
+cl_mem map_size_buffer;
+//==================
+
+cl_platform_id* platforms = NULL;
+cl_device_id* devices = NULL;
+
+cl_program vertex_shader_prog;
+cl_program fragment_shader_prog;
+cl_kernel vertex_shader_kernel;
+cl_kernel fragment_shader_kernel;
+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);
+
+	vertex_shader_prog = clCreateProgramWithSource(context, 1, (const char **)&vertex_shader_kernel_source, NULL, &err);
+	fragment_shader_prog = clCreateProgramWithSource(context, 1, (const char**)&fragment_shader_kernel_source, NULL, &err);
+	err = clBuildProgram(vertex_shader_prog, 1, devices, NULL, NULL, NULL);
+	err = clBuildProgram(fragment_shader_prog, 1, devices, NULL, NULL, NULL);
+
+	if (err != CL_SUCCESS)
+	{
+		size_t len;
+		char buffer[2048];
+		printf("Error: Failed to build the fragment shader prog!\n");
+		clGetProgramBuildInfo(fragment_shader_prog, devices[0], CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
+		printf(buffer);
+		//return;
+	}
+
+	vertex_shader_kernel = clCreateKernel(vertex_shader_prog, "vertex_shader", &err);
+	fragment_shader_kernel = clCreateKernel(fragment_shader_prog, "fragment_shader", &err);
+	err = 0;
+}
+
+void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices) {
+	if(!vertex_shader_buffers_initialized)
+	{
+		mat_z            = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * MATRIX_SIZE     , NULL, &err);
+		vertices_mem     = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * vertex_count * 3, NULL, &err);
+		new_vertices_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float4) * vertex_count, NULL, &err);
+
+		err = clSetKernelArg(vertex_shader_kernel, 0, sizeof(cl_mem), &mat_z);
+		err = clSetKernelArg(vertex_shader_kernel, 1, sizeof(cl_mem), &vertices_mem);
+		err = clSetKernelArg(vertex_shader_kernel, 2, sizeof(cl_mem), &new_vertices_mem);
+		vertex_shader_buffers_initialized = true;
+	}
+
+	err = clEnqueueWriteBuffer(commands, mat_z       , CL_TRUE, 0, sizeof(float) * MATRIX_SIZE, z, 0, NULL, NULL);
+	err = clEnqueueWriteBuffer(commands, vertices_mem, CL_TRUE, 0, sizeof(float) * vertex_count * 3, vertices, 0, NULL, NULL);
+
+	size_t vertex_shader_global[] = { vertex_count*4 };
+	size_t vertex_shader_local[] = { 4 };
+	err = clEnqueueNDRangeKernel(commands, vertex_shader_kernel, 1, NULL, vertex_shader_global, vertex_shader_local, 0, NULL, NULL);
+}
+
+void clear_pixel_buffer() {
+	for (int i = 0; i < screen_width * screen_height; i++) {
+		hidden_pixel_buffer[i] = 0;
+	}
+}
+
+void clear(cl_mem* buffer, size_t size, const int pattern) {
+	clEnqueueFillBuffer(commands, *buffer, &pattern, sizeof(int), 0, size, 0, NULL, NULL);
+}
+
+void fragment_shader(
+	cl_int3* faces,
+	int nfaces,
+	float* uv,
+	size_t uv_size,
+	float* uniform_m,
+	float* uniform_mit,
+	float* light_dir,
+	unsigned char* diffuse_map,
+	float* norms,
+	size_t norms_size,
+	unsigned char* spec_map,
+	int* map_size
+) {
+	if (!fragment_shader_buffers_initialized) {
+
+		faces_buffer       = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(cl_int3) * 3 * nfaces                , NULL, &err);
+		pixel_data_buffer  = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * screen_height * screen_width  , NULL, &err);
+		screen_width_mem   = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int)                                 , NULL, &err);
+		z_buffer_mem       = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * screen_height * screen_width, NULL, &err);
+		nfaces_mem         = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int)                                 , NULL, &err);
+		uv_buffer          = clCreateBuffer(context, CL_MEM_READ_ONLY , uv_size                                     , NULL, &err);
+		map_size_buffer    = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int) * 2                             , NULL, &err);
+		norms_mem          = clCreateBuffer(context, CL_MEM_READ_ONLY , norms_size                                  , NULL, &err);
+		light_dir_buffer   = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * 3                           , NULL, &err);
+		diffuse_map_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(char) * 3 * map_size[0] * map_size[1]    , NULL, &err);
+
+
+		err = clSetKernelArg(fragment_shader_kernel, 0, sizeof(cl_mem), &faces_buffer);
+		err = clSetKernelArg(fragment_shader_kernel, 1, sizeof(cl_mem), &new_vertices_mem);
+		err = clSetKernelArg(fragment_shader_kernel, 2, sizeof(cl_mem), &pixel_data_buffer);
+		err = clSetKernelArg(fragment_shader_kernel, 3, sizeof(cl_mem), &screen_width_mem);
+		err = clSetKernelArg(fragment_shader_kernel, 4, sizeof(cl_mem), &z_buffer_mem);
+		err = clSetKernelArg(fragment_shader_kernel, 5, sizeof(cl_mem), &nfaces_mem);
+		err = clSetKernelArg(fragment_shader_kernel, 6, sizeof(cl_mem), &uv_buffer);
+		err = clSetKernelArg(fragment_shader_kernel, 7, sizeof(cl_mem), &map_size_buffer);
+		err = clSetKernelArg(fragment_shader_kernel, 8, sizeof(cl_mem), &light_dir_buffer);
+		err = clSetKernelArg(fragment_shader_kernel, 9, sizeof(cl_mem), &norms_mem);
+		err = clSetKernelArg(fragment_shader_kernel, 10, sizeof(cl_mem), &diffuse_map_buffer);
+
+		err = clEnqueueWriteBuffer(commands, screen_width_mem  , CL_FALSE, 0, sizeof(int)                             , &screen_width , 0, NULL, NULL);
+		err = clEnqueueWriteBuffer(commands, faces_buffer      , CL_FALSE, 0, sizeof(cl_int3) * 3 * nfaces            , faces         , 0, NULL, NULL);
+		err = clEnqueueWriteBuffer(commands, nfaces_mem        , CL_FALSE, 0, sizeof(int)                             , &nfaces       , 0, NULL, NULL);
+		err = clEnqueueWriteBuffer(commands, uv_buffer         , CL_FALSE, 0, uv_size                                 , uv            , 0, NULL, NULL);
+		err = clEnqueueWriteBuffer(commands, map_size_buffer   , CL_FALSE, 0, sizeof(int) * 2                         , map_size      , 0, NULL, NULL);
+		err = clEnqueueWriteBuffer(commands, light_dir_buffer  , CL_FALSE, 0, sizeof(float) * 3                       , light_dir     , 0, NULL, NULL);
+		err = clEnqueueWriteBuffer(commands, norms_mem         , CL_FALSE, 0, norms_size                              , norms         , 0, NULL, NULL);
+		err = clEnqueueWriteBuffer(commands, diffuse_map_buffer, CL_FALSE, 0, sizeof(char) * 3 * map_size[0] * map_size[1], diffuse_map   , 0, NULL, NULL);
+
+		hidden_pixel_buffer = new int[screen_width*screen_height];
+		fragment_shader_buffers_initialized = true;
+	}
+
+	clear(&pixel_data_buffer, sizeof(int) * screen_width * screen_height, 0);
+	clear(&z_buffer_mem, sizeof(float) * screen_width * screen_height, 0);
+
+
+	size_t fragment_shader_global[] = { nfaces * 256 };
+	size_t framgent_shader_local[] = { 256 };
+
+	err = clEnqueueNDRangeKernel(commands, fragment_shader_kernel, 1, NULL, fragment_shader_global, NULL, 0, NULL, NULL);
+
+	err = 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()
+{
+	clReleaseProgram(vertex_shader_prog);
+	clReleaseKernel(vertex_shader_kernel);
+	if (vertex_shader_buffers_initialized) {
+		clReleaseMemObject(mat_z);
+		clReleaseMemObject(vertices_mem);
+		clReleaseMemObject(new_vertices_mem);
+	}
+
+	if (fragment_shader_buffers_initialized) {
+		clReleaseMemObject(faces_buffer);
+		clReleaseMemObject(screen_width_mem);
+		clReleaseMemObject(z_buffer_mem);
+		clReleaseMemObject(pixel_data_buffer);
+		clReleaseMemObject(screen_width_mem);
+		clReleaseMemObject(nfaces_mem);
+		clReleaseMemObject(uv_buffer);
+		clReleaseMemObject(map_size_buffer);
+		clReleaseMemObject(light_dir_buffer);
+		clReleaseMemObject(norms_mem);
+		clReleaseMemObject(diffuse_map_buffer);
+	}
+	clReleaseCommandQueue(commands);
+
+	free(hidden_pixel_buffer);
+
+
+	clReleaseContext(context);
+}

OpenWindow/kernels.h

@@ -0,0 +1,65 @@
+#ifndef _KERNEL_HEADER_
+#define _KERNEL_HEADER_
+#include <Windows.h>
+#include <CL/cl.h>
+#include "geometry.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;
+
+void init_kernels();
+void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices);
+void fragment_shader(
+		cl_int3* faces,
+		int nfaces,
+		float* uv,
+		size_t uv_size,
+		float* uniform_m,
+		float* uniform_mit,
+		float* light_dir,
+		unsigned char* diffuse_map,
+		float* norms,
+		size_t norms_size,
+		unsigned char* spec_map,
+		int* map_size
+);
+
+void destroy_kernels();
+
+
+
+
+
+
+
+
+
+
+
+
+#endif // !_KERNEL_HEADER_

OpenWindow/keyframe.cpp

@@ -1,17 +0,0 @@
-#include "keyframe.h"
-
-Keyframe::Keyframe(float timestamp, std::map<std::string, JointTransform> pose)
-{
-	timestamp_ = timestamp;
-	pose_ = pose;
-}
-
-float Keyframe::gettimestamp()
-{
-	return timestamp_;
-}
-
-std::map<std::string, JointTransform> Keyframe::getJointKeyFrames()
-{
-	return pose_;
-}

OpenWindow/keyframe.h

@@ -1,23 +0,0 @@
-#ifndef HI
-#define HI
-
-// your declarations (and certain types of definitions) here
-
-
-#include <map>
-#include "jointtransform.h"
-
-
-class Keyframe{
-private:
-	float timestamp_;
-	std::map<std::string, JointTransform> pose_;
-
-public:
-	Keyframe() = default;
-	Keyframe(float timestamp,std::map<std::string, JointTransform> pose);
-	float gettimestamp();
-	std::map<std::string, JointTransform> getJointKeyFrames();
-};
-
-#endif

OpenWindow/main.cpp

@@ -2,6 +2,7 @@
 #include "util_window.h"
 #include "renderer.h"
 #include "ctime"
+#include "kernels.h"
 
 const int screen_width = 1000;
 const int screen_height = 1000;
@@ -20,11 +21,22 @@ int WINAPI WinMain(HINSTANCE hInstance, HINSTANCE hPrevInstance, LPSTR lpCmdLine
 
    hwnd = create_window(hInstance);
    ShowCursor(false);
+
+    AllocConsole();
+	freopen("CONIN$", "r",stdin);
+	freopen("CONOUT$", "w",stdout);
+	freopen("CONOUT$", "w",stderr);
+
    ShowWindow(hwnd, nCmdShow);
 
    init_camera();
+   init_kernels();
 
 
+	//camera.ApplyChanges();
+	//render();
+	//Update();
+
    SetTimer(hwnd, NULL, 1000 / TARGET_FRAMERATE, (TIMERPROC)FixedUpdate);
 
    while (GetMessage(&Msg, NULL, 0, 0))
@@ -46,13 +58,13 @@ void HandleMouseMovement() {
 }
 
 bool HandleButtonPressed() {
-	if (GetAsyncKeyState(VK_UP) & 0x8000)
+	if (GetAsyncKeyState(0x57) & 0x8000)
 		camera.move_camera_forward();
-	if (GetAsyncKeyState(VK_DOWN) & 0x8000)
+	if (GetAsyncKeyState(0x53) & 0x8000)
 		camera.move_camera_backward();
-	if (GetAsyncKeyState(VK_RIGHT) & 0x8000)
+	if (GetAsyncKeyState(0x44) & 0x8000)
 		camera.move_camera_right();
-	if (GetAsyncKeyState(VK_LEFT) & 0x8000)
+	if (GetAsyncKeyState(0x41) & 0x8000)
 		camera.move_camera_left();
 	if (GetAsyncKeyState(VK_SPACE) & 0x8000)
 		camera.rise();
@@ -63,11 +75,14 @@ bool HandleButtonPressed() {
 
 void CALLBACK FixedUpdate(HWND hwnd, UINT message, UINT uInt, DWORD dWord)
 {
-	TIME += 0.167;
+	TIME += 0.03333333;
 	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)
@@ -83,6 +98,7 @@ LRESULT CALLBACK WndProc(HWND hwnd, UINT message, WPARAM wParam, LPARAM lParam)
 		break;
 	case WM_CLOSE:
 		DestroyWindow(hwnd);
+		destroy_kernels();
 		break;
 	case WM_DESTROY:
 		PostQuitMessage(0);

OpenWindow/model.h

@@ -7,15 +7,15 @@
 
 class Model {
 private:
+    void load_texture(std::string filename, const char *suffix, TGAImage &img);
+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);
     ~Model();
     int nverts();

OpenWindow/quaternion.cpp

@@ -1,104 +0,0 @@
-#include "quaternion.h"
-#include <math.h>
-
-Quaternion::Quaternion(float x, float y, float z, float w)
-{
-	x_ = x;
-	y_ = y;
-	z_ = z;
-	w_ = w;
-	normalize();
-}
-
-void Quaternion::normalize()
-{
-	float mag = (float)sqrt(w_ * w_ + x_ * x_ + y_ * y_ + z_ * z_);
-	w_ /= mag;
-	x_ /= mag;
-	y_ /= mag;
-	z_ /= mag;
-}
-
-Matrix Quaternion::toRotationMatrix()
-{
-	Matrix matrix;
-	float xy = x_ * y_;
-	float xz = x_ * z_;
-	float xw = x_ * w_;
-	float yz = y_ * z_;
-	float yw = y_ * w_;
-	float zw = z_ * w_;
-	float xSquared = x_ * x_;
-	float ySquared = y_ * y_;
-	float zSquared = z_ * z_;
-	
-	Vec4f c0 = Vec4f(1 - 2 * (ySquared + zSquared), 2 * (xy - zw), 2 * (xz + yw),0);
-	Vec4f c1 = Vec4f(2 * (xy + zw), 1 - 2 * (xSquared + zSquared), 2 * (yz - xw), 0);
-	Vec4f c2 = Vec4f(2 * (xz - yw), 2 * (yz + xw), 1 - 2 * (xSquared + ySquared), 0);
-	Vec4f c3 = Vec4f(0,0,0,1);
-
-	matrix.set_col(0,c0);
-	matrix.set_col(1, c1);
-	matrix.set_col(2, c2);
-	matrix.set_col(3, c3);
-	return matrix;
-}
-
-Quaternion Quaternion::fromMatrix(Matrix matrix)
-{
-	float w_, x_, y_, z_;
-	float diagonal = matrix[0][0] + matrix[1][1] + matrix[2][2];
-	if (diagonal > 0) {
-		float w4 = (float)(sqrt(diagonal + 1.f) * 2.f);
-		w_ = w4 / 4.f;
-		x_ = (matrix[2][1] - matrix[1][2]) / w4;
-		y_ = (matrix[0][2] - matrix[2][0]) / w4;
-		z_ = (matrix[1][0] - matrix[0][1]) / w4;
-	}
-	else if ((matrix[0][0] > matrix[1][1]) && (matrix[0][0] > matrix[2][2])) {
-		float x4 = (float)(sqrt(1.f + matrix[0][0] - matrix[1][1] - matrix[2][2]) * 2.f);
-		w_ = (matrix[2][1] - matrix[1][2]) / x4;
-		x_ = x4 / 4.f;
-		y_ = (matrix[0][1] + matrix[1][0]) / x4;
-		z_ = (matrix[0][2] + matrix[2][0]) / x4;
-	}
-	else if (matrix[1][1] > matrix[2][2]) {
-		float y4 = (float)(sqrt(1.f + matrix[1][1] - matrix[0][0] - matrix[2][2]) * 2.f);
-		w_ = (matrix[0][2] - matrix[2][0]) / y4;
-		x_ = (matrix[0][1] + matrix[1][0]) / y4;
-		y_ = y4 / 4.f;
-		z_ = (matrix[1][2] + matrix[2][1]) / y4;
-	}
-	else {
-		float z4 = (float)(sqrt(1.f + matrix[2][2] - matrix[0][0] - matrix[1][1]) * 2.f);
-		w_ = (matrix[1][0] - matrix[0][1]) / z4;
-		x_ = (matrix[0][2] + matrix[2][0]) / z4;
-		y_ = (matrix[1][2] + matrix[2][1]) / z4;
-		z_ = z4 / 4.f;
-	}
-	return Quaternion(x_, y_, z_, w_);
-}
-
-Quaternion Quaternion::interpolate(Quaternion a, Quaternion b, float blend)
- {
-	Quaternion result(0, 0, 0, 1);
-	float dot = a.w_ * b.w_ + a.x_ * b.x_ + a.y_ * b.y_ + a.z_ * b.z_;
-	float blendI = 1.f - blend;
-	if (dot < 0) {
-		result.w_ = blendI * a.w_ + blend * -b.w_;
-		result.x_ = blendI * a.x_ + blend * -b.x_;
-		result.y_ = blendI * a.y_ + blend * -b.y_;
-		result.z_ = blendI * a.z_ + blend * -b.z_;
-	}
-	else {
-		result.w_ = blendI * a.w_ + blend * b.w_;
-		result.x_ = blendI * a.x_ + blend * b.x_;
-		result.y_ = blendI * a.y_ + blend * b.y_;
-		result.z_ = blendI * a.z_ + blend * b.z_;
-	}
-	result.normalize();
-	return result;
-}
-
-
-

OpenWindow/quaternion.h

@@ -1,17 +0,0 @@
-#pragma once
-#include "geometry.h"
-class Quaternion
-{
-private:
-	float x_, y_, z_, w_;
-
-
-public:
-	Quaternion() = default;
-	Quaternion(float x, float y, float z, float w);
-	void normalize();
-	Matrix toRotationMatrix();
-	static Quaternion fromMatrix(Matrix matrix);
-	static Quaternion interpolate(Quaternion a, Quaternion b, float blend);
-};
-

OpenWindow/renderer.cpp

@@ -3,19 +3,23 @@
 #include "util_window.h"
 #include "camera.h"
 #include "util_renderer.h"
-#include "colladamodel.h"
+#include "CL/cl.h"
+#include "kernels.h"
+#include <ctime>
+
+#pragma comment (lib, "x86_64/opencl.lib")
 
 #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                     0 
+#define NEAR_CLIP_PLANE                    1.f 
-#define FAR_CLIP_PLANE                     10
+#define FAR_CLIP_PLANE                     2000.0f
-#define FOV                                30
+#define FOV                                50
-#define CAMERA_MOVEMENT_SPEED           0.05f
+#define CAMERA_MOVEMENT_SPEED              .7f
 #define DEFAULT_CAMERA_POS     Vec3f(0, 0, 5)
 #define DEFAULT_CAMERA_ROT     Vec3f(0, 0, 0)
-#define LIGHT_INTENSITY                   2.f
+#define LIGHT_INTENSITY                   1.5
 
 const TGAColor white = TGAColor(255, 255, 255, 255);
 const TGAColor red = TGAColor(255, 0, 0, 255);
@@ -26,11 +30,16 @@ Matrix ViewPort = Matrix::identity();
 Matrix ModelView = Matrix::identity();
 Matrix Projection = Matrix::identity();
 
-ColladaModel* model = new ColladaModel("sssssssssssss.dae");
+Model* model = new Model("sakura.obj");
 Camera camera;
 
 Vec3f light_dir = Vec3f(1, 1, 1).normalize();
 
+float* new_verts = (float*)malloc(4 * sizeof(float) * model->nverts());
+cl_int3* faces = (cl_int3*)malloc(3 * sizeof(cl_int3) * model->nfaces());
+
+bool init_flag = false;
+
 void init_camera() {
 	camera.SetPosition(DEFAULT_CAMERA_POS);
 	camera.SetRotation(DEFAULT_CAMERA_ROT);
@@ -56,11 +65,14 @@ struct TextureShader : public IShader {
 	mat<2, 3, float> varying_uv_coords;
 	Matrix uniform_mit;
 	Matrix uniform_m;
+	Matrix z;
 
 	virtual Vec4f vertex(int iface, int nthvert) {
-		varying_uv_coords.set_col(nthvert, model->uv(iface, nthvert));
+		//varying_uv_coords.set_col(nthvert, model->vert(iface, nthvert));
-		Vec4f gl_Vertex = embed<4>(model->vertix(iface, nthvert));
+		Vec4f gl_Vertex = embed<4>(model->verts_[model->faces_[iface][nthvert][0]]);
-		return ViewPort * Projection * ModelView * gl_Vertex;									// transform it to screen coordinates
+		//return ViewPort * Projection * ModelView * gl_Vertex;									// transform it to screen coordinates
+		return z * gl_Vertex;
+		//return Vec4f(0,0,0,0);
 	}
 
 	virtual bool fragment(Vec3f bar, TGAColor &color) {
@@ -73,7 +85,7 @@ struct TextureShader : public IShader {
 		TGAColor c = model->diffuse(uv);
 		color = c;
 		for (int i = 0; i < 3; i++)
-			color[i] = std::fmin(1 + c[i] * (diff_intensity + 0.8 * spec_intensity), 255) * LIGHT_INTENSITY;
+			color[i] = std::fmin(3 + c[i] * ((1 * diff_intensity + 0.1 * spec_intensity)), 255);// *LIGHT_INTENSITY;
 		return false;                         
 	}
 };
@@ -81,37 +93,102 @@ struct TextureShader : public IShader {
 
 void render()
 {
-	{
+	if (!init_flag) {
 		//light_dir = camera.GetForward().normalize() * -1;
+		viewport(0, 0, screen_width, screen_height, FAR_CLIP_PLANE, NEAR_CLIP_PLANE);
+		for (int i = 0; i < model->nfaces(); i++) {
+			for (int j = 0; j < 3; j++) {
+				faces[i * 3 + j].x = model->faces_[i][j][0];
+				faces[i * 3 + j].y = model->faces_[i][j][1];
+				faces[i * 3 + j].z = model->faces_[i][j][2];
+			}
+		}
+		init_flag = true;
 	}
 
 	{
-		viewport(0, 0, screen_width, screen_height, FAR_CLIP_PLANE, NEAR_CLIP_PLANE);
 		Projection = camera.GetProjectionMatrix();
 		ModelView = camera.GetModelViewMatrix();
 	}
 	
-
+	{
-	//Matrix z = ViewPort * Projection * ModelView * model->Transform;
+		//model->rotate(Vec3f(0, 0, 90));
-
+		//model->ApplyTransform();
-	clear_zbuffer();
-	TextureShader shader;
-	shader.uniform_m =   (Projection);
-	shader.uniform_mit = (Projection).invert_transpose();
-
-	#pragma omp parallel for
-	for (int i = 0; i < model->nfaces(); i++) {
-		Vec4f screen_coords[3];
-		bool out = true;
-		#pragma omp parallel for
-		for (int j = 0; j < 3; j++) {
-			screen_coords[j] = shader.vertex(i, j);
-			Vec3f screen3(screen_coords[j]);
-
-			if (screen3.x > 0 && screen3.x < screen_width && screen3.y > 0 && screen3.y < screen_height) out = false;
-		}
-		if(!out)
-			triangle(screen_coords, shader);
 	}
+
+//	float* normal_test = *(float**)((Vec3f*)&model->norms_);
+//
+//	Vec3f smth = Vec3f(
+//		normal_test[3 * faces[0 * 9 + 0 * 3 + 2] + 0],
+//		normal_test[3 * faces[0 * 9 + 0 * 3 + 2] + 1],
+//		normal_test[3 * faces[0 * 9 + 0 * 3 + 2] + 2]
+//	).normalize();
+//	printf("Real: %f, %f, %f\n", model->normal(0, 0)[0], model->normal(0, 0)[1], model->normal(0, 0)[2]);
+//	printf("Please: %f, %f, %f\n", smth[0], smth[1], smth[2]);
+
+	//TextureShader shader;
+	//shader.uniform_m =   (Projection);
+	//shader.uniform_mit = (Projection).invert_transpose();
+
+	Matrix z = ViewPort * Projection * ModelView * model->Transform;
+
+	Matrix uniform_m =   (Projection);
+	Matrix uniform_mit = (Projection).invert_transpose();
+
+	// Vertex Shader: Should be called per model
+	vertex_shader((float*)&z, *(float**)((Vec3f*) &model->verts_), model->nverts(), new_verts);
+
+	// Things needed in the GPU fragment shader
+	// [x] model->faces_
+	// [x] model->nfaces()
+	// [x] new_verts
+	// [x] model->nverts()
+	// [x] screen_width
+	// [x] screen_height
+	// [x] z_buffer
+	// [x] uniform_m
+	// [x] uniform_mit
+	// [x] light direction
+	// [ ] diffuse map
+	// [ ] normal map
+	// [ ] specular map
+	// [ ] pixel_data
+
+	int map_size[] = {model->diffusemap_.get_width(), model->diffusemap_.get_height()};
+
+	fragment_shader(
+		faces,
+		model->nfaces(),
+		*(float**)((Vec2f*)&model->uv_),
+		sizeof(float) * model->uv_.size() * 2,
+		(float*) &uniform_m,
+		(float*) &uniform_mit,
+		(float*) &light_dir,
+		model->diffusemap_.data,
+		*(float**)((Vec3f*)&model->norms_),
+		sizeof(float) * model->norms_.size() * 3,
+		model->specularmap_.data,
+		map_size
+	);
+
+	// Here starts the CPU fragment shader
+	//printf("Here starts the loop\n");
+	//#pragma omp parallel for
+	//for (int i = 0; i < model->nfaces(); i++) {
+	//	Vec4f screen_coords[3];
+	//	bool out = true;
+	//	#pragma omp parallel for
+	//	for (int j = 0; j < 3; j++) {
+	//		screen_coords[j] = ((Vec4f*)new_verts)[model->faces_[i][j][0]];
+	//		Vec3f screen3(screen_coords[j]);
+
+	//		shader.varying_uv_coords.set_col(j, model->uv(i, j));
+	//		if (screen3.x > 0 && screen3.x < screen_width && screen3.y > 0 && screen3.y < screen_height) out = false;
+	//	}
+	//	if(!out)
+	//		triangle(screen_coords, shader);
+	//}
+
+	//printf("that's it\n");
 }
 

OpenWindow/tgaimage.h

@@ -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
     };

OpenWindow/util_renderer.cpp

@@ -1,5 +1,6 @@
 #include "util_renderer.h"
 #include "util_window.h"
+#include "kernels.h"
 
 IShader::~IShader() {}
 
@@ -16,7 +17,7 @@ void viewport(int x, int y, int w, int h, int far_plane, int near_plane) {
 }
 
 int color_to_int(TGAColor col) {
-	return (col[2] << 16) | (col[1] << 8) | col[0];
+	return (col[0] << 24) | (col[1] << 16) | (col[2] << 8) | col[3];
 }
 
 Vec3f barycentric(Vec3f* pts, Vec3f P)
@@ -52,22 +53,38 @@ void triangle( Vec4f* pts, IShader &shader)
 			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);
+			//Vec3f bc_coord = barycentric(pts3, P);
-			float frag_depth = 0;
+			//float frag_depth = 0;
-			for (int i = 0; i < 3; i++)
+			//for (int i = 0; i < 3; i++)
-				frag_depth += pts3[i][2] * bc_coord[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;
+			//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;
+			//TGAColor color;
-			bool discard = shader.fragment(bc_coord, color);
+			//bool discard = shader.fragment(bc_coord, color);
-			if (!discard) {
+			//if (!discard) {
-				z_buffer[P.x + P.y * screen_width] = frag_depth;
+			//	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, color_to_int(color));
-			}
+			    set_pixel(P.x, P.y, 0x00ff00);
+			//}
 		}
 	}
 }

OpenWindow/util_window.cpp

@@ -95,9 +95,9 @@ 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;