Merge pull request #1 from mo7sener/make_it_usable

Make it usable

.gitignore

@@ -67,7 +67,6 @@ StyleCopReport.xml
 *_h.h
 *.ilk
 *.meta
-*.obj
 *.iobj
 *.pch
 *.pdb

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">
@@ -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,15 +127,28 @@
       <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="camera.cpp" />
     <ClCompile Include="geometry.cpp" />
+    <ClCompile Include="improv_gfx.cpp" />
+    <ClCompile Include="kernels.cpp">
+      <AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Release|x64'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+    </ClCompile>
+    <ClCompile Include="kernel_sources.cpp" />
     <ClCompile Include="main.cpp" />
     <ClCompile Include="model.cpp" />
     <ClCompile Include="renderer.cpp" />
@@ -127,6 +159,8 @@
   <ItemGroup>
     <ClInclude Include="geometry.h" />
     <ClInclude Include="camera.h" />
+    <ClInclude Include="improv_gfx.h" />
+    <ClInclude Include="kernels.h" />
     <ClInclude Include="model.h" />
     <ClInclude Include="renderer.h" />
     <ClInclude Include="tgaimage.h" />

OpenWindow/OpenWindow.vcxproj.filters

@@ -39,6 +39,15 @@
     <ClCompile Include="util_renderer.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
+    <ClCompile Include="kernels.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
+    <ClCompile Include="kernel_sources.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
+    <ClCompile Include="improv_gfx.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
   </ItemGroup>
   <ItemGroup>
     <ClInclude Include="util_window.h">
@@ -62,5 +71,11 @@
     <ClInclude Include="util_renderer.h">
       <Filter>Header Files</Filter>
     </ClInclude>
+    <ClInclude Include="kernels.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="improv_gfx.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
   </ItemGroup>
 </Project>

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/improv_gfx.cpp

@@ -0,0 +1,27 @@
+#include "improv_gfx.h"
+
+Model *octopus, *ocean, *sharky;
+
+void Setup() {
+	octopus = new Model("octojet.obj", 1);
+	octopus->scale(Vec3f(4, 4, 4));
+	//octopus->translate(Vec3f(0, -0.15, 0));
+	octopus->ApplyTransform();
+	AddModel(octopus);
+	ocean = new Model("ocean.obj", 1);
+	ocean->scale(Vec3f(1, 1, 1));
+	ocean->ApplyTransform();
+	AddModel(ocean);
+	sharky = new Model("sharky.obj", 0);
+	sharky->scale(Vec3f(0.02, 0.02, 0.02));
+	sharky->translate(Vec3f(0, 0.05, 0));
+	sharky->ApplyTransform();
+	AddModel(sharky);
+}
+
+void Draw() {
+	sharky->translate(Vec3f(sharky->Translation[0][3] + 0.01*cos(TIME), sharky->Translation[1][3], sharky->Translation[2][3] + 0.01*sin(TIME)));
+	sharky->rotate(Vec3f(0, (cos(TIME) + sin(TIME)) * -90, 0));
+	sharky->ApplyTransform();
+
+}

OpenWindow/improv_gfx.h

