Compare commits
4 Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 Robear Selwans
|
88091af7a5 | ||
|
Robear Selwans
|
2231dc3992 | ||
 mo7sen
|
9a2f6129bc | ||
|
mo7sen
|
5dd2de5d77 |
@@ -67,7 +67,6 @@ StyleCopReport.xml
|
|||||||
*_h.h
|
*_h.h
|
||||||
*.ilk
|
*.ilk
|
||||||
*.meta
|
*.meta
|
||||||
*.obj
|
|
||||||
*.iobj
|
*.iobj
|
||||||
*.pch
|
*.pch
|
||||||
*.pdb
|
*.pdb
|
||||||
|
|||||||
@@ -141,6 +141,7 @@
|
|||||||
<ItemGroup>
|
<ItemGroup>
|
||||||
<ClCompile Include="camera.cpp" />
|
<ClCompile Include="camera.cpp" />
|
||||||
<ClCompile Include="geometry.cpp" />
|
<ClCompile Include="geometry.cpp" />
|
||||||
|
<ClCompile Include="improv_gfx.cpp" />
|
||||||
<ClCompile Include="kernels.cpp">
|
<ClCompile Include="kernels.cpp">
|
||||||
<AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
|
<AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
|
||||||
<AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
|
<AdditionalIncludeDirectories Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
|
||||||
@@ -158,6 +159,7 @@
|
|||||||
<ItemGroup>
|
<ItemGroup>
|
||||||
<ClInclude Include="geometry.h" />
|
<ClInclude Include="geometry.h" />
|
||||||
<ClInclude Include="camera.h" />
|
<ClInclude Include="camera.h" />
|
||||||
|
<ClInclude Include="improv_gfx.h" />
|
||||||
<ClInclude Include="kernels.h" />
|
<ClInclude Include="kernels.h" />
|
||||||
<ClInclude Include="model.h" />
|
<ClInclude Include="model.h" />
|
||||||
<ClInclude Include="renderer.h" />
|
<ClInclude Include="renderer.h" />
|
||||||
|
|||||||
@@ -45,6 +45,9 @@
|
|||||||
<ClCompile Include="kernel_sources.cpp">
|
<ClCompile Include="kernel_sources.cpp">
|
||||||
<Filter>Source Files</Filter>
|
<Filter>Source Files</Filter>
|
||||||
</ClCompile>
|
</ClCompile>
|
||||||
|
<ClCompile Include="improv_gfx.cpp">
|
||||||
|
<Filter>Source Files</Filter>
|
||||||
|
</ClCompile>
|
||||||
</ItemGroup>
|
</ItemGroup>
|
||||||
<ItemGroup>
|
<ItemGroup>
|
||||||
<ClInclude Include="util_window.h">
|
<ClInclude Include="util_window.h">
|
||||||
@@ -71,5 +74,8 @@
|
|||||||
<ClInclude Include="kernels.h">
|
<ClInclude Include="kernels.h">
|
||||||
<Filter>Header Files</Filter>
|
<Filter>Header Files</Filter>
|
||||||
</ClInclude>
|
</ClInclude>
|
||||||
|
<ClInclude Include="improv_gfx.h">
|
||||||
|
<Filter>Header Files</Filter>
|
||||||
|
</ClInclude>
|
||||||
</ItemGroup>
|
</ItemGroup>
|
||||||
</Project>
|
</Project>
|
||||||
File diff suppressed because it is too large
Load Diff
Binary file not shown.
Binary file not shown.
@@ -0,0 +1,27 @@
|
|||||||
|
#include "improv_gfx.h"
|
||||||
|
|
||||||
|
Model *octopus, *ocean, *sharky;
|
||||||
|
|
||||||
|
void Setup() {
|
||||||
|
octopus = new Model("octojet.obj", 1);
|
||||||
|
octopus->scale(Vec3f(4, 4, 4));
|
||||||
|
//octopus->translate(Vec3f(0, -0.15, 0));
|
||||||
|
octopus->ApplyTransform();
|
||||||
|
AddModel(octopus);
|
||||||
|
ocean = new Model("ocean.obj", 1);
|
||||||
|
ocean->scale(Vec3f(1, 1, 1));
|
||||||
|
ocean->ApplyTransform();
|
||||||
|
AddModel(ocean);
|
||||||
|
sharky = new Model("sharky.obj", 0);
|
||||||
|
sharky->scale(Vec3f(0.02, 0.02, 0.02));
|
||||||
|
sharky->translate(Vec3f(0, 0.05, 0));
|
||||||
|
sharky->ApplyTransform();
|
||||||
|
AddModel(sharky);
|
||||||
|
}
|
||||||
|
|
||||||
|
void Draw() {
|
||||||
|
sharky->translate(Vec3f(sharky->Translation[0][3] + 0.01*cos(TIME), sharky->Translation[1][3], sharky->Translation[2][3] + 0.01*sin(TIME)));
|
||||||
|
sharky->rotate(Vec3f(0, (cos(TIME) + sin(TIME)) * -90, 0));
|
||||||
|
sharky->ApplyTransform();
|
||||||
|
|
||||||
|
}
|
||||||
@@ -0,0 +1,11 @@
|
|||||||
|
#ifndef __IMPROV_GFX__
|
||||||
|
#define __IMPROV_GFX__
|
||||||
|
#include "model.h"
|
||||||
|
#include "geometry.h"
|
||||||
|
#include "renderer.h"
|
||||||
|
|
||||||
|
void Setup();
|
||||||
|
void Draw();
|
||||||
|
void AddModel(Model* model);
|
||||||
|
|
||||||
|
#endif
|
||||||
+145
-145
@@ -1,145 +1,145 @@
|
|||||||
#include "kernels.h"
|
//#include "kernels.h"
|
||||||
|
//
|
||||||
const char* vertex_shader_kernel_source =
|
//const char* vertex_shader_kernel_source =
|
||||||
"__kernel \n"
|
//"__kernel \n"
|
||||||
"void vertex_shader( __global float* m, \n"
|
//"void vertex_shader( __global float* m, \n"
|
||||||
" __global float* VertexBuffer, \n"
|
//" __global float* VertexBuffer, \n"
|
||||||
" __global float* NewVertexBuffer) \n"
|
//" __global float* NewVertexBuffer) \n"
|
||||||
"{ \n"
|
//"{ \n"
|
||||||
" int local_index = get_local_id(0); \n"
|
//" int local_index = get_local_id(0); \n"
|
||||||
" int global_index = get_group_id(0); \n"
|
//" int global_index = get_group_id(0); \n"
|
||||||
" NewVertexBuffer[4*global_index+local_index] = \n"
|
//" NewVertexBuffer[4*global_index+local_index] = \n"
|
||||||
" m[local_index*4]*VertexBuffer[3*global_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 + 1]*VertexBuffer[3*global_index+1] \n"
|
||||||
" + m[local_index*4 + 2]*VertexBuffer[3*global_index+2] \n"
|
//" + m[local_index*4 + 2]*VertexBuffer[3*global_index+2] \n"
|
||||||
" + m[local_index*4 + 3]; \n"
|
//" + m[local_index*4 + 3]; \n"
|
||||||
"} \n";
|
//"} \n";
|
||||||
|
//
|
||||||
const char* fragment_shader_kernel_source =
|
//const char* fragment_shader_kernel_source =
|
||||||
"float3 barycentric(float3* pts, float3 P) \n"
|
//"float3 barycentric(float3* pts, float3 P) \n"
|
||||||
"{ \n"
|
//"{ \n"
|
||||||
" float3 u = cross( \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[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"
|
//" (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"
|
//" ); \n"
|
||||||
" if (fabs(u[2]) < 1) return (float3){-1, 1, 1}; \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"
|
//" return (float3){1.f - (u[0] + u[1]) / u[2], u[1] / u[2], u[0] / u[2]}; \n"
|
||||||
"} \n"
|
//"} \n"
|
||||||
" \n"
|
//" \n"
|
||||||
"__kernel void fragment_shader ( \n"
|
//"__kernel void fragment_shader ( \n"
|
||||||
" __global int3* faces, \n"
|
//" __global int3* faces, \n"
|
||||||
" __global float* vertices, \n"
|
//" __global float* vertices, \n"
|
||||||
" __global int* pixels, \n"
|
//" __global int* pixels, \n"
|
||||||
" __global int* screen_width, \n"
|
//" __global int* screen_width, \n"
|
||||||
" __global float* z_buffer, \n"
|
//" __global float* z_buffer, \n"
|
||||||
" __global int* nfaces, \n"
|
//" __global int* nfaces, \n"
|
||||||
" __global float* uv_buffer, \n"
|
//" __global float* uv_buffer, \n"
|
||||||
" __global int* map_size, \n"
|
//" __global int* map_size, \n"
|
||||||
" __global float* light_dir, \n"
|
//" __global float* light_dir, \n"
|
||||||
" __global float* norms_buff, \n"
|
//" __global float* norms_buff, \n"
|
||||||
" __global uchar* diffuse_map \n"
|
//" __global uchar* diffuse_map \n"
|
||||||
") { \n"
|
//") { \n"
|
||||||
" int GROUP_ID = get_group_id(0); \n"
|
//" int GROUP_ID = get_group_id(0); \n"
|
||||||
" int GROUP_SIZE = get_local_size(0); \n"
|
//" int GROUP_SIZE = get_local_size(0); \n"
|
||||||
" int LOCAL_ID = get_local_id(0); \n"
|
//" int LOCAL_ID = get_local_id(0); \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" bool out = true; \n"
|
//" bool out = true; \n"
|
||||||
" float3 vertices3[3]; \n"
|
//" float3 vertices3[3]; \n"
|
||||||
" float2 uv_coords[3]; \n"
|
//" float2 uv_coords[3]; \n"
|
||||||
" float3 norms[3]; \n"
|
//" float3 norms[3]; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" for(int i = 0; i < 3; i++) { \n"
|
//" for(int i = 0; i < 3; i++) { \n"
|
||||||
" float4 vertex;// = vertices[faces[GROUP_ID * 3 + i ][0]]; \n"
|
//" float4 vertex;// = vertices[faces[GROUP_ID * 3 + i ][0]]; \n"
|
||||||
" for(int j = 0; j < 4; j ++) { \n"
|
//" for(int j = 0; j < 4; j ++) { \n"
|
||||||
" vertex[j] = vertices[4 * faces[GROUP_ID*3 + i][0] + j]; \n"
|
//" vertex[j] = vertices[4 * faces[GROUP_ID*3 + i][0] + j]; \n"
|
||||||
" } \n"
|
//" } \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" for(int j = 0; j < 2; j++) { \n"
|
//" for(int j = 0; j < 2; j++) { \n"
|
||||||
" uv_coords[i][j] = uv_buffer [2 * faces[GROUP_ID * 3 + i][1] + j]; \n"
|
//" uv_coords[i][j] = uv_buffer [2 * faces[GROUP_ID * 3 + i][1] + j]; \n"
|
||||||
" } \n"
|
//" } \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" for( int j = 0; j < 3; j++ ) { \n"
|
//" for( int j = 0; j < 3; j++ ) { \n"
|
||||||
" vertices3[j][i] = (vertex[j]/vertex[3]); \n"
|
//" vertices3[j][i] = (vertex[j]/vertex[3]); \n"
|
||||||
" norms[j][i] = norms_buff[3 * (faces[GROUP_ID * 3 + i][2]) + j]; \n"
|
//" norms[j][i] = norms_buff[3 * (faces[GROUP_ID * 3 + i][2]) + j]; \n"
|
||||||
" } \n"
|
//" } \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" if ( vertices3[0][i] > 0 && vertices3[0][i] < *screen_width \n"
|
//" if ( vertices3[0][i] > 0 && vertices3[0][i] < *screen_width \n"
|
||||||
" && vertices3[1][i] > 0 && vertices3[1][i] < *screen_width ) \n"
|
//" && vertices3[1][i] > 0 && vertices3[1][i] < *screen_width ) \n"
|
||||||
" out = false; \n"
|
//" out = false; \n"
|
||||||
" } \n"
|
//" } \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" if(out) return; \n"
|
//" if(out) return; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" //if(vertices3[1][0] == vertices3[1][1] && vertices3[1][2] == vertices3[1][1]) return; \n"
|
//" //if(vertices3[1][0] == vertices3[1][1] && vertices3[1][2] == vertices3[1][1]) return; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" int2 bounding_box_min = (int2) { *screen_width - 1, *screen_width - 1 }; \n"
|
//" int2 bounding_box_min = (int2) { *screen_width - 1, *screen_width - 1 }; \n"
|
||||||
" int2 bounding_box_max = (int2) { 0, 0 }; \n"
|
//" int2 bounding_box_max = (int2) { 0, 0 }; \n"
|
||||||
" int2 clamper = (int2) { *screen_width - 1, *screen_width - 1 }; \n"
|
//" int2 clamper = (int2) { *screen_width - 1, *screen_width - 1 }; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" for(int i = 0; i < 3; i++) { \n"
|
//" for(int i = 0; i < 3; i++) { \n"
|
||||||
" for(int j = 0; j < 2; j++) { \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_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"
|
//" bounding_box_max[j] = min(clamper[j], max(bounding_box_max[j], (int)vertices3[j][i])); \n"
|
||||||
" } \n"
|
//" } \n"
|
||||||
" } \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"
|
//" 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"
|
||||||
" \n"
|
//" \n"
|
||||||
" int X_PER_ITEM = (int)(ceil((float)(bounding_box_max[0] - bounding_box_min[0]) / 16.f)); \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"
|
//" 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 STARTING_X = bounding_box_min[0] + X_PER_ITEM * (LOCAL_ID % 16); \n"
|
||||||
" float ENDING_X = STARTING_X + X_PER_ITEM; \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 STARTING_Y = bounding_box_min[1] + Y_PER_ITEM * (LOCAL_ID / 16); \n"
|
||||||
" float ENDING_Y = STARTING_Y + Y_PER_ITEM; \n"
|
//" float ENDING_Y = STARTING_Y + Y_PER_ITEM; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" float3 point; \n"
|
//" float3 point; \n"
|
||||||
" for(point[0] = STARTING_X; point[0] <= ENDING_X; point[0]++) { \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"
|
//" for(point[1] = STARTING_Y; point[1] <= ENDING_Y; point[1]++) { \n"
|
||||||
" if(point[1] >= *screen_width || point[1] >= *screen_width) break; \n"
|
//" if(point[1] >= *screen_width || point[1] >= *screen_width) break; \n"
|
||||||
" float3 bc_coord = barycentric(vertices3, point); \n"
|
//" float3 bc_coord = barycentric(vertices3, point); \n"
|
||||||
" if (bc_coord[0] < 0 || bc_coord[1] < 0 || bc_coord[2] < 0) continue; \n"
|
//" if (bc_coord[0] < 0 || bc_coord[1] < 0 || bc_coord[2] < 0) continue; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" float2 uv_vec = (float2){0, 0}; \n"
|
//" float2 uv_vec = (float2){0, 0}; \n"
|
||||||
" float3 normal = (float3){0, 0, 0}; \n"
|
//" float3 normal = (float3){0, 0, 0}; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" point[2] = dot(vertices3[2], bc_coord); \n"
|
//" point[2] = dot(vertices3[2], bc_coord); \n"
|
||||||
" if (z_buffer[(int)(point[0] + point[1] * *screen_width)] > point[2]) { \n"
|
//" if (z_buffer[(int)(point[0] + point[1] * *screen_width)] > point[2]) { \n"
|
||||||
" continue; \n"
|
//" continue; \n"
|
||||||
" } \n"
|
//" } \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" for (int i = 0; i < 3; i++) { \n"
|
//" for (int i = 0; i < 3; i++) { \n"
|
||||||
" uv_vec[0] += uv_coords[i][0] * bc_coord[i]; \n"
|
//" uv_vec[0] += uv_coords[i][0] * bc_coord[i]; \n"
|
||||||
" uv_vec[1] += uv_coords[i][1] * bc_coord[i]; \n"
|
//" uv_vec[1] += uv_coords[i][1] * bc_coord[i]; \n"
|
||||||
" normal[i] = dot(norms[i], bc_coord); \n"
|
//" normal[i] = dot(norms[i], bc_coord); \n"
|
||||||
" } \n"
|
//" } \n"
|
||||||
" int2 uv_point = (int2) { (int)(uv_vec[0] * map_size[0]), (int)(uv_vec[1] * map_size[1]) }; \n"
|
//" int2 uv_point = (int2) { (int)(uv_vec[0] * map_size[0]), (int)(uv_vec[1] * map_size[1]) }; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" int col_index = 3 * (uv_point[0] + uv_point[1] * map_size[0]); \n"
|
//" int col_index = 3 * (uv_point[0] + uv_point[1] * map_size[0]); \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" float3 normalized_norm = normalize(normal); \n"
|
//" float3 normalized_norm = normalize(normal); \n"
|
||||||
" \n"
|
//" \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"
|
//" float intensity = clamp((dot(normalized_norm , (float3){light_dir[0], light_dir[1], light_dir[2]})), 0.f, 1.f) + 0.2; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" int color = 0; \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 + 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 + 1]) * intensity, (float) 0xff)) << 8; \n"
|
||||||
" color |= ((int)fmin((float)(diffuse_map[col_index + 2]) * intensity, (float) 0xff)) << 0; \n"
|
//" color |= ((int)fmin((float)(diffuse_map[col_index + 2]) * intensity, (float) 0xff)) << 0; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 16; \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)) << 8; \n"
|
||||||
" //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 0; \n"
|
//" //color |= ((int)fmin((float)(0xff) * intensity, (float) 0xff)) << 0; \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" \n"
|
//" \n"
|
||||||
" z_buffer[(int) (point[0] + point[1] * *screen_width)] = point[2]; \n"
|
//" z_buffer[(int) (point[0] + point[1] * *screen_width)] = point[2]; \n"
|
||||||
" pixels [(int) (point[0] + point[1] * *screen_width)] = color;// & 0x00ffffff; \n"
|
//" pixels [(int) (point[0] + point[1] * *screen_width)] = color;// & 0x00ffffff; \n"
|
||||||
" } \n"
|
//" } \n"
|
||||||
" } \n"
|
//" } \n"
|
||||||
"} \n";
|
//"} \n";
|
||||||
|
|||||||
+14
-173
@@ -1,5 +1,5 @@
|
|||||||
#include "kernels.h"
|
#include "kernels.h"
|
||||||
#include "util_window.h"
|
#include "renderer.h"
|
||||||
#include <debugapi.h>
|
#include <debugapi.h>
|
||||||
#include <ctime>
|
#include <ctime>
|
||||||
#include <climits>
|
#include <climits>
|
||||||
@@ -11,65 +11,21 @@ cl_platform_id platform_id;
|
|||||||
cl_device_id device_id;
|
cl_device_id device_id;
|
||||||
cl_context context;
|
cl_context context;
|
||||||
|
|
||||||
int triangles_drawn = 0;
|
|
||||||
|
|
||||||
// Matrix Multiplication Variables
|
|
||||||
#define MATRIX_SIZE 16
|
|
||||||
|
|
||||||
// vertex shader
|
|
||||||
cl_mem mat_z;
|
|
||||||
cl_mem vertices_mem;
|
|
||||||
cl_mem new_vertices_mem;
|
|
||||||
bool vertex_shader_buffers_initialized = false;
|
|
||||||
//==================
|
|
||||||
// fragment shader
|
|
||||||
// int* faces,
|
|
||||||
// int nfaces,
|
|
||||||
// float* new_verts,
|
|
||||||
// int nverts,
|
|
||||||
// int screen_width,
|
|
||||||
// int screen_height,
|
|
||||||
// float* z_buffer,
|
|
||||||
// float* uniform_m,
|
|
||||||
// float* uniform_mit,
|
|
||||||
// float* light_dir,
|
|
||||||
// const char* diffuse_map,
|
|
||||||
// const char* normal_map,
|
|
||||||
// const char* spec_map,
|
|
||||||
// char* pixel_data
|
|
||||||
cl_mem faces_buffer;
|
|
||||||
cl_mem uniform_m;
|
|
||||||
cl_mem uniform_mit;
|
|
||||||
cl_mem light_dir_buffer;
|
|
||||||
cl_mem diffuse_map_buffer;
|
|
||||||
cl_mem norms_mem;
|
|
||||||
cl_mem spec_map_buffer;
|
|
||||||
int* hidden_pixel_buffer;
|
int* hidden_pixel_buffer;
|
||||||
bool fragment_shader_buffers_initialized = false;
|
|
||||||
cl_mem triangles_verts_mem;
|
|
||||||
cl_mem bounding_box_min_mem;
|
|
||||||
cl_mem z_buffer_mem;
|
cl_mem z_buffer_mem;
|
||||||
cl_mem pixel_data_buffer;
|
cl_mem pixel_data_buffer;
|
||||||
cl_mem screen_width_mem;
|
cl_mem time_buffer;
|
||||||
cl_mem nfaces_mem;
|
|
||||||
cl_mem uv_buffer;
|
|
||||||
cl_mem map_size_buffer;
|
|
||||||
//==================
|
//==================
|
||||||
|
|
||||||
cl_platform_id* platforms = NULL;
|
cl_platform_id* platforms = NULL;
|
||||||
cl_device_id* devices = NULL;
|
cl_device_id* devices = NULL;
|
||||||
|
|
||||||
cl_program vertex_shader_prog;
|
|
||||||
cl_program fragment_shader_prog;
|
|
||||||
cl_kernel vertex_shader_kernel;
|
|
||||||
cl_kernel fragment_shader_kernel;
|
|
||||||
cl_command_queue commands;
|
cl_command_queue commands;
|
||||||
|
|
||||||
int err;
|
int err;
|
||||||
|
|
||||||
|
|
||||||
void init_kernels() {
|
void init_kernels() {
|
||||||
|
|
||||||
err = clGetPlatformIDs(0, NULL, &numPlatforms);
|
err = clGetPlatformIDs(0, NULL, &numPlatforms);
|
||||||
platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * numPlatforms);
|
platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * numPlatforms);
|
||||||
err = clGetPlatformIDs(numPlatforms, platforms, NULL);
|
err = clGetPlatformIDs(numPlatforms, platforms, NULL);
|
||||||
@@ -82,156 +38,41 @@ void init_kernels() {
|
|||||||
|
|
||||||
commands = clCreateCommandQueue(context, devices[0], 0, &err);
|
commands = clCreateCommandQueue(context, devices[0], 0, &err);
|
||||||
|
|
||||||
vertex_shader_prog = clCreateProgramWithSource(context, 1, (const char **)&vertex_shader_kernel_source, NULL, &err);
|
hidden_pixel_buffer = new int[screen_width*screen_height];
|
||||||
fragment_shader_prog = clCreateProgramWithSource(context, 1, (const char**)&fragment_shader_kernel_source, NULL, &err);
|
pixel_data_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * screen_height * screen_width , NULL, &err);
|
||||||
err = clBuildProgram(vertex_shader_prog, 1, devices, NULL, NULL, NULL);
|
z_buffer_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * screen_height * screen_width, NULL, &err);
|
||||||
err = clBuildProgram(fragment_shader_prog, 1, devices, NULL, NULL, NULL);
|
time_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float), NULL, &err);
|
||||||
|
|
||||||
if (err != CL_SUCCESS)
|
|
||||||
{
|
|
||||||
size_t len;
|
|
||||||
char buffer[2048];
|
|
||||||
printf("Error: Failed to build the fragment shader prog!\n");
|
|
||||||
clGetProgramBuildInfo(fragment_shader_prog, devices[0], CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
|
|
||||||
printf(buffer);
|
|
||||||
//return;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
vertex_shader_kernel = clCreateKernel(vertex_shader_prog, "vertex_shader", &err);
|
|
||||||
fragment_shader_kernel = clCreateKernel(fragment_shader_prog, "fragment_shader", &err);
|
|
||||||
err = 0;
|
|
||||||
}
|
|
||||||
|
|
||||||
void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices) {
|
|
||||||
if(!vertex_shader_buffers_initialized)
|
|
||||||
{
|
|
||||||
mat_z = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * MATRIX_SIZE , NULL, &err);
|
|
||||||
vertices_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * vertex_count * 3, NULL, &err);
|
|
||||||
new_vertices_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float4) * vertex_count, NULL, &err);
|
|
||||||
|
|
||||||
err = clSetKernelArg(vertex_shader_kernel, 0, sizeof(cl_mem), &mat_z);
|
|
||||||
err = clSetKernelArg(vertex_shader_kernel, 1, sizeof(cl_mem), &vertices_mem);
|
|
||||||
err = clSetKernelArg(vertex_shader_kernel, 2, sizeof(cl_mem), &new_vertices_mem);
|
|
||||||
vertex_shader_buffers_initialized = true;
|
|
||||||
}
|
|
||||||
|
|
||||||
err = clEnqueueWriteBuffer(commands, mat_z , CL_TRUE, 0, sizeof(float) * MATRIX_SIZE, z, 0, NULL, NULL);
|
|
||||||
err = clEnqueueWriteBuffer(commands, vertices_mem, CL_TRUE, 0, sizeof(float) * vertex_count * 3, vertices, 0, NULL, NULL);
|
|
||||||
|
|
||||||
size_t vertex_shader_global[] = { vertex_count*4 };
|
|
||||||
size_t vertex_shader_local[] = { 4 };
|
|
||||||
err = clEnqueueNDRangeKernel(commands, vertex_shader_kernel, 1, NULL, vertex_shader_global, vertex_shader_local, 0, NULL, NULL);
|
|
||||||
}
|
|
||||||
|
|
||||||
void clear_pixel_buffer() {
|
|
||||||
for (int i = 0; i < screen_width * screen_height; i++) {
|
|
||||||
hidden_pixel_buffer[i] = 0;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
void clear(cl_mem* buffer, size_t size, const int pattern) {
|
void clear(cl_mem* buffer, size_t size, const int pattern) {
|
||||||
clEnqueueFillBuffer(commands, *buffer, &pattern, sizeof(int), 0, size, 0, NULL, NULL);
|
clEnqueueFillBuffer(commands, *buffer, &pattern, sizeof(int), 0, size, 0, NULL, NULL);
|
||||||
}
|
}
|
||||||
|
|
||||||
void fragment_shader(
|
void new_frame() {
|
||||||
cl_int3* faces,
|
|
||||||
int nfaces,
|
|
||||||
float* uv,
|
|
||||||
size_t uv_size,
|
|
||||||
float* uniform_m,
|
|
||||||
float* uniform_mit,
|
|
||||||
float* light_dir,
|
|
||||||
unsigned char* diffuse_map,
|
|
||||||
float* norms,
|
|
||||||
size_t norms_size,
|
|
||||||
unsigned char* spec_map,
|
|
||||||
int* map_size
|
|
||||||
) {
|
|
||||||
if (!fragment_shader_buffers_initialized) {
|
|
||||||
|
|
||||||
faces_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(cl_int3) * 3 * nfaces , NULL, &err);
|
|
||||||
pixel_data_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * screen_height * screen_width , NULL, &err);
|
|
||||||
screen_width_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int) , NULL, &err);
|
|
||||||
z_buffer_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * screen_height * screen_width, NULL, &err);
|
|
||||||
nfaces_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int) , NULL, &err);
|
|
||||||
uv_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , uv_size , NULL, &err);
|
|
||||||
map_size_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int) * 2 , NULL, &err);
|
|
||||||
norms_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , norms_size , NULL, &err);
|
|
||||||
light_dir_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * 3 , NULL, &err);
|
|
||||||
diffuse_map_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(char) * 3 * map_size[0] * map_size[1] , NULL, &err);
|
|
||||||
|
|
||||||
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 0, sizeof(cl_mem), &faces_buffer);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 1, sizeof(cl_mem), &new_vertices_mem);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 2, sizeof(cl_mem), &pixel_data_buffer);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 3, sizeof(cl_mem), &screen_width_mem);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 4, sizeof(cl_mem), &z_buffer_mem);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 5, sizeof(cl_mem), &nfaces_mem);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 6, sizeof(cl_mem), &uv_buffer);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 7, sizeof(cl_mem), &map_size_buffer);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 8, sizeof(cl_mem), &light_dir_buffer);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 9, sizeof(cl_mem), &norms_mem);
|
|
||||||
err = clSetKernelArg(fragment_shader_kernel, 10, sizeof(cl_mem), &diffuse_map_buffer);
|
|
||||||
|
|
||||||
err = clEnqueueWriteBuffer(commands, screen_width_mem , CL_FALSE, 0, sizeof(int) , &screen_width , 0, NULL, NULL);
|
|
||||||
err = clEnqueueWriteBuffer(commands, faces_buffer , CL_FALSE, 0, sizeof(cl_int3) * 3 * nfaces , faces , 0, NULL, NULL);
|
|
||||||
err = clEnqueueWriteBuffer(commands, nfaces_mem , CL_FALSE, 0, sizeof(int) , &nfaces , 0, NULL, NULL);
|
|
||||||
err = clEnqueueWriteBuffer(commands, uv_buffer , CL_FALSE, 0, uv_size , uv , 0, NULL, NULL);
|
|
||||||
err = clEnqueueWriteBuffer(commands, map_size_buffer , CL_FALSE, 0, sizeof(int) * 2 , map_size , 0, NULL, NULL);
|
|
||||||
err = clEnqueueWriteBuffer(commands, light_dir_buffer , CL_FALSE, 0, sizeof(float) * 3 , light_dir , 0, NULL, NULL);
|
|
||||||
err = clEnqueueWriteBuffer(commands, norms_mem , CL_FALSE, 0, norms_size , norms , 0, NULL, NULL);
|
|
||||||
err = clEnqueueWriteBuffer(commands, diffuse_map_buffer, CL_FALSE, 0, sizeof(char) * 3 * map_size[0] * map_size[1], diffuse_map , 0, NULL, NULL);
|
|
||||||
|
|
||||||
hidden_pixel_buffer = new int[screen_width*screen_height];
|
|
||||||
fragment_shader_buffers_initialized = true;
|
|
||||||
}
|
|
||||||
|
|
||||||
clear(&pixel_data_buffer, sizeof(int) * screen_width * screen_height, 0);
|
clear(&pixel_data_buffer, sizeof(int) * screen_width * screen_height, 0);
|
||||||
clear(&z_buffer_mem, sizeof(float) * 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() {
|
||||||
size_t fragment_shader_global[] = { nfaces * 256 };
|
clFinish(commands);
|
||||||
size_t framgent_shader_local[] = { 256 };
|
clEnqueueReadBuffer(commands, pixel_data_buffer, CL_TRUE, 0, sizeof(int) * screen_height * screen_width, hidden_pixel_buffer, 0, NULL, NULL);
|
||||||
|
|
||||||
err = clEnqueueNDRangeKernel(commands, fragment_shader_kernel, 1, NULL, fragment_shader_global, NULL, 0, NULL, NULL);
|
|
||||||
|
|
||||||
err = clEnqueueReadBuffer(commands, pixel_data_buffer, CL_TRUE, 0, sizeof(int) * screen_height * screen_width, hidden_pixel_buffer, 0, NULL, NULL);
|
|
||||||
|
|
||||||
for (int j = 0; j < screen_height; j++)
|
for (int j = 0; j < screen_height; j++)
|
||||||
for (int i = 0; i < screen_width; i++)
|
for (int i = 0; i < screen_width; i++)
|
||||||
set_pixel(i, j, hidden_pixel_buffer[j * screen_width + i]);
|
set_pixel(i, j, hidden_pixel_buffer[j * screen_width + i]);
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
}
|
}
|
||||||
|
|
||||||
void destroy_kernels()
|
void destroy_kernels()
|
||||||
{
|
{
|
||||||
clReleaseProgram(vertex_shader_prog);
|
|
||||||
clReleaseKernel(vertex_shader_kernel);
|
|
||||||
if (vertex_shader_buffers_initialized) {
|
|
||||||
clReleaseMemObject(mat_z);
|
|
||||||
clReleaseMemObject(vertices_mem);
|
|
||||||
clReleaseMemObject(new_vertices_mem);
|
|
||||||
}
|
|
||||||
|
|
||||||
if (fragment_shader_buffers_initialized) {
|
|
||||||
clReleaseMemObject(faces_buffer);
|
|
||||||
clReleaseMemObject(screen_width_mem);
|
|
||||||
clReleaseMemObject(z_buffer_mem);
|
clReleaseMemObject(z_buffer_mem);
|
||||||
clReleaseMemObject(pixel_data_buffer);
|
clReleaseMemObject(pixel_data_buffer);
|
||||||
clReleaseMemObject(screen_width_mem);
|
clReleaseMemObject(time_buffer);
|
||||||
clReleaseMemObject(nfaces_mem);
|
|
||||||
clReleaseMemObject(uv_buffer);
|
|
||||||
clReleaseMemObject(map_size_buffer);
|
|
||||||
clReleaseMemObject(light_dir_buffer);
|
|
||||||
clReleaseMemObject(norms_mem);
|
|
||||||
clReleaseMemObject(diffuse_map_buffer);
|
|
||||||
}
|
|
||||||
clReleaseCommandQueue(commands);
|
clReleaseCommandQueue(commands);
|
||||||
|
|
||||||
free(hidden_pixel_buffer);
|
free(hidden_pixel_buffer);
|
||||||
|
|
||||||
|
|
||||||
clReleaseContext(context);
|
clReleaseContext(context);
|
||||||
}
|
}
|
||||||
|
|||||||
+13
-16
@@ -3,6 +3,7 @@
|
|||||||
#include <Windows.h>
|
#include <Windows.h>
|
||||||
#include <CL/cl.h>
|
#include <CL/cl.h>
|
||||||
#include "geometry.h"
|
#include "geometry.h"
|
||||||
|
#include "util_window.h"
|
||||||
|
|
||||||
|
|
||||||
extern cl_uint numPlatforms;
|
extern cl_uint numPlatforms;
|
||||||
@@ -31,25 +32,21 @@ extern const char* matrix_mul_kernel_source;
|
|||||||
|
|
||||||
extern const char* vertex_shader_kernel_source;
|
extern const char* vertex_shader_kernel_source;
|
||||||
extern const char* fragment_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 init_kernels();
|
||||||
void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices);
|
|
||||||
void fragment_shader(
|
|
||||||
cl_int3* faces,
|
|
||||||
int nfaces,
|
|
||||||
float* uv,
|
|
||||||
size_t uv_size,
|
|
||||||
float* uniform_m,
|
|
||||||
float* uniform_mit,
|
|
||||||
float* light_dir,
|
|
||||||
unsigned char* diffuse_map,
|
|
||||||
float* norms,
|
|
||||||
size_t norms_size,
|
|
||||||
unsigned char* spec_map,
|
|
||||||
int* map_size
|
|
||||||
);
|
|
||||||
|
|
||||||
void destroy_kernels();
|
void destroy_kernels();
|
||||||
|
void new_frame();
|
||||||
|
void end_frame();
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|||||||
+132208
File diff suppressed because it is too large
Load Diff
+2
-1
@@ -75,7 +75,7 @@ bool HandleButtonPressed() {
|
|||||||
|
|
||||||
void CALLBACK FixedUpdate(HWND hwnd, UINT message, UINT uInt, DWORD dWord)
|
void CALLBACK FixedUpdate(HWND hwnd, UINT message, UINT uInt, DWORD dWord)
|
||||||
{
|
{
|
||||||
TIME += 0.03333333;
|
TIME += 1.f / (float)TARGET_FRAMERATE;
|
||||||
HandleButtonPressed();
|
HandleButtonPressed();
|
||||||
camera.ApplyChanges();
|
camera.ApplyChanges();
|
||||||
clock_t start = clock();
|
clock_t start = clock();
|
||||||
@@ -98,6 +98,7 @@ LRESULT CALLBACK WndProc(HWND hwnd, UINT message, WPARAM wParam, LPARAM lParam)
|
|||||||
break;
|
break;
|
||||||
case WM_CLOSE:
|
case WM_CLOSE:
|
||||||
DestroyWindow(hwnd);
|
DestroyWindow(hwnd);
|
||||||
|
free_renderer();
|
||||||
destroy_kernels();
|
destroy_kernels();
|
||||||
break;
|
break;
|
||||||
case WM_DESTROY:
|
case WM_DESTROY:
|
||||||
|
|||||||
+279
-6
@@ -6,7 +6,8 @@
|
|||||||
#define PI 3.14159265358979323846
|
#define PI 3.14159265358979323846
|
||||||
#define DEG2RAD PI/180
|
#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;
|
std::ifstream in;
|
||||||
Transform = Matrix::identity();
|
Transform = Matrix::identity();
|
||||||
Rotation = Matrix::identity();
|
Rotation = Matrix::identity();
|
||||||
@@ -45,13 +46,13 @@ Model::Model(const char *filename) : verts_(), faces_(), norms_(), uv_(), diffus
|
|||||||
faces_.push_back(f);
|
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, "_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() {
|
void Model::ApplyTransform() {
|
||||||
Transform = Translation * Scale * Rotation;
|
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) {
|
if (dot!=std::string::npos) {
|
||||||
texfile = texfile.substr(0,dot) + std::string(suffix);
|
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;
|
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();
|
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";
|
||||||
|
|||||||
+25
-1
@@ -4,10 +4,13 @@
|
|||||||
#include <string>
|
#include <string>
|
||||||
#include "geometry.h"
|
#include "geometry.h"
|
||||||
#include "tgaimage.h"
|
#include "tgaimage.h"
|
||||||
|
#include "kernels.h"
|
||||||
|
#include "util_window.h"
|
||||||
|
|
||||||
class Model {
|
class Model {
|
||||||
private:
|
private:
|
||||||
void load_texture(std::string filename, const char *suffix, TGAImage &img);
|
void load_texture(std::string filename, const char *suffix, TGAImage &img);
|
||||||
|
int fluid;
|
||||||
public:
|
public:
|
||||||
TGAImage diffusemap_;
|
TGAImage diffusemap_;
|
||||||
TGAImage normalmap_;
|
TGAImage normalmap_;
|
||||||
@@ -16,7 +19,23 @@ public:
|
|||||||
std::vector<std::vector<Vec3i> > faces_; // attention, this Vec3i means vertex/uv/normal
|
std::vector<std::vector<Vec3i> > faces_; // attention, this Vec3i means vertex/uv/normal
|
||||||
std::vector<Vec3f> norms_;
|
std::vector<Vec3f> norms_;
|
||||||
std::vector<Vec2f> uv_;
|
std::vector<Vec2f> uv_;
|
||||||
Model(const char *filename);
|
cl_program vertex_shader_prog;
|
||||||
|
cl_program fragment_shader_prog;
|
||||||
|
cl_kernel vertex_shader_kernel;
|
||||||
|
cl_kernel fragment_shader_kernel;
|
||||||
|
cl_mem vertex_shader_matz;
|
||||||
|
cl_mem vertex_shader_vertices;
|
||||||
|
cl_mem new_vertices_mem;
|
||||||
|
cl_mem fragment_shader_faces;
|
||||||
|
cl_mem fragment_shader_screen_width;
|
||||||
|
cl_mem fragment_shader_uv;
|
||||||
|
cl_mem fragment_shader_map_size;
|
||||||
|
cl_mem fragment_shader_norms;
|
||||||
|
cl_mem fragment_shader_light_dir;
|
||||||
|
cl_mem fragment_shader_diffuse_map;
|
||||||
|
cl_int3* faces;
|
||||||
|
|
||||||
|
Model(const char *filename, int watery);
|
||||||
~Model();
|
~Model();
|
||||||
int nverts();
|
int nverts();
|
||||||
int nfaces();
|
int nfaces();
|
||||||
@@ -33,6 +52,11 @@ public:
|
|||||||
void rotate(Vec3f rot);
|
void rotate(Vec3f rot);
|
||||||
void scale(Vec3f scl);
|
void scale(Vec3f scl);
|
||||||
void ApplyTransform();
|
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);
|
TGAColor diffuse(Vec2f uv);
|
||||||
float specular(Vec2f uv);
|
float specular(Vec2f uv);
|
||||||
std::vector<int> face(int idx);
|
std::vector<int> face(int idx);
|
||||||
|
|||||||
+304054
File diff suppressed because it is too large
Load Diff
Binary file not shown.
|
After Width: | Height: | Size: 3.0 MiB |
+335413
File diff suppressed because it is too large
Load Diff
Binary file not shown.
|
After Width: | Height: | Size: 50 MiB |
+20
-131
@@ -5,6 +5,7 @@
|
|||||||
#include "util_renderer.h"
|
#include "util_renderer.h"
|
||||||
#include "CL/cl.h"
|
#include "CL/cl.h"
|
||||||
#include "kernels.h"
|
#include "kernels.h"
|
||||||
|
#include "improv_gfx.h"
|
||||||
#include <ctime>
|
#include <ctime>
|
||||||
|
|
||||||
#pragma comment (lib, "x86_64/opencl.lib")
|
#pragma comment (lib, "x86_64/opencl.lib")
|
||||||
@@ -16,28 +17,19 @@
|
|||||||
#define NEAR_CLIP_PLANE 1.f
|
#define NEAR_CLIP_PLANE 1.f
|
||||||
#define FAR_CLIP_PLANE 2000.0f
|
#define FAR_CLIP_PLANE 2000.0f
|
||||||
#define FOV 50
|
#define FOV 50
|
||||||
#define CAMERA_MOVEMENT_SPEED .7f
|
#define CAMERA_MOVEMENT_SPEED .1f
|
||||||
#define DEFAULT_CAMERA_POS Vec3f(0, 0, 5)
|
#define DEFAULT_CAMERA_POS Vec3f(0, 0, 5)
|
||||||
#define DEFAULT_CAMERA_ROT Vec3f(0, 0, 0)
|
#define DEFAULT_CAMERA_ROT Vec3f(0, 0, 0)
|
||||||
#define LIGHT_INTENSITY 1.5
|
#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 ViewPort = Matrix::identity();
|
||||||
Matrix ModelView = Matrix::identity();
|
Matrix ModelView = Matrix::identity();
|
||||||
Matrix Projection = Matrix::identity();
|
Matrix Projection = Matrix::identity();
|
||||||
|
|
||||||
Model* model = new Model("sakura.obj");
|
Model* model, *model2;
|
||||||
Camera camera;
|
Camera camera;
|
||||||
|
|
||||||
Vec3f light_dir = Vec3f(1, 1, 1).normalize();
|
Vec3f light_dir = Vec3f(0, 1, 1).normalize();
|
||||||
|
|
||||||
float* new_verts = (float*)malloc(4 * sizeof(float) * model->nverts());
|
|
||||||
cl_int3* faces = (cl_int3*)malloc(3 * sizeof(cl_int3) * model->nfaces());
|
|
||||||
|
|
||||||
bool init_flag = false;
|
bool init_flag = false;
|
||||||
|
|
||||||
void init_camera() {
|
void init_camera() {
|
||||||
@@ -54,55 +46,15 @@ void init_camera() {
|
|||||||
camera.ApplyChanges();
|
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;
|
|
||||||
Matrix z;
|
|
||||||
|
|
||||||
virtual Vec4f vertex(int iface, int nthvert) {
|
|
||||||
//varying_uv_coords.set_col(nthvert, model->vert(iface, nthvert));
|
|
||||||
Vec4f gl_Vertex = embed<4>(model->verts_[model->faces_[iface][nthvert][0]]);
|
|
||||||
//return ViewPort * Projection * ModelView * gl_Vertex; // transform it to screen coordinates
|
|
||||||
return z * gl_Vertex;
|
|
||||||
//return Vec4f(0,0,0,0);
|
|
||||||
}
|
|
||||||
|
|
||||||
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(3 + c[i] * ((1 * diff_intensity + 0.1 * spec_intensity)), 255);// *LIGHT_INTENSITY;
|
|
||||||
return false;
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
|
|
||||||
void render()
|
void render()
|
||||||
{
|
{
|
||||||
if (!init_flag) {
|
if (!init_flag) {
|
||||||
//light_dir = camera.GetForward().normalize() * -1;
|
|
||||||
viewport(0, 0, screen_width, screen_height, FAR_CLIP_PLANE, NEAR_CLIP_PLANE);
|
viewport(0, 0, screen_width, screen_height, FAR_CLIP_PLANE, NEAR_CLIP_PLANE);
|
||||||
for (int i = 0; i < model->nfaces(); i++) {
|
Setup();
|
||||||
for (int j = 0; j < 3; j++) {
|
|
||||||
faces[i * 3 + j].x = model->faces_[i][j][0];
|
|
||||||
faces[i * 3 + j].y = model->faces_[i][j][1];
|
|
||||||
faces[i * 3 + j].z = model->faces_[i][j][2];
|
|
||||||
}
|
|
||||||
}
|
|
||||||
init_flag = true;
|
init_flag = true;
|
||||||
}
|
}
|
||||||
|
|
||||||
@@ -111,84 +63,21 @@ void render()
|
|||||||
ModelView = camera.GetModelViewMatrix();
|
ModelView = camera.GetModelViewMatrix();
|
||||||
}
|
}
|
||||||
|
|
||||||
{
|
Matrix z = ViewPort * Projection * ModelView;
|
||||||
//model->rotate(Vec3f(0, 0, 90));
|
|
||||||
//model->ApplyTransform();
|
new_frame();
|
||||||
|
Draw();
|
||||||
|
for(Model* model : models_in_scene)
|
||||||
|
model->render(&z, (float*)&light_dir);
|
||||||
|
end_frame();
|
||||||
|
|
||||||
}
|
}
|
||||||
|
|
||||||
// float* normal_test = *(float**)((Vec3f*)&model->norms_);
|
void AddModel(Model* model) {
|
||||||
//
|
models_in_scene.push_back(model);
|
||||||
// Vec3f smth = Vec3f(
|
|
||||||
// normal_test[3 * faces[0 * 9 + 0 * 3 + 2] + 0],
|
|
||||||
// normal_test[3 * faces[0 * 9 + 0 * 3 + 2] + 1],
|
|
||||||
// normal_test[3 * faces[0 * 9 + 0 * 3 + 2] + 2]
|
|
||||||
// ).normalize();
|
|
||||||
// printf("Real: %f, %f, %f\n", model->normal(0, 0)[0], model->normal(0, 0)[1], model->normal(0, 0)[2]);
|
|
||||||
// printf("Please: %f, %f, %f\n", smth[0], smth[1], smth[2]);
|
|
||||||
|
|
||||||
//TextureShader shader;
|
|
||||||
//shader.uniform_m = (Projection);
|
|
||||||
//shader.uniform_mit = (Projection).invert_transpose();
|
|
||||||
|
|
||||||
Matrix z = ViewPort * Projection * ModelView * model->Transform;
|
|
||||||
|
|
||||||
Matrix uniform_m = (Projection);
|
|
||||||
Matrix uniform_mit = (Projection).invert_transpose();
|
|
||||||
|
|
||||||
// Vertex Shader: Should be called per model
|
|
||||||
vertex_shader((float*)&z, *(float**)((Vec3f*) &model->verts_), model->nverts(), new_verts);
|
|
||||||
|
|
||||||
// Things needed in the GPU fragment shader
|
|
||||||
// [x] model->faces_
|
|
||||||
// [x] model->nfaces()
|
|
||||||
// [x] new_verts
|
|
||||||
// [x] model->nverts()
|
|
||||||
// [x] screen_width
|
|
||||||
// [x] screen_height
|
|
||||||
// [x] z_buffer
|
|
||||||
// [x] uniform_m
|
|
||||||
// [x] uniform_mit
|
|
||||||
// [x] light direction
|
|
||||||
// [ ] diffuse map
|
|
||||||
// [ ] normal map
|
|
||||||
// [ ] specular map
|
|
||||||
// [ ] pixel_data
|
|
||||||
|
|
||||||
int map_size[] = {model->diffusemap_.get_width(), model->diffusemap_.get_height()};
|
|
||||||
|
|
||||||
fragment_shader(
|
|
||||||
faces,
|
|
||||||
model->nfaces(),
|
|
||||||
*(float**)((Vec2f*)&model->uv_),
|
|
||||||
sizeof(float) * model->uv_.size() * 2,
|
|
||||||
(float*) &uniform_m,
|
|
||||||
(float*) &uniform_mit,
|
|
||||||
(float*) &light_dir,
|
|
||||||
model->diffusemap_.data,
|
|
||||||
*(float**)((Vec3f*)&model->norms_),
|
|
||||||
sizeof(float) * model->norms_.size() * 3,
|
|
||||||
model->specularmap_.data,
|
|
||||||
map_size
|
|
||||||
);
|
|
||||||
|
|
||||||
// Here starts the CPU fragment shader
|
|
||||||
//printf("Here starts the loop\n");
|
|
||||||
//#pragma omp parallel for
|
|
||||||
//for (int i = 0; i < model->nfaces(); i++) {
|
|
||||||
// Vec4f screen_coords[3];
|
|
||||||
// bool out = true;
|
|
||||||
// #pragma omp parallel for
|
|
||||||
// for (int j = 0; j < 3; j++) {
|
|
||||||
// screen_coords[j] = ((Vec4f*)new_verts)[model->faces_[i][j][0]];
|
|
||||||
// Vec3f screen3(screen_coords[j]);
|
|
||||||
|
|
||||||
// shader.varying_uv_coords.set_col(j, model->uv(i, j));
|
|
||||||
// if (screen3.x > 0 && screen3.x < screen_width && screen3.y > 0 && screen3.y < screen_height) out = false;
|
|
||||||
// }
|
|
||||||
// if(!out)
|
|
||||||
// triangle(screen_coords, shader);
|
|
||||||
//}
|
|
||||||
|
|
||||||
//printf("that's it\n");
|
|
||||||
}
|
}
|
||||||
|
|
||||||
|
void free_renderer() {
|
||||||
|
for(Model* model : models_in_scene)
|
||||||
|
delete model;
|
||||||
|
}
|
||||||
|
|||||||
@@ -8,7 +8,9 @@ extern float* z_buffer;
|
|||||||
extern Camera camera;
|
extern Camera camera;
|
||||||
extern float TIME;
|
extern float TIME;
|
||||||
|
|
||||||
|
|
||||||
void init_camera();
|
void init_camera();
|
||||||
void render();
|
void render();
|
||||||
|
void free_renderer();
|
||||||
int color_to_int(TGAColor col);
|
int color_to_int(TGAColor col);
|
||||||
#endif
|
#endif
|
||||||
+135072
File diff suppressed because it is too large
Load Diff
Binary file not shown.
+47676
File diff suppressed because it is too large
Load Diff
Binary file not shown.
|
After Width: | Height: | Size: 12 MiB |
Reference in New Issue
Block a user