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

Vertex Shader is now GPU-Accelerated

Browse Source
This commit is contained in:
mo7sener
2019-12-18 19:09:00 +02:00
parent a47ccbc3cd
commit e57702936e
6 changed files with 51 additions and 126 deletions
Download Patch File Download Diff File Expand all files Collapse all files
OpenWindow/OpenWindow.vcxproj
+8
View File
@@ -78,9 +78,11 @@
<ConformanceMode>true</ConformanceMode> <ConformanceMode>true</ConformanceMode>
<AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories> <AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories>
<DisableSpecificWarnings>4996;</DisableSpecificWarnings> <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
<PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
</ClCompile> </ClCompile>
<Link> <Link>
<AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories> <AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
<SubSystem>NotSet</SubSystem>
</Link> </Link>
</ItemDefinitionGroup> </ItemDefinitionGroup>
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'"> <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">
@@ -91,9 +93,11 @@
<ConformanceMode>true</ConformanceMode> <ConformanceMode>true</ConformanceMode>
<AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories> <AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories>
<DisableSpecificWarnings>4996;</DisableSpecificWarnings> <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
<PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
</ClCompile> </ClCompile>
<Link> <Link>
<AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories> <AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
<SubSystem>NotSet</SubSystem>
</Link> </Link>
</ItemDefinitionGroup> </ItemDefinitionGroup>
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'"> <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
@@ -106,11 +110,13 @@
<ConformanceMode>true</ConformanceMode> <ConformanceMode>true</ConformanceMode>
<AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories> <AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories>
<DisableSpecificWarnings>4996;</DisableSpecificWarnings> <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
<PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
</ClCompile> </ClCompile>
<Link> <Link>
<EnableCOMDATFolding>true</EnableCOMDATFolding> <EnableCOMDATFolding>true</EnableCOMDATFolding>
<OptimizeReferences>true</OptimizeReferences> <OptimizeReferences>true</OptimizeReferences>
<AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories> <AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
<SubSystem>NotSet</SubSystem>
</Link> </Link>
</ItemDefinitionGroup> </ItemDefinitionGroup>
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'"> <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'">
@@ -123,11 +129,13 @@
<ConformanceMode>true</ConformanceMode> <ConformanceMode>true</ConformanceMode>
<AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories> <AdditionalIncludeDirectories>$(ProjectDir)\include</AdditionalIncludeDirectories>
<DisableSpecificWarnings>4996;</DisableSpecificWarnings> <DisableSpecificWarnings>4996;</DisableSpecificWarnings>
<PreprocessorDefinitions>_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
</ClCompile> </ClCompile>
<Link> <Link>
<EnableCOMDATFolding>true</EnableCOMDATFolding> <EnableCOMDATFolding>true</EnableCOMDATFolding>
<OptimizeReferences>true</OptimizeReferences> <OptimizeReferences>true</OptimizeReferences>
<AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories> <AdditionalLibraryDirectories>$(ProjectDir)\lib;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
<SubSystem>NotSet</SubSystem>
</Link> </Link>
</ItemDefinitionGroup> </ItemDefinitionGroup>
<ItemGroup> <ItemGroup>
OpenWindow/kernel_sources.cpp
+13 -43
View File
@@ -1,46 +1,16 @@
#include "kernels.h" #include "kernels.h"
const char* matrix_mul_kernel_source =
"__kernel void mmul ( \n"
" __global float* A, \n"
" __global float* B, \n"
" __global float* C) \n"
"{ \n"
" int k; \n"
" int i = get_global_id(0); \n"
" int j = get_global_id(1); \n"
" float tmp; \n"
" if( (i < 4) && (j < 4) ) \n"
" { \n"
" tmp = 0.0; \n"
" for(k = 0; k < 4; k++) \n"
" tmp += A[i*4+k] * B[k*4+j]; \n"
" C[i*4+j] = tmp; \n"
" } \n"
"} \n" ;
const char* vertex_shader_kernel_source = const char* vertex_shader_kernel_source =
"__kernel \n" "__kernel \n"
"void vertex_shader( __global float* z_matrix, \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 index = 3 * get_global_id(0); \n" " int local_index = get_local_id(0); \n"
" float new_x = 0; \n" " int global_index = get_group_id(0); \n"
" float new_y = 0; \n" " NewVertexBuffer[4*global_index+local_index] = \n"
" float new_z = 0; \n" " m[local_index*4]*VertexBuffer[3*global_index] \n"
" float new_w = 0; \n" " + m[local_index*4 + 1]*VertexBuffer[3*global_index+1] \n"
" for(int i = 0; i < 3; i++) { \n" " + m[local_index*4 + 2]*VertexBuffer[3*global_index+2] \n"
" new_x = new_x + VertexBuffer[index+i]*z_matrix[0*4 + i]; \n" " + m[local_index*4 + 3]; \n"
" new_y = new_y + VertexBuffer[index+i]*z_matrix[1*4 + i]; \n" "} \n";
" new_z = new_z + VertexBuffer[index+i]*z_matrix[2*4 + i]; \n"
" new_w = new_w + VertexBuffer[index+i]*z_matrix[3*4 + i]; \n"
" } \n"
" new_x = new_x + z_matrix[0*4 +3]; \n"
" new_y = new_y + z_matrix[1*4 +3]; \n"
" new_z = new_z + z_matrix[2*4 +3]; \n"
" new_w = new_w + z_matrix[3*4 +3]; \n"
" NewVertexBuffer[index+0] = new_x / new_w; \n"
" NewVertexBuffer[index+1] = new_y / new_w; \n"
" NewVertexBuffer[index+2] = new_z / new_w; \n"
"} \n";
OpenWindow/kernels.cpp
+14 -64
View File
@@ -11,12 +11,7 @@ cl_context context;
// Matrix Multiplication Variables // Matrix Multiplication Variables
#define ORDER 4
#define MATRIX_SIZE 16 #define MATRIX_SIZE 16
int Mdim, Pdim, Ndim;
cl_mem mat_a;
cl_mem mat_b;
cl_mem mat_c;
// vertex shader // vertex shader
cl_mem mat_z; cl_mem mat_z;
@@ -27,15 +22,10 @@ cl_mem new_vertices_mem;
cl_platform_id* platforms = NULL; cl_platform_id* platforms = NULL;
cl_device_id* devices = NULL; cl_device_id* devices = NULL;
cl_program mat_mul_prog;
cl_program vertex_shader_prog; cl_program vertex_shader_prog;
cl_kernel mat_mul_kernel;
cl_kernel vertex_shader_kernel; cl_kernel vertex_shader_kernel;
cl_command_queue commands; cl_command_queue commands;
size_t mat_mul_global[DIM];
size_t mat_mul_local[DIM];
int err; int err;
@@ -53,78 +43,43 @@ void init_kernels() {
commands = clCreateCommandQueue(context, devices[0], 0, &err); commands = clCreateCommandQueue(context, devices[0], 0, &err);
mat_a = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * MATRIX_SIZE, NULL, &err);
mat_b = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * MATRIX_SIZE, NULL, &err);
mat_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * MATRIX_SIZE, NULL, &err);
mat_mul_prog = clCreateProgramWithSource(context, 1, (const char **)&matrix_mul_kernel_source, NULL, &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);
err = clBuildProgram(mat_mul_prog, 1, devices, NULL, NULL, NULL);
err = clBuildProgram(vertex_shader_prog, 1, devices, NULL, NULL, NULL); err = clBuildProgram(vertex_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];
OutputDebugString("Error: Failed to build program executable!\n"); printf("Error: Failed to build program executable!\n");
clGetProgramBuildInfo(mat_mul_prog, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); //clGetProgramBuildInfo(mat_mul_prog, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
OutputDebugString(buffer); printf(buffer);
return; return;
} }
mat_mul_kernel = clCreateKernel(mat_mul_prog, "mmul", &err);
vertex_shader_kernel = clCreateKernel(vertex_shader_prog, "vertex_shader", &err); vertex_shader_kernel = clCreateKernel(vertex_shader_prog, "vertex_shader", &err);
err = 0; err = 0;
err |= clSetKernelArg(mat_mul_kernel, 0, sizeof(cl_mem), &mat_a);
err |= clSetKernelArg(mat_mul_kernel, 1, sizeof(cl_mem), &mat_b);
err |= clSetKernelArg(mat_mul_kernel, 2, sizeof(cl_mem), &mat_c);
} }
void mat4_mul(float* A, float* B, float* C) void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices) {
{ mat_z = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * MATRIX_SIZE , NULL, &err);
err = clEnqueueWriteBuffer(commands, mat_a, CL_TRUE, 0, sizeof(float) * MATRIX_SIZE, A, 0, NULL, NULL); vertices_mem = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(float) * vertex_count * 3, NULL, &err);
err = clEnqueueWriteBuffer(commands, mat_b, CL_TRUE, 0, sizeof(float) * MATRIX_SIZE, B, 0, NULL, NULL); new_vertices_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * vertex_count * 4, NULL, &err);
mat_mul_global[0] = ORDER;
mat_mul_global[1] = ORDER;
err = clEnqueueNDRangeKernel(commands, mat_mul_kernel, 2, NULL, mat_mul_global, NULL, 0, NULL, NULL);
clFlush(commands);
clFinish(commands);
err = clEnqueueReadBuffer(commands, mat_c, CL_TRUE, 0, sizeof(float) * MATRIX_SIZE, C, 0, NULL, NULL);
}
void vertex_shader(Matrix* ViewPort, Matrix* Projection, Matrix* ModelView, Matrix* ModelTransform, float* vertices, int vertex_count, float* new_vertices) {
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(float) * vertex_count * 3, NULL, &err);
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);
err = clEnqueueWriteBuffer(commands, mat_z , CL_TRUE, 0, sizeof(float) * MATRIX_SIZE, z, 0, NULL, NULL);
Matrix intermediate_z;
{ // Calculate the Z to be multiplied by each vertex
mat4_mul((float*)ViewPort , (float*)Projection , (float*) &intermediate_z);
mat4_mul((float*)&intermediate_z, (float*)ModelView , (float*) &intermediate_z);
mat4_mul((float*)&intermediate_z, (float*)ModelTransform, (float*) &intermediate_z);
}
err = clEnqueueWriteBuffer(commands, mat_z, CL_TRUE, 0, sizeof(float) * MATRIX_SIZE, (float*)&intermediate_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 }; size_t vertex_shader_global[] = { vertex_count*4 };
size_t vertex_shader_local[] = { 16 }; size_t vertex_shader_local[] = { 4 };
err = clEnqueueNDRangeKernel(commands, vertex_shader_kernel, 1, NULL, vertex_shader_global, NULL, 0, NULL, NULL); err = clEnqueueNDRangeKernel(commands, vertex_shader_kernel, 1, NULL, vertex_shader_global, vertex_shader_local, 0, NULL, NULL);
err = clEnqueueReadBuffer(commands, new_vertices_mem, CL_TRUE, 0, sizeof(float) * vertex_count * 3, new_vertices, 0, NULL, NULL); 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(mat_z);
clReleaseMemObject(vertices_mem); clReleaseMemObject(vertices_mem);
@@ -133,13 +88,8 @@ void vertex_shader(Matrix* ViewPort, Matrix* Projection, Matrix* ModelView, Matr
void destroy_kernels() void destroy_kernels()
{ {
clReleaseProgram(mat_mul_prog);
clReleaseProgram(vertex_shader_prog); clReleaseProgram(vertex_shader_prog);
clReleaseKernel(mat_mul_kernel);
clReleaseKernel(vertex_shader_kernel); clReleaseKernel(vertex_shader_kernel);
clReleaseMemObject(mat_a);
clReleaseMemObject(mat_b);
clReleaseMemObject(mat_c);
// clReleaseMemObject(mat_z); // clReleaseMemObject(mat_z);
// clReleaseMemObject(vertices_mem); // clReleaseMemObject(vertices_mem);
// clReleaseMemObject(new_vertices_mem); // clReleaseMemObject(new_vertices_mem);
OpenWindow/kernels.h
+1 -1
View File
@@ -39,7 +39,7 @@ extern const char* vertex_shader_kernel_source;
void init_kernels(); void init_kernels();
void mat4_mul(float* A, float* B, float* C); void mat4_mul(float* A, float* B, float* C);
void vertex_shader(Matrix* ViewPort, Matrix* Projection, Matrix* ModelView, Matrix* ModelTransform, float* vertices, int vertex_count, float* new_vertices); void vertex_shader(float* z, float* vertices, int vertex_count, float* new_vertices);
void destroy_kernels(); void destroy_kernels();
OpenWindow/main.cpp
+6
View File
@@ -21,6 +21,12 @@ int WINAPI WinMain(HINSTANCE hInstance, HINSTANCE hPrevInstance, LPSTR lpCmdLine
hwnd = create_window(hInstance); hwnd = create_window(hInstance);
ShowCursor(false); ShowCursor(false);
AllocConsole();
freopen("CONIN$", "r",stdin);
freopen("CONOUT$", "w",stdout);
freopen("CONOUT$", "w",stderr);
ShowWindow(hwnd, nCmdShow); ShowWindow(hwnd, nCmdShow);
init_camera(); init_camera();
OpenWindow/renderer.cpp
+9 -18
View File
@@ -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 <ctime>
#pragma comment (lib, "x86_64/opencl.lib") #pragma comment (lib, "x86_64/opencl.lib")
@@ -34,6 +35,8 @@ 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());
void init_camera() { void init_camera() {
camera.SetPosition(DEFAULT_CAMERA_POS); camera.SetPosition(DEFAULT_CAMERA_POS);
camera.SetRotation(DEFAULT_CAMERA_ROT); camera.SetRotation(DEFAULT_CAMERA_ROT);
@@ -98,28 +101,20 @@ void render()
} }
{ {
model->rotate(Vec3f(0, 0, 90)); //model->rotate(Vec3f(0, 0, 90));
model->ApplyTransform(); //model->ApplyTransform();
} }
//Matrix z = ViewPort * Projection * ModelView * model->Transform;
clear_zbuffer(); clear_zbuffer();
TextureShader shader; TextureShader shader;
shader.uniform_m = (Projection); shader.uniform_m = (Projection);
shader.uniform_mit = (Projection).invert_transpose(); shader.uniform_mit = (Projection).invert_transpose();
//Matrix intermediate_z; Matrix z = ViewPort * Projection * ModelView * model->Transform;
//mat4_mul((float*) &ViewPort, (float*) &Projection, (float*) &intermediate_z);
//mat4_mul((float*)&intermediate_z, (float*)&ModelView, (float*) &intermediate_z);
//mat4_mul((float*)&intermediate_z, (float*)&model->Transform, (float*) &intermediate_z);
//shader.z = intermediate_z;
//std::vector<Vec3f> new_verts = model->verts_; // Vertex Shader: Should be called per model
Vec3f* new_verts = (Vec3f*)malloc(3 * sizeof(float) * model->nverts()); vertex_shader((float*)&z, *(float**)((Vec3f*) &model->verts_), model->nverts(), new_verts);
vertex_shader(&ViewPort, &Projection, &ModelView, &model->Transform, (float*)&model->verts_, model->nverts(), (float*)new_verts);
#pragma omp parallel for #pragma omp parallel for
for (int i = 0; i < model->nfaces(); i++) { for (int i = 0; i < model->nfaces(); i++) {
@@ -127,9 +122,7 @@ void render()
bool out = true; bool out = true;
#pragma omp parallel for #pragma omp parallel for
for (int j = 0; j < 3; j++) { for (int j = 0; j < 3; j++) {
//screen_coords[j] = model->vert(i, j); screen_coords[j] = ((Vec4f*)new_verts)[model->faces_[i][j][0]];
//screen_coords[j] = shader.vertex(i, j);
screen_coords[j] = new_verts[model->faces_[i][j][0]];
Vec3f screen3(screen_coords[j]); Vec3f screen3(screen_coords[j]);
shader.varying_uv_coords.set_col(j, model->uv(i, j)); shader.varying_uv_coords.set_col(j, model->uv(i, j));
@@ -138,7 +131,5 @@ void render()
if(!out) if(!out)
triangle(screen_coords, shader); triangle(screen_coords, shader);
} }
free(new_verts);
} }