@@ -0,0 +1,11 @@
+#ifndef __IMPROV_GFX__
+#define __IMPROV_GFX__
+#include "model.h"
+#include "geometry.h"
+#include "renderer.h"
+
+void Setup();
+void Draw();
+void AddModel(Model* model);
+
+#endif

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/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,78 @@
+#include "kernels.h"
+#include "renderer.h"
+#include <debugapi.h>
+#include <ctime>
+#include <climits>
+
+cl_uint numPlatforms;
+cl_uint numDevices;
+
+cl_platform_id platform_id;
+cl_device_id device_id;
+cl_context context;
+
+int* hidden_pixel_buffer;
+cl_mem z_buffer_mem;
+cl_mem pixel_data_buffer;
+cl_mem time_buffer;
+//==================
+
+cl_platform_id* platforms = NULL;
+cl_device_id* devices = NULL;
+
+cl_command_queue commands;
+
+int err;
+
+
+void init_kernels() {
+	err = clGetPlatformIDs(0, NULL, &numPlatforms);
+	platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * numPlatforms);
+	err = clGetPlatformIDs(numPlatforms, platforms, NULL);
+
+	err = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices);
+	devices = (cl_device_id*)malloc(sizeof(cl_device_id)*numDevices);
+	err = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
+
+	context = clCreateContext(NULL, numDevices, devices, NULL, NULL, &err);
+
+	commands = clCreateCommandQueue(context, devices[0], 0, &err);
+
+	hidden_pixel_buffer = new int[screen_width*screen_height];
+	pixel_data_buffer  = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * screen_height * screen_width  , NULL, &err);
+	z_buffer_mem       = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * screen_height * screen_width, NULL, &err);
+	time_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float), NULL, &err);
+}
+
+
+void clear(cl_mem* buffer, size_t size, const int pattern) {
+	clEnqueueFillBuffer(commands, *buffer, &pattern, sizeof(int), 0, size, 0, NULL, NULL);
+}
+
+void new_frame() {
+	clear(&pixel_data_buffer, sizeof(int) * screen_width * screen_height, 0);
+	clear(&z_buffer_mem, sizeof(float) * screen_width * screen_height, 0);
+	clEnqueueWriteBuffer(commands, time_buffer, CL_TRUE, 0, sizeof(float), &TIME, 0, NULL, NULL);
+	clFinish(commands);
+}
+
+void end_frame() {
+	clFinish(commands);
+	clEnqueueReadBuffer(commands, pixel_data_buffer, CL_TRUE, 0, sizeof(int) * screen_height * screen_width, hidden_pixel_buffer, 0, NULL, NULL);
+
+	for (int j = 0; j < screen_height; j++)
+		for (int i = 0; i < screen_width; i++)
+			set_pixel(i, j, hidden_pixel_buffer[j * screen_width + i]);
+}
+
+void destroy_kernels()
+{
+	clReleaseMemObject(z_buffer_mem);
+	clReleaseMemObject(pixel_data_buffer);
+	clReleaseMemObject(time_buffer);
+	clReleaseCommandQueue(commands);
+
+	free(hidden_pixel_buffer);
+
+	clReleaseContext(context);
+}

OpenWindow/kernels.h

@@ -0,0 +1,62 @@
+#ifndef _KERNEL_HEADER_
+#define _KERNEL_HEADER_
+#include <Windows.h>
+#include <CL/cl.h>
+#include "geometry.h"
+#include "util_window.h"
+
+
+extern cl_uint numPlatforms;
+extern cl_platform_id firstPlatformId;
+extern cl_device_id device_id;
+extern cl_context context;
+
+
+
+// Matrix Multiplication Variables
+#define ORDER 4
+#define MATRIX_SIZE 16
+extern int Mdim, Pdim, Ndim;
+extern cl_mem mat_a;
+extern cl_mem mat_b;
+extern cl_mem mat_c;
+
+extern cl_program mat_mul_prog;
+extern cl_kernel mat_mul_kernel;
+extern cl_command_queue commands;
+
+extern int err;
+
+extern const char* matrix_mul_kernel_source;
+// End of: Matrix Multiplication Variables
+
+extern const char* vertex_shader_kernel_source;
+extern const char* fragment_shader_kernel_source;
+extern const char* fluid_vertex_shader_kernel_source;
+
+extern cl_command_queue commands;
+extern cl_platform_id* platforms;
+extern cl_device_id* devices;
+extern cl_context context;
+
+extern cl_mem z_buffer_mem;
+extern cl_mem pixel_data_buffer;
+extern cl_mem time_buffer;
+
+void init_kernels();
+void destroy_kernels();
+void new_frame();
+void end_frame();
+
+
+
+
+
+
+
+
+
+
+
+
+#endif // !_KERNEL_HEADER_

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 += 1.f / (float)TARGET_FRAMERATE;
 	HandleButtonPressed();
 	camera.ApplyChanges();
+	clock_t start = clock();
 	render();
 	Update();
+	clock_t end = clock();
+	printf("FPS: %f\n", 1/((float)(end-start)/CLOCKS_PER_SEC));
 }
 
 LRESULT CALLBACK WndProc(HWND hwnd, UINT message, WPARAM wParam, LPARAM lParam)
@@ -83,6 +98,8 @@ LRESULT CALLBACK WndProc(HWND hwnd, UINT message, WPARAM wParam, LPARAM lParam)
 		break;
 	case WM_CLOSE:
 		DestroyWindow(hwnd);
+		free_renderer();
+		destroy_kernels();
 		break;
 	case WM_DESTROY:
 		PostQuitMessage(0);

