diff --git a/OpenWindow/OpenWindow.vcxproj b/OpenWindow/OpenWindow.vcxproj index 7b45e66..99901a3 100644 --- a/OpenWindow/OpenWindow.vcxproj +++ b/OpenWindow/OpenWindow.vcxproj @@ -141,6 +141,7 @@ + $(ProjectDir)\include;%(AdditionalIncludeDirectories) $(ProjectDir)\include;%(AdditionalIncludeDirectories) @@ -158,6 +159,7 @@ + diff --git a/OpenWindow/OpenWindow.vcxproj.filters b/OpenWindow/OpenWindow.vcxproj.filters index fe8b030..3b7b9ae 100644 --- a/OpenWindow/OpenWindow.vcxproj.filters +++ b/OpenWindow/OpenWindow.vcxproj.filters @@ -45,6 +45,9 @@ Source Files + + Source Files + @@ -71,5 +74,8 @@ Header Files + + Header Files + \ No newline at end of file diff --git a/OpenWindow/african_head_nm_tangent.tga b/OpenWindow/african_head_nm_tangent.tga deleted file mode 100644 index 46ee251..0000000 Binary files a/OpenWindow/african_head_nm_tangent.tga and /dev/null differ diff --git a/OpenWindow/african_head_spec.tga b/OpenWindow/african_head_spec.tga deleted file mode 100644 index 12e02ce..0000000 Binary files a/OpenWindow/african_head_spec.tga and /dev/null differ diff --git a/OpenWindow/improv_gfx.cpp b/OpenWindow/improv_gfx.cpp new file mode 100644 index 0000000..5ed55ad --- /dev/null +++ b/OpenWindow/improv_gfx.cpp @@ -0,0 +1,18 @@ +#include "improv_gfx.h" + +void Setup() { + + AddModel(new Model("sakura.obj")); + + AddModel(new Model("african_head.obj")); + + Model* model2 = new Model("african_head.obj"); + model2->translate(Vec3f(1, 0, 0)); + model2->ApplyTransform(); + AddModel(model2); + + Model* model3 = new Model("african_head.obj"); + model3->translate(Vec3f(-1, 0, 0)); + model3->ApplyTransform(); + AddModel(model3); +} diff --git a/OpenWindow/improv_gfx.h b/OpenWindow/improv_gfx.h new file mode 100644 index 0000000..e96b7d3 --- /dev/null +++ b/OpenWindow/improv_gfx.h @@ -0,0 +1,9 @@ +#ifndef __IMPROV_GFX__ +#define __IMPROV_GFX__ +#include "model.h" +#include "geometry.h" + +void Setup(); +void AddModel(Model* model); + +#endif \ No newline at end of file diff --git a/OpenWindow/kernel_sources.cpp b/OpenWindow/kernel_sources.cpp index ad68bc1..8f50247 100644 --- a/OpenWindow/kernel_sources.cpp +++ b/OpenWindow/kernel_sources.cpp @@ -1,145 +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"; +//#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"; diff --git a/OpenWindow/kernels.cpp b/OpenWindow/kernels.cpp index 88e0a06..8c4a37d 100644 --- a/OpenWindow/kernels.cpp +++ b/OpenWindow/kernels.cpp @@ -11,65 +11,20 @@ cl_platform_id platform_id; cl_device_id device_id; cl_context context; -int triangles_drawn = 0; - -// Matrix Multiplication Variables -#define MATRIX_SIZE 16 - -// vertex shader -cl_mem mat_z; -cl_mem vertices_mem; -cl_mem new_vertices_mem; -bool vertex_shader_buffers_initialized = false; -//================== -// fragment shader -// int* faces, -// int nfaces, -// float* new_verts, -// int nverts, -// int screen_width, -// int screen_height, -// float* z_buffer, -// float* uniform_m, -// float* uniform_mit, -// float* light_dir, -// const char* diffuse_map, -// const char* normal_map, -// const char* spec_map, -// char* pixel_data -cl_mem faces_buffer; -cl_mem uniform_m; -cl_mem uniform_mit; -cl_mem light_dir_buffer; -cl_mem diffuse_map_buffer; -cl_mem norms_mem; -cl_mem spec_map_buffer; int* hidden_pixel_buffer; -bool fragment_shader_buffers_initialized = false; -cl_mem triangles_verts_mem; -cl_mem bounding_box_min_mem; cl_mem z_buffer_mem; cl_mem pixel_data_buffer; -cl_mem screen_width_mem; -cl_mem nfaces_mem; -cl_mem uv_buffer; -cl_mem map_size_buffer; //================== cl_platform_id* platforms = NULL; cl_device_id* devices = NULL; -cl_program vertex_shader_prog; -cl_program fragment_shader_prog; -cl_kernel vertex_shader_kernel; -cl_kernel fragment_shader_kernel; cl_command_queue commands; int err; void init_kernels() { - err = clGetPlatformIDs(0, NULL, &numPlatforms); platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * numPlatforms); err = clGetPlatformIDs(numPlatforms, platforms, NULL); @@ -82,156 +37,37 @@ void init_kernels() { commands = clCreateCommandQueue(context, devices[0], 0, &err); - vertex_shader_prog = clCreateProgramWithSource(context, 1, (const char **)&vertex_shader_kernel_source, NULL, &err); - fragment_shader_prog = clCreateProgramWithSource(context, 1, (const char**)&fragment_shader_kernel_source, NULL, &err); - err = clBuildProgram(vertex_shader_prog, 1, devices, NULL, NULL, NULL); - err = clBuildProgram(fragment_shader_prog, 1, devices, NULL, NULL, NULL); - - if (err != CL_SUCCESS) - { - size_t len; - char buffer[2048]; - printf("Error: Failed to build the fragment shader prog!\n"); - clGetProgramBuildInfo(fragment_shader_prog, devices[0], CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); - printf(buffer); - //return; - } - - vertex_shader_kernel = clCreateKernel(vertex_shader_prog, "vertex_shader", &err); - fragment_shader_kernel = clCreateKernel(fragment_shader_prog, "fragment_shader", &err); - err = 0; + 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); } -void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices) { - if(!vertex_shader_buffers_initialized) - { - mat_z = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * MATRIX_SIZE , NULL, &err); - vertices_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * vertex_count * 3, NULL, &err); - new_vertices_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float4) * vertex_count, NULL, &err); - - err = clSetKernelArg(vertex_shader_kernel, 0, sizeof(cl_mem), &mat_z); - err = clSetKernelArg(vertex_shader_kernel, 1, sizeof(cl_mem), &vertices_mem); - err = clSetKernelArg(vertex_shader_kernel, 2, sizeof(cl_mem), &new_vertices_mem); - vertex_shader_buffers_initialized = true; - } - - err = clEnqueueWriteBuffer(commands, mat_z , CL_TRUE, 0, sizeof(float) * MATRIX_SIZE, z, 0, NULL, NULL); - err = clEnqueueWriteBuffer(commands, vertices_mem, CL_TRUE, 0, sizeof(float) * vertex_count * 3, vertices, 0, NULL, NULL); - - size_t vertex_shader_global[] = { vertex_count*4 }; - size_t vertex_shader_local[] = { 4 }; - err = clEnqueueNDRangeKernel(commands, vertex_shader_kernel, 1, NULL, vertex_shader_global, vertex_shader_local, 0, NULL, NULL); -} - -void clear_pixel_buffer() { - for (int i = 0; i < screen_width * screen_height; i++) { - hidden_pixel_buffer[i] = 0; - } -} void clear(cl_mem* buffer, size_t size, const int pattern) { clEnqueueFillBuffer(commands, *buffer, &pattern, sizeof(int), 0, size, 0, NULL, NULL); } -void fragment_shader( - cl_int3* faces, - int nfaces, - float* uv, - size_t uv_size, - float* uniform_m, - float* uniform_mit, - float* light_dir, - unsigned char* diffuse_map, - float* norms, - size_t norms_size, - unsigned char* spec_map, - int* map_size -) { - if (!fragment_shader_buffers_initialized) { - - faces_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(cl_int3) * 3 * nfaces , NULL, &err); - pixel_data_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * screen_height * screen_width , NULL, &err); - screen_width_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int) , NULL, &err); - z_buffer_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * screen_height * screen_width, NULL, &err); - nfaces_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int) , NULL, &err); - uv_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , uv_size , NULL, &err); - map_size_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(int) * 2 , NULL, &err); - norms_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , norms_size , NULL, &err); - light_dir_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * 3 , NULL, &err); - diffuse_map_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(char) * 3 * map_size[0] * map_size[1] , NULL, &err); - - - err = clSetKernelArg(fragment_shader_kernel, 0, sizeof(cl_mem), &faces_buffer); - err = clSetKernelArg(fragment_shader_kernel, 1, sizeof(cl_mem), &new_vertices_mem); - err = clSetKernelArg(fragment_shader_kernel, 2, sizeof(cl_mem), &pixel_data_buffer); - err = clSetKernelArg(fragment_shader_kernel, 3, sizeof(cl_mem), &screen_width_mem); - err = clSetKernelArg(fragment_shader_kernel, 4, sizeof(cl_mem), &z_buffer_mem); - err = clSetKernelArg(fragment_shader_kernel, 5, sizeof(cl_mem), &nfaces_mem); - err = clSetKernelArg(fragment_shader_kernel, 6, sizeof(cl_mem), &uv_buffer); - err = clSetKernelArg(fragment_shader_kernel, 7, sizeof(cl_mem), &map_size_buffer); - err = clSetKernelArg(fragment_shader_kernel, 8, sizeof(cl_mem), &light_dir_buffer); - err = clSetKernelArg(fragment_shader_kernel, 9, sizeof(cl_mem), &norms_mem); - err = clSetKernelArg(fragment_shader_kernel, 10, sizeof(cl_mem), &diffuse_map_buffer); - - err = clEnqueueWriteBuffer(commands, screen_width_mem , CL_FALSE, 0, sizeof(int) , &screen_width , 0, NULL, NULL); - err = clEnqueueWriteBuffer(commands, faces_buffer , CL_FALSE, 0, sizeof(cl_int3) * 3 * nfaces , faces , 0, NULL, NULL); - err = clEnqueueWriteBuffer(commands, nfaces_mem , CL_FALSE, 0, sizeof(int) , &nfaces , 0, NULL, NULL); - err = clEnqueueWriteBuffer(commands, uv_buffer , CL_FALSE, 0, uv_size , uv , 0, NULL, NULL); - err = clEnqueueWriteBuffer(commands, map_size_buffer , CL_FALSE, 0, sizeof(int) * 2 , map_size , 0, NULL, NULL); - err = clEnqueueWriteBuffer(commands, light_dir_buffer , CL_FALSE, 0, sizeof(float) * 3 , light_dir , 0, NULL, NULL); - err = clEnqueueWriteBuffer(commands, norms_mem , CL_FALSE, 0, norms_size , norms , 0, NULL, NULL); - err = clEnqueueWriteBuffer(commands, diffuse_map_buffer, CL_FALSE, 0, sizeof(char) * 3 * map_size[0] * map_size[1], diffuse_map , 0, NULL, NULL); - - hidden_pixel_buffer = new int[screen_width*screen_height]; - fragment_shader_buffers_initialized = true; - } - +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); +} - - size_t fragment_shader_global[] = { nfaces * 256 }; - size_t framgent_shader_local[] = { 256 }; - - err = clEnqueueNDRangeKernel(commands, fragment_shader_kernel, 1, NULL, fragment_shader_global, NULL, 0, NULL, NULL); - - err = clEnqueueReadBuffer(commands, pixel_data_buffer, CL_TRUE, 0, sizeof(int) * screen_height * screen_width, hidden_pixel_buffer, 0, NULL, NULL); +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() { - clReleaseProgram(vertex_shader_prog); - clReleaseKernel(vertex_shader_kernel); - if (vertex_shader_buffers_initialized) { - clReleaseMemObject(mat_z); - clReleaseMemObject(vertices_mem); - clReleaseMemObject(new_vertices_mem); - } - - if (fragment_shader_buffers_initialized) { - clReleaseMemObject(faces_buffer); - clReleaseMemObject(screen_width_mem); - clReleaseMemObject(z_buffer_mem); - clReleaseMemObject(pixel_data_buffer); - clReleaseMemObject(screen_width_mem); - clReleaseMemObject(nfaces_mem); - clReleaseMemObject(uv_buffer); - clReleaseMemObject(map_size_buffer); - clReleaseMemObject(light_dir_buffer); - clReleaseMemObject(norms_mem); - clReleaseMemObject(diffuse_map_buffer); - } + clReleaseMemObject(z_buffer_mem); + clReleaseMemObject(pixel_data_buffer); clReleaseCommandQueue(commands); free(hidden_pixel_buffer); - clReleaseContext(context); } diff --git a/OpenWindow/kernels.h b/OpenWindow/kernels.h index 0465e3a..7d532f6 100644 --- a/OpenWindow/kernels.h +++ b/OpenWindow/kernels.h @@ -32,24 +32,18 @@ extern const char* matrix_mul_kernel_source; extern const char* vertex_shader_kernel_source; extern const char* fragment_shader_kernel_source; -void init_kernels(); -void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices); -void fragment_shader( - cl_int3* faces, - int nfaces, - float* uv, - size_t uv_size, - float* uniform_m, - float* uniform_mit, - float* light_dir, - unsigned char* diffuse_map, - float* norms, - size_t norms_size, - unsigned char* spec_map, - int* map_size -); +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; + +void init_kernels(); void destroy_kernels(); +void new_frame(); +void end_frame(); diff --git a/OpenWindow/main.cpp b/OpenWindow/main.cpp index 39e05ac..db8ed76 100644 --- a/OpenWindow/main.cpp +++ b/OpenWindow/main.cpp @@ -98,6 +98,7 @@ 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: diff --git a/OpenWindow/model.cpp b/OpenWindow/model.cpp index 8ddf0cc..8a279f0 100644 --- a/OpenWindow/model.cpp +++ b/OpenWindow/model.cpp @@ -45,13 +45,16 @@ 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; + //std::cerr << "# v# " << verts_.size() << " f# " << faces_.size() << " vt# " << uv_.size() << " vn# " << norms_.size() << std::endl; load_texture(filename, "_diffuse.tga", diffusemap_); - load_texture(filename, "_nm_tangent.tga", normalmap_); - load_texture(filename, "_spec.tga", specularmap_); + //load_texture(filename, "_nm_tangent.tga", normalmap_); + //load_texture(filename, "_spec.tga", specularmap_); + init_kernels(); } -Model::~Model() {} +Model::~Model() { + release_kernels(); +} void Model::ApplyTransform() { Transform = Translation * Scale * Rotation; @@ -141,3 +144,244 @@ Vec3f Model::normal(int iface, int nthvert) { return norms_[idx].normalize(); } +void Model::init_kernels() { + 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); + 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* 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 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"; diff --git a/OpenWindow/model.h b/OpenWindow/model.h index 196dd6b..7617e61 100644 --- a/OpenWindow/model.h +++ b/OpenWindow/model.h @@ -4,6 +4,8 @@ #include #include "geometry.h" #include "tgaimage.h" +#include "kernels.h" +#include "util_window.h" class Model { private: @@ -16,6 +18,22 @@ public: std::vector > faces_; // attention, this Vec3i means vertex/uv/normal std::vector norms_; std::vector uv_; + 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); ~Model(); int nverts(); @@ -33,6 +51,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 face(int idx); diff --git a/OpenWindow/renderer.cpp b/OpenWindow/renderer.cpp index 2b376fa..0c2f446 100644 --- a/OpenWindow/renderer.cpp +++ b/OpenWindow/renderer.cpp @@ -5,6 +5,7 @@ #include "util_renderer.h" #include "CL/cl.h" #include "kernels.h" +#include "improv_gfx.h" #include #pragma comment (lib, "x86_64/opencl.lib") @@ -21,23 +22,14 @@ #define DEFAULT_CAMERA_ROT Vec3f(0, 0, 0) #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("sakura.obj"); +Model* model, *model2; Camera camera; Vec3f light_dir = Vec3f(1, 1, 1).normalize(); - -float* new_verts = (float*)malloc(4 * sizeof(float) * model->nverts()); -cl_int3* faces = (cl_int3*)malloc(3 * sizeof(cl_int3) * model->nfaces()); - bool init_flag = false; void init_camera() { @@ -54,55 +46,15 @@ void init_camera() { camera.ApplyChanges(); } -void clear_zbuffer() -{ - for (int i = 0; i < screen_width * screen_height; i++) - z_buffer[i] = 0; -} - -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; - } -}; +std::vector models_in_scene; void render() { if (!init_flag) { - //light_dir = camera.GetForward().normalize() * -1; viewport(0, 0, screen_width, screen_height, FAR_CLIP_PLANE, NEAR_CLIP_PLANE); - for (int i = 0; i < model->nfaces(); i++) { - for (int j = 0; j < 3; j++) { - faces[i * 3 + j].x = model->faces_[i][j][0]; - faces[i * 3 + j].y = model->faces_[i][j][1]; - faces[i * 3 + j].z = model->faces_[i][j][2]; - } - } + Setup(); init_flag = true; } @@ -110,85 +62,21 @@ void render() Projection = camera.GetProjectionMatrix(); ModelView = camera.GetModelViewMatrix(); } - - { - //model->rotate(Vec3f(0, 0, 90)); - //model->ApplyTransform(); - } -// float* normal_test = *(float**)((Vec3f*)&model->norms_); -// -// Vec3f smth = Vec3f( -// normal_test[3 * faces[0 * 9 + 0 * 3 + 2] + 0], -// normal_test[3 * faces[0 * 9 + 0 * 3 + 2] + 1], -// normal_test[3 * faces[0 * 9 + 0 * 3 + 2] + 2] -// ).normalize(); -// printf("Real: %f, %f, %f\n", model->normal(0, 0)[0], model->normal(0, 0)[1], model->normal(0, 0)[2]); -// printf("Please: %f, %f, %f\n", smth[0], smth[1], smth[2]); + Matrix z = ViewPort * Projection * ModelView; - //TextureShader shader; - //shader.uniform_m = (Projection); - //shader.uniform_mit = (Projection).invert_transpose(); + new_frame(); + for(Model* model : models_in_scene) + model->render(&z, (float*)&light_dir); + end_frame(); - 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 AddModel(Model* model) { + models_in_scene.push_back(model); +} + +void free_renderer() { + for(Model* model : models_in_scene) + delete model; +} diff --git a/OpenWindow/renderer.h b/OpenWindow/renderer.h index ae4152d..5029116 100644 --- a/OpenWindow/renderer.h +++ b/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 \ No newline at end of file diff --git a/OpenWindow/sakura_diffuse.tga b/OpenWindow/sakura_diffuse.tga index 9d7d13d..bdd085f 100644 Binary files a/OpenWindow/sakura_diffuse.tga and b/OpenWindow/sakura_diffuse.tga differ