mo7sen/ImprovGFX
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Adding basic texture fragment shader kernel

Browse Source
This commit is contained in:
mo7sener
2019-12-25 00:09:14 +02:00
parent e57702936e
commit 2e6ddd2ff2
11 changed files with 412 additions and 89 deletions
Download Patch File Download Diff File Expand all files Collapse all files
OpenWindow/OpenWindow.vcxproj
+4 -4
View File
@@ -76,7 +76,7 @@
<Optimization>Disabled</Optimization> <Optimization>Disabled</Optimization>
<SDLCheck>true</SDLCheck> <SDLCheck>true</SDLCheck>
<ConformanceMode>true</ConformanceMode> <ConformanceMode>true</ConformanceMode>
<AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories> <AdditionalIncludeDirectories>$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
<DisableSpecificWarnings>4996;</DisableSpecificWarnings> <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
<PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions> <PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
</ClCompile> </ClCompile>
@@ -91,7 +91,7 @@
<Optimization>Disabled</Optimization> <Optimization>Disabled</Optimization>
<SDLCheck>true</SDLCheck> <SDLCheck>true</SDLCheck>
<ConformanceMode>true</ConformanceMode> <ConformanceMode>true</ConformanceMode>
<AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories> <AdditionalIncludeDirectories>$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
<DisableSpecificWarnings>4996;</DisableSpecificWarnings> <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
<PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions> <PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
</ClCompile> </ClCompile>
@@ -108,7 +108,7 @@
<IntrinsicFunctions>true</IntrinsicFunctions> <IntrinsicFunctions>true</IntrinsicFunctions>
<SDLCheck>true</SDLCheck> <SDLCheck>true</SDLCheck>
<ConformanceMode>true</ConformanceMode> <ConformanceMode>true</ConformanceMode>
<AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories> <AdditionalIncludeDirectories>$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
<DisableSpecificWarnings>4996;</DisableSpecificWarnings> <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
<PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions> <PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
</ClCompile> </ClCompile>
@@ -127,7 +127,7 @@
<IntrinsicFunctions>true</IntrinsicFunctions> <IntrinsicFunctions>true</IntrinsicFunctions>
<SDLCheck>true</SDLCheck> <SDLCheck>true</SDLCheck>
<ConformanceMode>true</ConformanceMode> <ConformanceMode>true</ConformanceMode>
<AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories> <AdditionalIncludeDirectories>$(ProjectDir)\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
<DisableSpecificWarnings>4996;</DisableSpecificWarnings> <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
<PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions> <PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
</ClCompile> </ClCompile>
OpenWindow/camera.cpp
+7 -14
View File
@@ -101,31 +101,24 @@ Matrix Camera::GetModelViewMatrix() {
Vec3f center = position + forward; Vec3f center = position + forward;
Vec3f z = forward * -1; Vec3f z = forward * -1;
Matrix Minv = Matrix::identity(); Matrix Minv = Matrix::identity();
Matrix Tr = Matrix::identity();
for (int i = 0; i < 3; i++) { for (int i = 0; i < 3; i++) {
Minv[0][i] = right[i]; Minv[0][i] = right[i];
Minv[1][i] = up[i]; Minv[1][i] = up[i];
Minv[2][i] = z[i]; Minv[2][i] = z[i];
Tr[i][3] = -center[i];
} }
// Minv[3][0] = (right * position); Minv[0][3] = -(right * position);
// Minv[3][1] = (up * position); Minv[1][3] = -(up * position);
// Minv[3][2] = (z * position); Minv[2][3] = -(z * position);
return Minv * Tr; return Minv;
} }
Matrix Camera::GetProjectionMatrix() { Matrix Camera::GetProjectionMatrix() {
Matrix Projection = Matrix::identity(); Matrix Projection = Matrix::identity();
Projection[0][0] = 1 / tan(fov * DEG2RAD); Projection[0][0] = 1 / tan(fov * DEG2RAD / 2);
Projection[1][1] = 1 / tan(fov * DEG2RAD); Projection[1][1] = 1 / tan(fov * DEG2RAD / 2);
Projection[2][2] = (far_plane + near_plane) / (far_plane - near_plane); Projection[2][2] = (far_plane + near_plane) / (far_plane - near_plane);
Projection[2][3] = (2 * far_plane * near_plane) / (far_plane - near_plane); Projection[2][3] = (2 * far_plane * near_plane) / (far_plane - near_plane);
//Projection[2][2] = -(far_plane) / (far_plane - near_plane);
//Projection[2][3] = -1;
//Projection[3][2] = -(far_plane * near_plane) / (far_plane - near_plane);
//Projection[2][2] = -2 / (far_plane - near_plane);
//Projection[2][3] = (far_plane + near_plane) / (far_plane - near_plane);
Projection[3][2] = -1; Projection[3][2] = -1;
Projection[3][3] = 1; Projection[3][3] = 0;
return Projection; return Projection;
} }
OpenWindow/kernel_sources.cpp
+112
View File
@@ -14,3 +14,115 @@ const char* vertex_shader_kernel_source =
" + 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 =
"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 int* 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 index = get_global_id(0); \n"
" \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[(index * 9) + (3 * i) + (0)]]; \n"
" for(int j = 0; j < 4; j ++) { \n"
" vertex[j] = vertices[4 * faces[(index*9) + (3 * i) + (0)] + j]; \n"
" } \n"
" \n"
" for(int j = 0; j < 2; j++) { \n"
" uv_coords[i][j] = uv_buffer [2 * faces[(index * 9) + (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[(index * 9) + (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"
" float3 point; \n"
" for(point[0] = bounding_box_min[0]; point[0] <= bounding_box_max[0]; point[0]++) { \n"
" for(point[1] = bounding_box_min[1]; point[1] <= bounding_box_max[1]; point[1]++) { \n"
" \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); \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"
" z_buffer[(int) (point[0] + point[1] * *screen_width)] = point[2]; \n"
" pixels [(int) (point[0] + point[1] * *screen_width)] = color;// & 0x00ffffff; \n"
" } \n"
" } \n"
"} \n";
OpenWindow/kernels.cpp
+149 -15
View File
@@ -1,5 +1,8 @@
#include "kernels.h" #include "kernels.h"
#include "util_window.h"
#include <debugapi.h> #include <debugapi.h>
#include <ctime>
#include <climits>
cl_uint numPlatforms; cl_uint numPlatforms;
cl_uint numDevices; cl_uint numDevices;
@@ -8,7 +11,7 @@ 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 // Matrix Multiplication Variables
#define MATRIX_SIZE 16 #define MATRIX_SIZE 16
@@ -17,13 +20,49 @@ cl_context context;
cl_mem mat_z; cl_mem mat_z;
cl_mem vertices_mem; cl_mem vertices_mem;
cl_mem new_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_platform_id* platforms = NULL;
cl_device_id* devices = NULL; cl_device_id* devices = NULL;
cl_program vertex_shader_prog; cl_program vertex_shader_prog;
cl_program fragment_shader_prog;
cl_kernel vertex_shader_kernel; cl_kernel vertex_shader_kernel;
cl_kernel fragment_shader_kernel;
cl_command_queue commands; cl_command_queue commands;
int err; int err;
@@ -44,23 +83,28 @@ 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); 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(vertex_shader_prog, 1, devices, NULL, NULL, NULL);
err = clBuildProgram(fragment_shader_prog, 1, devices, NULL, NULL, NULL);
if (err != CL_SUCCESS) if (err != CL_SUCCESS)
{ {
size_t len; size_t len;
char buffer[2048]; char buffer[2048];
printf("Error: Failed to build program executable!\n"); printf("Error: Failed to build the fragment shader prog!\n");
//clGetProgramBuildInfo(mat_mul_prog, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); clGetProgramBuildInfo(fragment_shader_prog, devices[0], CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
printf(buffer); printf(buffer);
return; //return;
} }
vertex_shader_kernel = clCreateKernel(vertex_shader_prog, "vertex_shader", &err); vertex_shader_kernel = clCreateKernel(vertex_shader_prog, "vertex_shader", &err);
fragment_shader_kernel = clCreateKernel(fragment_shader_prog, "fragment_shader", &err);
err = 0; err = 0;
} }
void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices) { 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); 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); 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(float) * vertex_count * 4, NULL, &err); new_vertices_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * vertex_count * 4, NULL, &err);
@@ -68,34 +112,124 @@ void vertex_shader(float* z, float* vertices, int vertex_count, float* new_verti
err = clSetKernelArg(vertex_shader_kernel, 0, sizeof(cl_mem), &mat_z); 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, 1, sizeof(cl_mem), &vertices_mem);
err = clSetKernelArg(vertex_shader_kernel, 2, sizeof(cl_mem), &new_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, 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); 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_global[] = { vertex_count*4 };
size_t vertex_shader_local[] = { 4 }; size_t vertex_shader_local[] = { 4 };
err = clEnqueueNDRangeKernel(commands, vertex_shader_kernel, 1, NULL, vertex_shader_global, vertex_shader_local, 0, NULL, NULL); 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(
int* 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(int) * 3 * 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(int) * 3 * 3 * nfaces , faces , 0, NULL, NULL);
err = clEnqueueWriteBuffer(commands, nfaces_mem , CL_FALSE, 0, sizeof(int) , &nfaces , 0, NULL, NULL);
err = clEnqueueWriteBuffer(commands, uv_buffer , CL_FALSE, 0, uv_size , uv , 0, NULL, NULL);
err = clEnqueueWriteBuffer(commands, map_size_buffer , CL_FALSE, 0, sizeof(int) * 2 , map_size , 0, NULL, NULL);
err = clEnqueueWriteBuffer(commands, light_dir_buffer , CL_FALSE, 0, sizeof(float) * 3 , light_dir , 0, NULL, NULL);
err = clEnqueueWriteBuffer(commands, norms_mem , CL_FALSE, 0, norms_size , norms , 0, NULL, NULL);
err = clEnqueueWriteBuffer(commands, diffuse_map_buffer, CL_FALSE, 0, sizeof(char) * 3 * map_size[0] * map_size[1], diffuse_map , 0, NULL, NULL);
hidden_pixel_buffer = new int[screen_width*screen_height];
fragment_shader_buffers_initialized = true;
}
clear(&pixel_data_buffer, sizeof(int) * screen_width * screen_height, 0);
clear(&z_buffer_mem, sizeof(float) * screen_width * screen_height, 0);
size_t fragment_shader_global[] = { nfaces };
err = clEnqueueNDRangeKernel(commands, fragment_shader_kernel, 1, NULL, fragment_shader_global, NULL, 0, NULL, NULL);
err = clEnqueueReadBuffer(commands, pixel_data_buffer, CL_TRUE, 0, sizeof(int) * screen_height * screen_width, hidden_pixel_buffer, 0, NULL, NULL);
for (int j = 0; j < screen_height; j++)
for (int i = 0; i < screen_width; i++)
set_pixel(i, j, hidden_pixel_buffer[j * screen_width + i]);
clFlush(commands);
clFinish(commands);
err = clEnqueueReadBuffer(commands, new_vertices_mem, CL_TRUE, 0, sizeof(float) * vertex_count * 4, new_vertices, 0, NULL, NULL);
clReleaseMemObject(mat_z);
clReleaseMemObject(vertices_mem);
clReleaseMemObject(new_vertices_mem);
} }
void destroy_kernels() void destroy_kernels()
{ {
clReleaseProgram(vertex_shader_prog); clReleaseProgram(vertex_shader_prog);
clReleaseKernel(vertex_shader_kernel); clReleaseKernel(vertex_shader_kernel);
// clReleaseMemObject(mat_z); if (vertex_shader_buffers_initialized) {
// clReleaseMemObject(vertices_mem); clReleaseMemObject(mat_z);
// clReleaseMemObject(new_vertices_mem); clReleaseMemObject(vertices_mem);
clReleaseMemObject(new_vertices_mem);
}
if (fragment_shader_buffers_initialized) {
clReleaseMemObject(faces_buffer);
clReleaseMemObject(screen_width_mem);
clReleaseMemObject(z_buffer_mem);
clReleaseMemObject(pixel_data_buffer);
clReleaseMemObject(screen_width_mem);
clReleaseMemObject(nfaces_mem);
clReleaseMemObject(uv_buffer);
clReleaseMemObject(map_size_buffer);
clReleaseMemObject(light_dir_buffer);
clReleaseMemObject(norms_mem);
clReleaseMemObject(diffuse_map_buffer);
}
clReleaseCommandQueue(commands); clReleaseCommandQueue(commands);
free(hidden_pixel_buffer);
clReleaseContext(context); clReleaseContext(context);
OpenWindow/kernels.h
+16 -7
View File
@@ -26,20 +26,29 @@ extern cl_command_queue commands;
extern int err; extern int err;
#define DIM 2
extern size_t mat_mul_global[DIM];
extern size_t mat_mul_local[DIM];
extern const char* matrix_mul_kernel_source; extern const char* matrix_mul_kernel_source;
// End of: Matrix Multiplication Variables // End of: Matrix Multiplication Variables
extern const char* vertex_shader_kernel_source; extern const char* vertex_shader_kernel_source;
extern const char* fragment_shader_kernel_source;
void init_kernels(); void init_kernels();
void mat4_mul(float* A, float* B, float* C);
void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices); void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices);
void fragment_shader(
int* 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();
OpenWindow/main.cpp
+3
View File
@@ -78,7 +78,10 @@ void CALLBACK FixedUpdate(HWND hwnd, UINT message, UINT uInt, DWORD dWord)
TIME += 0.167; TIME += 0.167;
HandleButtonPressed(); HandleButtonPressed();
camera.ApplyChanges(); camera.ApplyChanges();
clock_t start = clock();
render(); render();
clock_t end = clock();
printf("FPS: %f\n", 1/((float)(end-start)/CLOCKS_PER_SEC));
Update(); Update();
} }
OpenWindow/model.h
+2 -2
View File
@@ -7,11 +7,11 @@
class Model { class Model {
private: private:
void load_texture(std::string filename, const char *suffix, TGAImage &img);
public:
TGAImage diffusemap_; TGAImage diffusemap_;
TGAImage normalmap_; TGAImage normalmap_;
TGAImage specularmap_; TGAImage specularmap_;
void load_texture(std::string filename, const char *suffix, TGAImage &img);
public:
std::vector<Vec3f> verts_; std::vector<Vec3f> verts_;
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_;
OpenWindow/renderer.cpp
+81 -25
View File
@@ -13,10 +13,10 @@
#define VERTICAL_CAMERA_SPEED 0.1 #define VERTICAL_CAMERA_SPEED 0.1
#define VERTICAL_CAMERA_CLAMP_UP 90 #define VERTICAL_CAMERA_CLAMP_UP 90
#define VERTICAL_CAMERA_CLAMP_DOWN -90 #define VERTICAL_CAMERA_CLAMP_DOWN -90
#define NEAR_CLIP_PLANE 0.3f #define NEAR_CLIP_PLANE 1.f
#define FAR_CLIP_PLANE 10.0f #define FAR_CLIP_PLANE 200.0f
#define FOV 30 #define FOV 50
#define CAMERA_MOVEMENT_SPEED 0.05f #define CAMERA_MOVEMENT_SPEED .1f
#define DEFAULT_CAMERA_POS Vec3f(0, 0, 5) #define DEFAULT_CAMERA_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
@@ -36,6 +36,9 @@ Camera camera;
Vec3f light_dir = Vec3f(1, 1, 1).normalize(); Vec3f light_dir = Vec3f(1, 1, 1).normalize();
float* new_verts = (float*)malloc(4 * sizeof(float) * model->nverts()); float* new_verts = (float*)malloc(4 * sizeof(float) * model->nverts());
int* faces = (int*)malloc(3 * 3 * sizeof(int) * model->nfaces());
bool init_flag = false;
void init_camera() { void init_camera() {
camera.SetPosition(DEFAULT_CAMERA_POS); camera.SetPosition(DEFAULT_CAMERA_POS);
@@ -82,7 +85,7 @@ struct TextureShader : public IShader {
TGAColor c = model->diffuse(uv); TGAColor c = model->diffuse(uv);
color = c; color = c;
for (int i = 0; i < 3; i++) for (int i = 0; i < 3; i++)
color[i] = std::fmin(2 + c[i] * (( 1 * diff_intensity + 1 * spec_intensity)), 255) * LIGHT_INTENSITY; color[i] = std::fmin(3 + c[i] * ((1 * diff_intensity + 0.1 * spec_intensity)), 255);// *LIGHT_INTENSITY;
return false; return false;
} }
}; };
@@ -90,12 +93,17 @@ struct TextureShader : public IShader {
void render() void render()
{ {
{ if (!init_flag) {
//light_dir = camera.GetForward().normalize() * -1; //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++)
for (int k = 0; k < 3; k++)
faces[i * 9 + j * 3 + k] = model->faces_[i][j][k];
init_flag = true;
} }
{ {
viewport(0, 0, screen_width, screen_height, FAR_CLIP_PLANE, NEAR_CLIP_PLANE);
Projection = camera.GetProjectionMatrix(); Projection = camera.GetProjectionMatrix();
ModelView = camera.GetModelViewMatrix(); ModelView = camera.GetModelViewMatrix();
} }
@@ -105,31 +113,79 @@ void render()
//model->ApplyTransform(); //model->ApplyTransform();
} }
clear_zbuffer(); // float* normal_test = *(float**)((Vec3f*)&model->norms_);
TextureShader shader; //
shader.uniform_m = (Projection); // Vec3f smth = Vec3f(
shader.uniform_mit = (Projection).invert_transpose(); // 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 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: Should be called per model
vertex_shader((float*)&z, *(float**)((Vec3f*) &model->verts_), model->nverts(), new_verts); vertex_shader((float*)&z, *(float**)((Vec3f*) &model->verts_), model->nverts(), new_verts);
#pragma omp parallel for // Things needed in the GPU fragment shader
for (int i = 0; i < model->nfaces(); i++) { // [x] model->faces_
Vec4f screen_coords[3]; // [x] model->nfaces()
bool out = true; // [x] new_verts
#pragma omp parallel for // [x] model->nverts()
for (int j = 0; j < 3; j++) { // [x] screen_width
screen_coords[j] = ((Vec4f*)new_verts)[model->faces_[i][j][0]]; // [x] screen_height
Vec3f screen3(screen_coords[j]); // [x] z_buffer
// [x] uniform_m
// [x] uniform_mit
// [x] light direction
// [ ] diffuse map
// [ ] normal map
// [ ] specular map
// [ ] pixel_data
shader.varying_uv_coords.set_col(j, model->uv(i, j)); int map_size[] = {model->diffusemap_.get_width(), model->diffusemap_.get_height()};
if (screen3.x > 0 && screen3.x < screen_width && screen3.y > 0 && screen3.y < screen_height) out = false;
} fragment_shader(
if(!out) faces,
triangle(screen_coords, shader); model->nfaces(),
} *(float**)((Vec2f*)&model->uv_),
sizeof(float) * model->uv_.size() * 2,
(float*) &uniform_m,
(float*) &uniform_mit,
(float*) &light_dir,
model->diffusemap_.data,
*(float**)((Vec3f*)&model->norms_),
sizeof(float) * model->norms_.size() * 3,
model->specularmap_.data,
map_size
);
// Here starts the CPU fragment shader
//printf("Here starts the loop\n");
//#pragma omp parallel for
//for (int i = 0; i < model->nfaces(); i++) {
// Vec4f screen_coords[3];
// bool out = true;
// #pragma omp parallel for
// for (int j = 0; j < 3; j++) {
// screen_coords[j] = ((Vec4f*)new_verts)[model->faces_[i][j][0]];
// Vec3f screen3(screen_coords[j]);
// shader.varying_uv_coords.set_col(j, model->uv(i, j));
// if (screen3.x > 0 && screen3.x < screen_width && screen3.y > 0 && screen3.y < screen_height) out = false;
// }
// if(!out)
// triangle(screen_coords, shader);
//}
//printf("that's it\n");
} }
OpenWindow/tgaimage.h
+1 -2
View File
@@ -62,14 +62,13 @@ struct TGAColor {
class TGAImage { class TGAImage {
protected: protected:
unsigned char* data;
int width; int width;
int height; int height;
int bytespp; int bytespp;
bool load_rle_data(std::ifstream &in); bool load_rle_data(std::ifstream &in);
bool unload_rle_data(std::ofstream &out); bool unload_rle_data(std::ofstream &out);
public: public:
unsigned char* data;
enum Format { enum Format {
GRAYSCALE=1, RGB=3, RGBA=4 GRAYSCALE=1, RGB=3, RGBA=4
}; };
OpenWindow/util_renderer.cpp
+29 -12
View File
@@ -1,5 +1,6 @@
#include "util_renderer.h" #include "util_renderer.h"
#include "util_window.h" #include "util_window.h"
#include "kernels.h"
IShader::~IShader() {} IShader::~IShader() {}
@@ -16,7 +17,7 @@ void viewport(int x, int y, int w, int h, int far_plane, int near_plane) {
} }
int color_to_int(TGAColor col) { int color_to_int(TGAColor col) {
return (col[2] << 16) | (col[1] << 8) | col[0]; return (col[0] << 24) | (col[1] << 16) | (col[2] << 8) | col[3];
} }
Vec3f barycentric(Vec3f* pts, Vec3f P) Vec3f barycentric(Vec3f* pts, Vec3f P)
@@ -52,22 +53,38 @@ void triangle( Vec4f* pts, IShader &shader)
bounding_box_max[j] = std::fmin(clamp[j], std::fmax(bounding_box_max[j], (int)pts3[i][j])); bounding_box_max[j] = std::fmin(clamp[j], std::fmax(bounding_box_max[j], (int)pts3[i][j]));
} }
//printf("%f, %f, %f\n", pts3[0][0], pts3[0][1], pts3[0][2]);
//float* points_test = (float*)&pts3;
//printf("%f, %f, %f\n", points_test[0], points_test[1], points_test[2]);
//printf("%d, %d\n", bounding_box_max[0], bounding_box_max[1]);
//fragment_shader_pixelwise(
// *(float**)((Vec3f*)&pts3),
// *(int**)((Vec2i*)&bounding_box_min),
// *(int**)((Vec2i*)&bounding_box_max));
//fragment_shader_pixelwise(
// (float*)(&pts3),
// (int*)(&bounding_box_min),
// (int*)(&bounding_box_max));
Vec3i P; Vec3i P;
#pragma omp parallel for #pragma omp parallel for
for (P.x = bounding_box_min.x; P.x <= bounding_box_max.x; P.x++) { for (P.x = bounding_box_min.x; P.x <= bounding_box_max.x; P.x++) {
for (P.y = bounding_box_min.y; P.y <= bounding_box_max.y; P.y++) { for (P.y = bounding_box_min.y; P.y <= bounding_box_max.y; P.y++) {
Vec3f bc_coord = barycentric(pts3, P); //Vec3f bc_coord = barycentric(pts3, P);
float frag_depth = 0; //float frag_depth = 0;
for (int i = 0; i < 3; i++) //for (int i = 0; i < 3; i++)
frag_depth += pts3[i][2] * bc_coord[i]; // frag_depth += pts3[i][2] * bc_coord[i];
if (bc_coord.x < 0 || bc_coord.y < 0 || bc_coord.z < 0 || z_buffer[ P.x + P.y * screen_width ]>frag_depth) continue; //if (bc_coord.x < 0 || bc_coord.y < 0 || bc_coord.z < 0 || z_buffer[ P.x + P.y * screen_width ]>frag_depth) continue;
TGAColor color; //TGAColor color;
bool discard = shader.fragment(bc_coord, color); //bool discard = shader.fragment(bc_coord, color);
if (!discard) { //if (!discard) {
z_buffer[P.x + P.y * screen_width] = frag_depth; // z_buffer[P.x + P.y * screen_width] = frag_depth;
set_pixel(P.x, P.y, color_to_int(color)); // set_pixel(P.x, P.y, color_to_int(color));
} set_pixel(P.x, P.y, 0x00ff00);
//}
} }
} }
} }
OpenWindow/util_window.cpp
+2 -2
View File
@@ -95,9 +95,9 @@ void destroy_window() {
void set_pixel(unsigned int x, unsigned int y, unsigned int color) { void set_pixel(unsigned int x, unsigned int y, unsigned int color) {
unsigned long pixel_index = y * bytes_per_row + x * 3; unsigned long pixel_index = y * bytes_per_row + x * 3;
pixel_data[pixel_index + 0] = (char)(color >> 0); pixel_data[pixel_index + 2] = (char)(color >> 0);
pixel_data[pixel_index + 1] = (char)(color >> 8); pixel_data[pixel_index + 1] = (char)(color >> 8);
pixel_data[pixel_index + 2] = (char)(color >> 16); pixel_data[pixel_index + 0] = (char)(color >> 16);
if (!screen_changed) if (!screen_changed)
screen_changed = true; screen_changed = true;