OpenWindow/model.cpp

@@ -6,7 +6,8 @@
 #define PI 3.14159265358979323846
 #define DEG2RAD PI/180
 
-Model::Model(const char *filename) : verts_(), faces_(), norms_(), uv_(), diffusemap_(), normalmap_(), specularmap_() {
+Model::Model(const char *filename, int watery) : verts_(), faces_(), norms_(), uv_(), diffusemap_(), normalmap_(), specularmap_() {
+	fluid = watery;
     std::ifstream in;
 	Transform = Matrix::identity();
 	Rotation = Matrix::identity();
@@ -45,13 +46,13 @@ Model::Model(const char *filename) : verts_(), faces_(), norms_(), uv_(), diffus
             faces_.push_back(f);
         }
     }
-    std::cerr << "# v# " << verts_.size() << " f# "  << faces_.size() << " vt# " << uv_.size() << " vn# " << norms_.size() << std::endl;
     load_texture(filename, "_diffuse.tga", diffusemap_);
-    load_texture(filename, "_nm_tangent.tga",      normalmap_);
+	init_kernels();
-    load_texture(filename, "_spec.tga",    specularmap_);
 }
 
-Model::~Model() {}
+Model::~Model() {
+	release_kernels();
+}
 
 void Model::ApplyTransform() {
 	Transform = Translation * Scale * Rotation;
@@ -109,7 +110,7 @@ void Model::load_texture(std::string filename, const char *suffix, TGAImage &img
     if (dot!=std::string::npos) {
         texfile = texfile.substr(0,dot) + std::string(suffix);
         std::cerr << "texture file " << texfile << " loading " << (img.read_tga_file(texfile.c_str()) ? "ok" : "failed") << std::endl;
-        img.flip_vertically();
+        //img.flip_vertically();
     }
 }
 
@@ -141,3 +142,275 @@ Vec3f Model::normal(int iface, int nthvert) {
     return norms_[idx].normalize();
 }
 
+void Model::init_kernels() {
+	if(fluid)
+		vertex_shader_prog = clCreateProgramWithSource(context, 1, (const char **)&fluid_vertex_shader_kernel_source, NULL, &err);
+	else
+		vertex_shader_prog = clCreateProgramWithSource(context, 1, (const char **)&vertex_shader_kernel_source, NULL, &err);
+	err = clBuildProgram(vertex_shader_prog, 1, devices, NULL, NULL, NULL);
+	vertex_shader_kernel = clCreateKernel(vertex_shader_prog, "vertex_shader", &err);
+	vertex_shader_matz = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * 16, NULL, &err);
+	vertex_shader_vertices = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * nverts() * 3, NULL, &err);
+	new_vertices_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float4) * nverts(), NULL, &err);
+
+	clSetKernelArg(vertex_shader_kernel, 0, sizeof(cl_mem), &vertex_shader_matz);
+	clSetKernelArg(vertex_shader_kernel, 1, sizeof(cl_mem), &vertex_shader_vertices);
+	clSetKernelArg(vertex_shader_kernel, 2, sizeof(cl_mem), &new_vertices_mem);
+	clSetKernelArg(vertex_shader_kernel, 3, sizeof(cl_mem), &time_buffer);
+	clEnqueueWriteBuffer(commands, vertex_shader_vertices, CL_TRUE, 0, sizeof(float) * nverts() * 3, *(float**)((Vec3f*) &verts_), 0, NULL, NULL);
+
+	int map_size[] = { diffusemap_.get_width(), diffusemap_.get_height() };
+	faces = (cl_int3*)malloc(3 * sizeof(cl_int3) * nfaces());
+	for (int i = 0; i < nfaces(); i++) {
+		for (int j = 0; j < 3; j++) {
+			faces[i * 3 + j].x = faces_[i][j][0];
+			faces[i * 3 + j].y = faces_[i][j][1];
+			faces[i * 3 + j].z = faces_[i][j][2];
+		}
+	}
+
+	fragment_shader_prog = clCreateProgramWithSource(context, 1, (const char**)&fragment_shader_kernel_source, NULL, &err);
+	err = clBuildProgram(fragment_shader_prog, 1, devices, NULL, NULL, NULL);
+	fragment_shader_kernel = clCreateKernel(fragment_shader_prog, "fragment_shader", &err);
+
+	fragment_shader_faces        = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(cl_int3) * 3 * nfaces()               , NULL, &err);
+	fragment_shader_screen_width = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int)                                  , NULL, &err);
+	fragment_shader_uv           = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * uv_.size() * 2               , NULL, &err);
+	fragment_shader_map_size     = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int) * 2                              , NULL, &err);
+	fragment_shader_norms        = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * norms_.size() * 3            , NULL, &err);
+	fragment_shader_light_dir    = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * 3                            , NULL, &err);
+	fragment_shader_diffuse_map  = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(char) * 3 * map_size[0] * map_size[1] , NULL, &err);
+
+	clSetKernelArg(fragment_shader_kernel, 0, sizeof(cl_mem), &fragment_shader_faces);
+	clSetKernelArg(fragment_shader_kernel, 1, sizeof(cl_mem), &new_vertices_mem);
+	clSetKernelArg(fragment_shader_kernel, 2, sizeof(cl_mem), &pixel_data_buffer);
+	clSetKernelArg(fragment_shader_kernel, 3, sizeof(cl_mem), &fragment_shader_screen_width);
+	clSetKernelArg(fragment_shader_kernel, 4, sizeof(cl_mem), &z_buffer_mem);
+	clSetKernelArg(fragment_shader_kernel, 5, sizeof(cl_mem), &fragment_shader_uv);
+	clSetKernelArg(fragment_shader_kernel, 6, sizeof(cl_mem), &fragment_shader_map_size);
+	clSetKernelArg(fragment_shader_kernel, 7, sizeof(cl_mem), &fragment_shader_light_dir);
+	clSetKernelArg(fragment_shader_kernel, 8, sizeof(cl_mem), &fragment_shader_norms);
+	clSetKernelArg(fragment_shader_kernel, 9, sizeof(cl_mem), &fragment_shader_diffuse_map);
+
+	clEnqueueWriteBuffer(commands, fragment_shader_screen_width, CL_FALSE, 0, sizeof(int)                                 , &screen_width              , 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_faces       , CL_FALSE, 0, sizeof(cl_int3) * 3 * nfaces()              , faces                      , 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_uv          , CL_FALSE, 0, sizeof(float) * 2 * uv_.size()              , *(float**)((Vec2f*)&uv_)   , 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_map_size    , CL_FALSE, 0, sizeof(int) * 2                             , map_size                   , 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_norms       , CL_FALSE, 0, sizeof(float) * norms_.size() * 3           , *(float**)((Vec3f*)&norms_), 0, NULL, NULL);
+	clEnqueueWriteBuffer(commands, fragment_shader_diffuse_map , CL_TRUE, 0, sizeof(char) * 3 * map_size[0] * map_size[1], diffusemap_.data           , 0, NULL, NULL);
+}
+
+void Model::vertex(float* z) {
+	size_t vertex_shader_global[] = { nverts() * 4 };
+	size_t vertex_shader_local[] = { 4 };
+
+	clEnqueueWriteBuffer(commands, vertex_shader_matz, CL_TRUE, 0, sizeof(float) * 16, z, 0, NULL, NULL);
+
+	clEnqueueNDRangeKernel(commands, vertex_shader_kernel, 1, NULL, vertex_shader_global, vertex_shader_local, 0, NULL, NULL);
+}
+void Model::fragment(float* light_dir) {
+	clEnqueueWriteBuffer(commands, fragment_shader_light_dir, CL_FALSE, 0, sizeof(float) * 3, light_dir, 0, NULL, NULL);
+	size_t fragment_shader_global[] = { nfaces() * 256 };
+	size_t fragment_shader_local[] = { 256 };
+
+	clEnqueueNDRangeKernel(commands, fragment_shader_kernel, 1, NULL, fragment_shader_global, fragment_shader_local, 0, NULL, NULL);
+
+}
+void Model::render(Matrix* z, float* light_dir) {
+	Matrix m = *z * Transform;
+	vertex((float*)&m);
+	clFinish(commands);
+	fragment(light_dir);
+}
+void Model::release_kernels() {
+
+	clReleaseProgram(vertex_shader_prog);
+	clReleaseKernel(vertex_shader_kernel);
+	clReleaseMemObject(vertex_shader_matz);
+	clReleaseMemObject(vertex_shader_vertices);
+	clReleaseMemObject(new_vertices_mem);
+
+	clReleaseProgram(fragment_shader_prog);
+	clReleaseKernel(fragment_shader_kernel);
+	clReleaseMemObject(fragment_shader_faces);
+	clReleaseMemObject(fragment_shader_screen_width);
+	clReleaseMemObject(fragment_shader_uv);
+	clReleaseMemObject(fragment_shader_map_size);
+	clReleaseMemObject(fragment_shader_norms);
+	clReleaseMemObject(fragment_shader_light_dir);
+	clReleaseMemObject(fragment_shader_diffuse_map);
+
+	free(faces);
+}
+
+
+const char* fluid_vertex_shader_kernel_source =
+"__kernel                                                                                                                    \n"
+"void vertex_shader( __global float* m,                                                                                      \n"
+"                    __global float* VertexBuffer,                                                                           \n"
+"                    __global float* NewVertexBuffer,                                                                        \n"
+"                    __global float* TIME           )                                                                        \n"
+"{                                                                                                                           \n"
+"    int local_index = get_local_id(0);                                                                                      \n"
+"    int global_index = get_group_id(0);                                                                                     \n"
+"    float AMPLITUDE = 5;                                                                                     \n"
+"    float FREQ = 2.0f;                                                                                     \n"
+"                                                                                                                            \n"
+"    float offset = 0.f;                                                                                                     \n"
+"    if(local_index == 1) {                                                                                                  \n"
+"    offset = (sin(TIME[0]*(fmod(10-VertexBuffer[3*global_index],5.f) * FREQ)) + sin(TIME[0]*(fmod(10-VertexBuffer[3*global_index+2], 5.f))*FREQ)) * AMPLITUDE;                                                                                                \n"
+"    }                                                                                                                       \n"
+"                                                                                                                            \n"
+"    NewVertexBuffer[4*global_index+local_index] =                                                                           \n"
+"						  m[local_index*4]*VertexBuffer[3*global_index]                                                      \n"
+"	                    + m[local_index*4 + 1]*VertexBuffer[3*global_index+1]                                                \n"
+"	                    + m[local_index*4 + 2]*VertexBuffer[3*global_index+2]                                                \n"
+"	                    + m[local_index*4 + 3]                                                                               \n"
+"	                    + offset;                                                                                            \n"
+"}                                                                                                                           \n";
+
+
+const char* vertex_shader_kernel_source =
+"__kernel                                                                                                                    \n"
+"void vertex_shader( __global float* m,                                                                                      \n"
+"                    __global float* VertexBuffer,                                                                           \n"
+"                    __global float* NewVertexBuffer,                                                                        \n"
+"                    __global float* TIME           )                                                                        \n"
+"{                                                                                                                           \n"
+"    int local_index = get_local_id(0);                                                                                      \n"
+"    int global_index = get_group_id(0);                                                                                     \n"
+"                                                                                                                            \n"
+"    NewVertexBuffer[4*global_index+local_index] =                                                                           \n"
+"						  m[local_index*4]*VertexBuffer[3*global_index]                                                      \n"
+"	                    + m[local_index*4 + 1]*VertexBuffer[3*global_index+1]                                                \n"
+"	                    + m[local_index*4 + 2]*VertexBuffer[3*global_index+2]                                                \n"
+"	                    + m[local_index*4 + 3];                                                                              \n"
+"}                                                                                                                           \n";
+
+const char* fragment_shader_kernel_source =
+"float3 barycentric(float3* pts, float3 P)                                                                                                                                                                    \n"
+"{                                                                                                                                                                                                          \n"
+"    float3 u = cross(                                                                                                                                                                                      \n"
+"        (float3){pts[0][2] - pts[0][0], pts[0][1] - pts[0][0], pts[0][0] - P[0]}, // AC_x, AB_x, distance_x                                                                                                \n"
+"        (float3){pts[1][2] - pts[1][0], pts[1][1] - pts[1][0], pts[1][0] - P[1]}  // AC_y, AB_y, distance_y                                                                                                \n"
+"    );                                                                                                                                                                                                     \n"
+"	 if (fabs(u[2]) < 1) return (float3){-1, 1, 1};                                                                                                                                                         \n"
+"	 return (float3){1.f - (u[0] + u[1]) / u[2], u[1] / u[2], u[0] / u[2]};                                                                                                                                 \n"
+"}                                                                                                                                                                                                          \n"
+"                                                                                                                                                                                                    \n"
+"__kernel void fragment_shader (                                                                                                                                                                \n"
+"    __global int3* faces,                                                                                                                                                                        \n"
+"    __global float* vertices,                                                                                                                                                                             \n"
+"    __global int* pixels,                                                                                                                                                                  \n"
+"    __global int* screen_width,                                                                                                                                                                  \n"
+"    __global float* z_buffer,                                                                                                                                                                         \n"
+"    __global float* uv_buffer,                                                                                                                                                                        \n"
+"    __global int* map_size,                                                                                                                                                                       \n"
+"    __global float* light_dir,                                                                                                                                                                           \n"
+"    __global float* norms_buff,                                                                                                                                                                           \n"
+"    __global uchar* diffuse_map                                                                                                                                                                              \n"
+") {                                                                                                                                                                                                        \n"
+"    int GROUP_ID = get_group_id(0);                                                                                                                                                                          \n"
+"    int GROUP_SIZE = get_local_size(0);                                                                                                                                                                          \n"
+"    int LOCAL_ID = get_local_id(0);                                                                                                                                                                          \n"
+"                                                                                                                                                                                                           \n"
+"    bool out = true;                                                                                                                                                                                       \n"
+"    float3 vertices3[3];                                                                                                                                                                                   \n"
+"    float2 uv_coords[3];                                 \n"
+"    float3 norms[3];                                                                                                                                                                                   \n"
+"                                                                                                                                                                                                           \n"
+"    for(int i = 0; i < 3; i++) {                                                                                                                                                                           \n"
+"        float4 vertex;// =  vertices[faces[GROUP_ID * 3 + i ][0]];                                                                                                                                                \n"
+"        for(int j = 0; j < 4; j ++) {                                                                                                                                                                                                   \n"
+"            vertex[j] = vertices[4 * faces[GROUP_ID*3 + i][0] + j];                                                                                                                                          \n"
+"        }                                                                                                                                                                                                  \n"
+"                                                                                                                                                                                                           \n"
+"        for(int j = 0; j < 2; j++) {                                                                                                                                                                       \n"
+"            uv_coords[i][j]  = uv_buffer [2 * faces[GROUP_ID * 3 + i][1] + j];                                                                                                                                                \n"
+"        }                                                                                                                                                                                                  \n"
+"                                                                                                                                                                                                           \n"
+"        for( int j = 0; j < 3; j++ ) {																		\n"
+"            vertices3[j][i] = (vertex[j]/vertex[3]);                                                                                                                                                       \n"
+"		     norms[j][i] = norms_buff[3 * (faces[GROUP_ID * 3 + i][2]) + j];                                                                                                                                          \n"
+"        }                                                                                                                                                                                                  \n"
+"                                                                                                                                                                                                           \n"
+"        if (   vertices3[0][i] > 0 && vertices3[0][i] < *screen_width                                                                                                                                      \n"
+"            && vertices3[1][i] > 0 && vertices3[1][i] < *screen_width   )                                                                                                                                  \n"
+"                out = false;                                                                                                                                                                               \n"
+"    }                                                                                                                                                                                                      \n"
+"                                                                                                                                                                                                           \n"
+"    if(out) return;                                                                                                                                                                                        \n"
+"                                                                                                                                                                                                           \n"
+"    //if(vertices3[1][0] == vertices3[1][1] && vertices3[1][2] == vertices3[1][1]) return;                                                                                                                   \n"
+"                                                                                                                                                                                                           \n"
+"    int2 bounding_box_min = (int2) { *screen_width - 1, *screen_width - 1 };                                                                                                                                    \n"
+"    int2 bounding_box_max = (int2) { 0, 0 };                                                                                                                                                                    \n"
+"    int2 clamper = (int2) { *screen_width - 1, *screen_width - 1 };                                                                                                                                               \n"
+"                                                                                                                                                                                                           \n"
+"    for(int i = 0; i < 3; i++) {                                                                                                                                                                           \n"
+"        for(int j = 0; j < 2; j++) {                                                                                                                                                                       \n"
+"            bounding_box_min[j] = max(0, min(bounding_box_min[j], (int)vertices3[j][i]));                                                                                                                  \n"
+"            bounding_box_max[j] = min(clamper[j], max(bounding_box_max[j], (int)vertices3[j][i]));                                                                                                           \n"
+"        }                                                                                                                                                                                                  \n"
+"    }                                                                                                                                                                                                      \n"
+"                                                                                                                                                                                                           \n"
+"    if(bounding_box_min[0] > *screen_width || bounding_box_max[0] < 0 || bounding_box_min[1] > *screen_width || bounding_box_max[1] < 0) return;                                                           \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"    int X_PER_ITEM = (int)(ceil((float)(bounding_box_max[0] - bounding_box_min[0]) / 16.f));                                                                                                                                                                                            \n"
+"    int Y_PER_ITEM = (int)(ceil((float)(bounding_box_max[1] - bounding_box_min[1]) / 16.f));                                                                                                                                                                                            \n"
+"    float STARTING_X = bounding_box_min[0] + X_PER_ITEM * (LOCAL_ID % 16);                                                                                                                                                                                            \n"
+"    float ENDING_X = STARTING_X + X_PER_ITEM;                                                                                                                                                                                            \n"
+"    float STARTING_Y = bounding_box_min[1] + Y_PER_ITEM * (LOCAL_ID / 16);                                                                                                                                                                                            \n"
+"    float ENDING_Y = STARTING_Y + Y_PER_ITEM;                                                                                                                                                                                            \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"    float3 point;                                                                                                                                                                                            \n"
+"    for(point[0] = STARTING_X; point[0] <= ENDING_X; point[0]++) {                                                                                                                     \n"
+"        for(point[1] = STARTING_Y; point[1] <= ENDING_Y; point[1]++) {                                                                                                                 \n"
+"           if(point[1] >= *screen_width || point[1] >= *screen_width) break;                                                                                                                                                                                                \n"
+"			float3 bc_coord = barycentric(vertices3, point);                                                                                                                                                \n"
+"           if (bc_coord[0] < 0 || bc_coord[1] < 0 || bc_coord[2] < 0) continue;                                                                                                                                                                                               \n"
+"                                                                                                                                                                                                           \n"
+"           float2 uv_vec = (float2){0, 0};                                                                                                                                                                                  \n"
+"           float3 normal = (float3){0, 0, 0};                                                                                                                                                                \n"
+"                                                                                                                                                                                                           \n"
+"           point[2] = dot(vertices3[2], bc_coord);                                                                                                                                                                                   \n"
+"			if (z_buffer[(int)(point[0] + point[1] * *screen_width)] > point[2]) {                                                                        \n"
+"               continue;                                                                                                                                                                                   \n"
+"           }                                                                                                                                                                                               \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"			for (int i = 0; i < 3; i++) {                                                                                                                                                                   \n"
+"               uv_vec[0]  += uv_coords[i][0] * bc_coord[i];                                                                                                                                                 \n"
+"               uv_vec[1]  += uv_coords[i][1] * bc_coord[i];                                                                                                                                                 \n"
+"               normal[i]  = dot(norms[i], bc_coord);                                                                                                                                                 \n"
+"           }                                                                                                                                                                                               \n"
+"           int2 uv_point = (int2) { (int)(uv_vec[0] * map_size[0]), (int)(uv_vec[1] * map_size[1]) };                                                                                                                    \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"           int col_index = 3 * (uv_point[0] + uv_point[1] * map_size[0]);                                                                                                                                       \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"           float3 normalized_norm = normalize(normal);                                                                                                                                                                                                \n"
+"                                                                                                                                                                                                           \n"
+"                                                                                                                                                                                                           \n"
+"           float intensity = clamp((dot(normalized_norm , (float3){light_dir[0], light_dir[1], light_dir[2]})), 0.f, 1.f) + 0.2;                                                                                                                                                                                                \n"
+"                                                                                                                                                                                                           \n"
+"           int color = 0;                                                                                                                                                                                  \n"
+"           color |= ((int)fmin((float)(diffuse_map[col_index + 0]) * intensity, (float) 0xff)) << 16;                                                                                                                                                               \n"
+"           color |= ((int)fmin((float)(diffuse_map[col_index + 1]) * intensity, (float) 0xff)) <<  8;                                                                                                                                                               \n"
+"           color |= ((int)fmin((float)(diffuse_map[col_index + 2]) * intensity, (float) 0xff)) <<  0;                                                                                                                                                               \n"
+"                                                                                                                                                                                                          \n"
+"           //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 16;                                                                                                                                                               \n"
+"           //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) <<  8;                                                                                                                                                               \n"
+"           //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) <<  0;                                                                                                                                                               \n"
+"                                                                                                                                                                                                          \n"
+"                                                                                                                                                                                                          \n"
+"           z_buffer[(int) (point[0] + point[1] * *screen_width)] = point[2];                                                                                                                                 \n"
+"			pixels  [(int) (point[0] + point[1] * *screen_width)] = color;// & 0x00ffffff;                                                                                                                                     \n"
+"        }                                                                                                                                                                                                  \n"
+"    }                                                                                                                                                                                                      \n"
+"}                                                                                                                                                                                                          \n";

OpenWindow/model.h

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

OpenWindow/renderer.cpp

@@ -3,32 +3,34 @@
 #include "util_window.h"
 #include "camera.h"
 #include "util_renderer.h"
+#include "CL/cl.h"
+#include "kernels.h"
+#include "improv_gfx.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              .1f
 #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);
-const TGAColor green = TGAColor(0, 255, 0, 255);
-const TGAColor blue = TGAColor(0, 0, 255, 255);
 
 Matrix ViewPort = Matrix::identity();
 Matrix ModelView = Matrix::identity();
 Matrix Projection = Matrix::identity();
 
-Model* model = new Model("african_head.obj");
+Model* model, *model2;
 Camera camera;
 
-Vec3f light_dir = Vec3f(1, 1, 1).normalize();
+Vec3f light_dir = Vec3f(0, 1, 1).normalize();
+bool init_flag = false;
 
 void init_camera() {
 	camera.SetPosition(DEFAULT_CAMERA_POS);
@@ -44,73 +46,38 @@ void init_camera() {
 	camera.ApplyChanges();
 }
 
-void clear_zbuffer()
-{
-	for (int i = 0; i < screen_width * screen_height; i++)
-		z_buffer[i] = 0;
-}
 
-
+std::vector<Model*> models_in_scene;
-struct TextureShader : public IShader {
-	mat<2, 3, float> varying_uv_coords;
-	Matrix uniform_mit;
-	Matrix uniform_m;
-
-	virtual Vec4f vertex(int iface, int nthvert) {
-		varying_uv_coords.set_col(nthvert, model->uv(iface, nthvert));
-		Vec4f gl_Vertex = embed<4>(model->vert(iface, nthvert));
-		return ViewPort * Projection * ModelView * gl_Vertex;									// transform it to screen coordinates
-	}
-
-	virtual bool fragment(Vec3f bar, TGAColor &color) {
-		Vec2f uv = varying_uv_coords * bar;
-		Vec3f normal = Vec3f(uniform_mit * Vec4f(model->normal(uv))).normalize();
-		Vec3f light =  Vec3f(uniform_m * Vec4f(light_dir)).normalize();
-		Vec3f reflection = (normal * (normal*light*2.f) - light).normalize();
-		float spec_intensity = pow(std::fmax(reflection.z, 0.f), model->specular(uv));
-		float diff_intensity = std::fmax(0.f, (normal*light));
-		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;
-		return false;                         
-	}
-};
 
 
 void render()
 {
-	{
+	if (!init_flag) {
-		//light_dir = camera.GetForward().normalize() * -1;
+		viewport(0, 0, screen_width, screen_height, FAR_CLIP_PLANE, NEAR_CLIP_PLANE);
+		Setup();
+		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;
+	Matrix z = ViewPort * Projection * ModelView;
 
-	clear_zbuffer();
+	new_frame();
-	TextureShader shader;
+	Draw();
-	shader.uniform_m =   (Projection);
+	for(Model* model : models_in_scene)
-	shader.uniform_mit = (Projection).invert_transpose();
+		model->render(&z, (float*)&light_dir);
+	end_frame();
 
-	#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);
-	}
 }
 
+void AddModel(Model* model) {
+	models_in_scene.push_back(model);
+}
+
+void free_renderer() {
+	for(Model* model : models_in_scene)
+		delete model;
+}

OpenWindow/renderer.h

@@ -8,7 +8,9 @@ extern float* z_buffer;
 extern Camera camera;
 extern float TIME;
 
+
 void init_camera();
 void render();
+void free_renderer();
 int color_to_int(TGAColor col);
 #endif

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;

README.md

@@ -1,2 +1,2 @@
-# OhShitgl
+# ImprovGFX
-A graphics library built from scratch.
+A well written documentation for the project.