diff --git a/PROJ3_WIN/565Rasterizer.sdf b/PROJ3_WIN/565Rasterizer.sdf new file mode 100644 index 0000000..c9917b6 Binary files /dev/null and b/PROJ3_WIN/565Rasterizer.sdf differ diff --git a/PROJ3_WIN/565Rasterizer/565Rasterizer.vcxproj b/PROJ3_WIN/565Rasterizer/565Rasterizer.vcxproj index 1077f39..816c9c4 100755 --- a/PROJ3_WIN/565Rasterizer/565Rasterizer.vcxproj +++ b/PROJ3_WIN/565Rasterizer/565Rasterizer.vcxproj @@ -1,120 +1,123 @@ - - - - - Debug - Win32 - - - Release - Win32 - - - - - - - - - - - - - - - - - - - - - - - - {FF21CA49-522E-4E86-B508-EE515B248FC4} - Win32Proj - 565Rasterizer - 565Rasterizer - - - - Application - true - Unicode - - - Application - false - true - Unicode - - - - - - - - - - - - - - true - - - false - - - - - - Level3 - Disabled - WIN32;_DEBUG;_CONSOLE;%(PreprocessorDefinitions) - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) - - - Console - true - ../shared/glew/lib;../shared/freeglut/lib;%(AdditionalLibraryDirectories) - cudart.lib; glew32.lib;glu32.lib;opengl32.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies) - mainCRTStartup - - - - - $(ProjectDir)$(Platform)/$(Configuration)/%(Filename)%(Extension).obj - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.0/C/common/inc;../shared/glew/includes;../shared/freeglut/includes - - - - - Level3 - - - MaxSpeed - true - true - WIN32;NDEBUG;_CONSOLE;%(PreprocessorDefinitions) - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) - - - Console - true - true - true - ../shared/glew/lib;../shared/freeglut/lib;%(AdditionalLibraryDirectories) - cudart.lib; glew32.lib;glu32.lib;opengl32.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies) - mainCRTStartup - - - $(ProjectDir)$(Platform)/$(Configuration)/%(Filename)%(Extension).obj - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.0/C/common/inc;../shared/glew/includes;../shared/freeglut/includes - - - - - - + + + + + Debug + Win32 + + + Release + Win32 + + + + + + + + + + + + + + + + + + + + + + + + + {FF21CA49-522E-4E86-B508-EE515B248FC4} + Win32Proj + 565Rasterizer + 565Rasterizer + + + + Application + true + Unicode + + + Application + false + true + Unicode + + + + + + + + + + + + + + true + + + false + + + + + + Level3 + Disabled + WIN32;_DEBUG;_CONSOLE;%(PreprocessorDefinitions) + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) + + + Console + true + ../shared/glew/lib;../shared/freeglut/lib;%(AdditionalLibraryDirectories) + cudart.lib; glew32.lib;glu32.lib;opengl32.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies) + mainCRTStartup + + + + + $(ProjectDir)$(Platform)/$(Configuration)/%(Filename)%(Extension).obj + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.0/C/common/inc;../shared/glew/includes;../shared/freeglut/includes + compute_20,sm_20 + + + + + Level3 + + + MaxSpeed + true + true + WIN32;NDEBUG;_CONSOLE;%(PreprocessorDefinitions) + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) + + + Console + true + true + true + ../shared/glew/lib;../shared/freeglut/lib;%(AdditionalLibraryDirectories) + cudart.lib; glew32.lib;glu32.lib;opengl32.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies) + mainCRTStartup + + + $(ProjectDir)$(Platform)/$(Configuration)/%(Filename)%(Extension).obj + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.0/C/common/inc;../shared/glew/includes;../shared/freeglut/includes + compute_20,sm_20 + + + + + + \ No newline at end of file diff --git a/PROJ3_WIN/565Rasterizer/565Rasterizer.vcxproj.filters b/PROJ3_WIN/565Rasterizer/565Rasterizer.vcxproj.filters index 5beb7fa..508311f 100755 --- a/PROJ3_WIN/565Rasterizer/565Rasterizer.vcxproj.filters +++ b/PROJ3_WIN/565Rasterizer/565Rasterizer.vcxproj.filters @@ -1,36 +1,37 @@ - - - - - - - - - - - ObjCore - - - ObjCore - - - - - - - - ObjCore - - - ObjCore - - - - - - - - {c37932d5-6627-43c3-9d6b-84866740745c} - - + + + + + + + + + + + ObjCore + + + ObjCore + + + + + + + + + ObjCore + + + ObjCore + + + + + + + + {c37932d5-6627-43c3-9d6b-84866740745c} + + \ No newline at end of file diff --git a/README.md b/README.md index dd7d45a..1664cbe 100644 --- a/README.md +++ b/README.md @@ -1,10 +1,47 @@ -------------------------------------------------------------------------------- +------------------------------------------------------------------------------- CIS565: Project 3: CUDA Rasterizer ------------------------------------------------------------------------------- Fall 2012 ------------------------------------------------------------------------------- Due Monday 11/05/2012 ------------------------------------------------------------------------------- +Blog: http://cudarasterizer.blogspot.com/ + +The results below are on the NVidia 650M card +Implementation Details: +The following have been implemented: +* Back Face Culling +* Stencil buffer +* Z-Buffer visual +* Anti-alias +* Camera interaction using mouse +* Correct color interpretetion between points +* Interpolation of normals (Geometry Shader, but done as part of the rasterizer itself) +* Specular and diffuse surfaces +* Multiple directional lights + +** Back Face Culling +This is implemented as part of the rasterizer, which ignores the faces facing away from the viewer using the normal data. For flat shading of a scaled up, rotated version of the cow I get the same fps with or without back face culling. Just over 2700 triangles were culled in this view. + +** Stencil Buffer +This too did not speed up or slow down the rasterizer. The fragment shader works on lesser number of fragment; only thise fragments that are in the visible part of the stencil are shaded. +** Z-Buffer Visual +I got t +he minimum and maximum values of the depth and shaded from 0.15 to 0.85 depending on depth from camera with 0.15 being the closest. + +** Anti-alias +I performed antialiasing on each and every primitive. This worked fine for a single primitive, but on a model, I get lines at the junctions where the primitives meet. This was done in the vertex shader itself. The image is anti-aliased but artifacts are disturbing. This runs at 12 fps same as the normal case. +So I performed anti aliasing using higher resolution (4 times the normal resolution) in the rasterizer and performed shading on the higher resolution before writing to the frame buffer. I get a 0.5 to 1 fps with this method instead of 12 fps. + +** Camera interaction using mouse +Use the left mouse button and drag to rotate the camera. Press down the "ALT" key and the left mouse down and then mouse the mouse for zooming in and out. + +** Color interpretation did not slow down the rasterizer + +** Shading the surfaces as well as lights also did not have much impact on the speed + + +All the options can be enabled by setting bool values in the cuda Rasterizer core function. ------------------------------------------------------------------------------- NOTE: @@ -13,7 +50,7 @@ This project requires an NVIDIA graphics card with CUDA capability! Any card wit ------------------------------------------------------------------------------- INTRODUCTION: -------------------------------------------------------------------------------- +--------------zbuffer----------------------------------------------------------------- In this project, you will implement a simplified CUDA based implementation of a standard rasterized graphics pipeline, similar to the OpenGL pipeline. In this project, you will implement vertex shading, primitive assembly, perspective transformation, rasterization, fragment shading, and write the resulting fragments to a framebuffer. More information about the rasterized graphics pipeline can be found in the 10/15 class slides and in your notes from CIS560. The basecode provided includes an OBJ loader and much of the mundane I/O and bookkeeping code. The basecode also includes some functions that you may find useful, described below. The core rasterization pipeline is left for you to implement. @@ -71,7 +108,7 @@ IMPORTANT: For each of these stages implemented, you must also add a section to * Correct color interpretation between points on a primitive * Texture mapping WITH texture filtering and perspective correct texture coordinates -* Support for additional primitices. Each one of these can count as HALF of a feature. +* Support for additional primitives. Each one of these can count as HALF of a feature. * Lines * Line strips * Triangle fans diff --git a/objs/cube.obj b/objs/cube.obj new file mode 100644 index 0000000..eba9294 --- /dev/null +++ b/objs/cube.obj @@ -0,0 +1,33 @@ +# cube.obj +# + +g cube + +v 0.0 0.0 0.0 +v 0.0 0.0 1.0 +v 0.0 1.0 0.0 +v 0.0 1.0 1.0 +v 1.0 0.0 0.0 +v 1.0 0.0 1.0 +v 1.0 1.0 0.0 +v 1.0 1.0 1.0 + +vn 0.0 0.0 1.0 +vn 0.0 0.0 -1.0 +vn 0.0 1.0 0.0 +vn 0.0 -1.0 0.0 +vn 1.0 0.0 0.0 +vn -1.0 0.0 0.0 + +f 1//2 7//2 5//2 +f 1//2 3//2 7//2 +f 1//6 4//6 3//6 +f 1//6 2//6 4//6 +f 3//3 8//3 7//3 +f 3//3 4//3 8//3 +f 5//5 7//5 8//5 +f 5//5 8//5 6//5 +f 1//4 5//4 6//4 +f 1//4 6//4 2//4 +f 2//1 6//1 8//1 +f 2//1 8//1 4//1 diff --git a/objs/cubeTemp.obj b/objs/cubeTemp.obj new file mode 100644 index 0000000..5e7c518 --- /dev/null +++ b/objs/cubeTemp.obj @@ -0,0 +1,24 @@ +# cube.obj +# + +g cube + +v 0.0 0.0 0.0 +v 0.0 0.0 1.0 +v 0.0 1.0 0.0 +v 0.0 1.0 1.0 +v 1.0 0.0 0.0 +v 1.0 0.0 1.0 +v 1.0 1.0 0.0 +v 1.0 1.0 1.0 + +vn 0.0 0.0 1.0 +vn 0.0 0.0 -1.0 +vn 0.0 1.0 0.0 +vn 0.0 -1.0 0.0 +vn 1.0 0.0 0.0 +vn -1.0 0.0 0.0 + +f 2//1 8//1 4//1 + + diff --git a/renders/Antialias.PNG b/renders/Antialias.PNG new file mode 100644 index 0000000..db28fe2 Binary files /dev/null and b/renders/Antialias.PNG differ diff --git a/renders/Antialias_1.PNG b/renders/Antialias_1.PNG new file mode 100644 index 0000000..53cd17c Binary files /dev/null and b/renders/Antialias_1.PNG differ diff --git a/renders/Antialias_Primitives.PNG b/renders/Antialias_Primitives.PNG new file mode 100644 index 0000000..df839fe Binary files /dev/null and b/renders/Antialias_Primitives.PNG differ diff --git a/renders/Cow_1.PNG b/renders/Cow_1.PNG new file mode 100644 index 0000000..2a6345f Binary files /dev/null and b/renders/Cow_1.PNG differ diff --git a/renders/Cow_Antialias.PNG b/renders/Cow_Antialias.PNG new file mode 100644 index 0000000..3b88713 Binary files /dev/null and b/renders/Cow_Antialias.PNG differ diff --git a/renders/Cow_Antialias_Compare.PNG b/renders/Cow_Antialias_Compare.PNG new file mode 100644 index 0000000..15844fd Binary files /dev/null and b/renders/Cow_Antialias_Compare.PNG differ diff --git a/renders/Cow_Antialias_Depth.PNG b/renders/Cow_Antialias_Depth.PNG new file mode 100644 index 0000000..3b88713 Binary files /dev/null and b/renders/Cow_Antialias_Depth.PNG differ diff --git a/renders/Cow_Depth.PNG b/renders/Cow_Depth.PNG new file mode 100644 index 0000000..62718b8 Binary files /dev/null and b/renders/Cow_Depth.PNG differ diff --git a/renders/Cow_Diffuse.PNG b/renders/Cow_Diffuse.PNG new file mode 100644 index 0000000..9943a6c Binary files /dev/null and b/renders/Cow_Diffuse.PNG differ diff --git a/renders/Cow_Flat.PNG b/renders/Cow_Flat.PNG new file mode 100644 index 0000000..5f6fead Binary files /dev/null and b/renders/Cow_Flat.PNG differ diff --git a/renders/Cow_MultipleLights.PNG b/renders/Cow_MultipleLights.PNG new file mode 100644 index 0000000..dd19cb3 Binary files /dev/null and b/renders/Cow_MultipleLights.PNG differ diff --git a/renders/Cow_MultipleLights_1.PNG b/renders/Cow_MultipleLights_1.PNG new file mode 100644 index 0000000..d1aa319 Binary files /dev/null and b/renders/Cow_MultipleLights_1.PNG differ diff --git a/renders/Cow_PrimitiveAntialias.PNG b/renders/Cow_PrimitiveAntialias.PNG new file mode 100644 index 0000000..739ce68 Binary files /dev/null and b/renders/Cow_PrimitiveAntialias.PNG differ diff --git a/renders/Cow_Specular.PNG b/renders/Cow_Specular.PNG new file mode 100644 index 0000000..3598915 Binary files /dev/null and b/renders/Cow_Specular.PNG differ diff --git a/renders/Cow_Stencil.PNG b/renders/Cow_Stencil.PNG new file mode 100644 index 0000000..3b6080a Binary files /dev/null and b/renders/Cow_Stencil.PNG differ diff --git a/renders/Cube_1.PNG b/renders/Cube_1.PNG new file mode 100644 index 0000000..8727ca6 Binary files /dev/null and b/renders/Cube_1.PNG differ diff --git a/renders/Primitive_Antialias_Zoomed.PNG b/renders/Primitive_Antialias_Zoomed.PNG new file mode 100644 index 0000000..a15e42f Binary files /dev/null and b/renders/Primitive_Antialias_Zoomed.PNG differ diff --git a/src/Eye.h b/src/Eye.h new file mode 100644 index 0000000..64ad2ce --- /dev/null +++ b/src/Eye.h @@ -0,0 +1,105 @@ +#ifndef EYE_H +#define EYE_H + +#include "glm/glm.hpp" +#include "cudaMat4.h" +#include "utilities.h" +#include "glm/gtc/matrix_transform.hpp" +#include "glm/gtc/matrix_inverse.hpp" + +class Eye +{ +public: + glm::vec3 pos; // wrt the world + glm::vec3 rot; // wrt the world + glm::vec3 up; + glm::vec3 viewDir; + glm::vec2 fov; + glm::vec2 resolution; + + float l; // left plane + float r; // right plane + float b; // bottom plane + float t; // top plane + float n; // near plane + float f; // far plane + + glm::mat4 transformIntoPerspective; + + Eye() + { + pos = glm::vec3(0, 0, 20); + rot = glm::vec3(0,0,0); + up = glm::vec3(0, 1, 0); + viewDir = glm::vec3(0, 0, -1); + resolution = glm::vec2(800, 800); + fov = glm::vec2(30, 30); + } + + void SetResolution(glm::vec2 fResolution) + { + resolution = fResolution; + } + + void SetBoundariesOfView(float f_l, float f_r, float f_b, float f_t, float f_n, float f_f) + { + l = f_l; + r = f_r; + b = f_b; + t = f_t; + n = f_n; + f = f_f; + TransformIntoPerspectiveView(); + } + + glm::mat4 TransformWorldToEye() + { + // ModelView matrix + glm::mat4 transformWorld; + + glm::mat4 translationMat = glm::translate(glm::mat4(), -pos); + glm::mat4 rotationMat = glm::rotate(glm::mat4(), rot.x, glm::vec3(1,0,0)); + rotationMat = rotationMat*glm::rotate(glm::mat4(), rot.y, glm::vec3(0,1,0)); + rotationMat = rotationMat*glm::rotate(glm::mat4(), rot.z, glm::vec3(0,0,1)); + glm::mat4 transformWorld1 = translationMat*rotationMat; + transformWorld[0] = glm::vec4(transformWorld1[0].x, transformWorld1[1].x, transformWorld1[2].x, transformWorld1[3].x); + transformWorld[1] = glm::vec4(transformWorld1[0].y, transformWorld1[1].y, transformWorld1[2].y, transformWorld1[3].y); + transformWorld[2] = glm::vec4(transformWorld1[0].z, transformWorld1[1].z, transformWorld1[2].z, transformWorld1[3].z); + transformWorld[3] = glm::vec4(transformWorld1[0].w, transformWorld1[1].w, transformWorld1[2].w, transformWorld1[3].w); + + return transformWorld; + } + + void TransformIntoPerspectiveView() + { + transformIntoPerspective[0] = glm::vec4(2*n/(r-l), 0, (r+l)/(r-l), 0); + transformIntoPerspective[1] = glm::vec4(0, 2*n/(t-b), (t+b)/(t-b), 0); + transformIntoPerspective[2] = glm::vec4(0, 0, -(f+n)/(f-n), -2*(f*n)/(f-n)); + transformIntoPerspective[3] = glm::vec4(0, 0, -1, 0); + } + + glm::mat4 GetTransformWorldToPerspective() + { + return TransformWorldToEye()*transformIntoPerspective; + } + + void GetParametersForScreenTransform(glm::vec2 resolution, glm::vec3& M, glm::vec3& A, glm::vec3& B, float& distImagePlaneFromCamera) + { + viewDir = glm::normalize(viewDir); + up = glm::normalize(up); + + A = glm::normalize(glm::cross(viewDir, up)); + B = glm::normalize(glm::cross(A, viewDir)); + + float tanVert = tan(fov.y*PI/180); + + float camDistFromScreen = (float)((resolution.y/2.0)/tanVert); + fov.x = atan(resolution.x/(2*camDistFromScreen)) * 180/PI; + glm::vec3 C = viewDir*camDistFromScreen; + M = pos + C; + + distImagePlaneFromCamera = camDistFromScreen; + } +}; + +#endif diff --git a/src/ObjCore/obj.cpp b/src/ObjCore/obj.cpp index e748574..90f3c38 100755 --- a/src/ObjCore/obj.cpp +++ b/src/ObjCore/obj.cpp @@ -1,331 +1,375 @@ -//OBJCORE- A Obj Mesh Library by Yining Karl Li -//This file is part of OBJCORE, Coyright (c) 2012 Yining Karl Li - -#include "obj.h" -#include -#include - -#define EPSILON std::numeric_limits::epsilon() - -using namespace std; - -obj::obj(){ - vbosize = 0; - nbosize = 0; - cbosize = 0; - ibosize = 0; - top = 0; - defaultColor = glm::vec3(0,0,0); - boundingbox = new float[32]; - maxminSet = false; - xmax=0; xmin=0; ymax=0; ymin=0; zmax=0; zmin=0; - -} - -obj::~obj(){ - /*delete vbo; - delete nbo; - delete cbo; - delete ibo;*/ - delete boundingbox; - for(int i=0; i VBOvec; - vector NBOvec; - vector IBOvec; - int index = 0; - bool genNormals = false; - if(faces.size()!=facenormals.size()){ - genNormals = true; - } - for(int k = 0; k face = faces[k]; - - glm::vec4 p0 = points[face[0]]; - - for(int i=2; i facenormal = facenormals[k]; - NBOvec.push_back(normals[facenormal[0]][0]); NBOvec.push_back(normals[facenormal[0]][1]); NBOvec.push_back(normals[facenormal[0]][2]); //NBOvec.push_back(0.0f); - NBOvec.push_back(normals[facenormal[i-1]][0]); NBOvec.push_back(normals[facenormal[i-1]][1]); NBOvec.push_back(normals[facenormal[i-1]][2]); //NBOvec.push_back(0.0f); - NBOvec.push_back(normals[facenormal[i]][0]); NBOvec.push_back(normals[facenormal[i]][1]); NBOvec.push_back(normals[facenormal[i]][2]); //NBOvec.push_back(0.0f); - }else{ - - glm::vec3 a = glm::vec3(p1[0], p1[1], p1[2]) - glm::vec3(p0[0], p0[1], p0[2]); - glm::vec3 b = glm::vec3(p2[0], p2[1], p2[2]) - glm::vec3(p0[0], p0[1], p0[2]); - glm::vec3 n = glm::normalize(glm::cross(a,b)); - NBOvec.push_back(n[0]); NBOvec.push_back(n[1]); NBOvec.push_back(n[2]); //NBOvec.push_back(0.0f); - NBOvec.push_back(n[0]); NBOvec.push_back(n[1]); NBOvec.push_back(n[2]); //NBOvec.push_back(0.0f); - NBOvec.push_back(n[0]); NBOvec.push_back(n[1]); NBOvec.push_back(n[2]); //NBOvec.push_back(0.0f); - } - - IBOvec.push_back(index+0); IBOvec.push_back(index+1); IBOvec.push_back(index+2); - - index=index+3; - } - } - } - - vbo = new float[VBOvec.size()]; - nbo = new float[NBOvec.size()]; - ibo = new int[IBOvec.size()]; - vbosize = (int)VBOvec.size(); - nbosize = (int)NBOvec.size(); - ibosize = (int)IBOvec.size(); - for(int i=0; ixmax){ xmax = x; } - if(xymax){ ymax = y; } - if(yzmax){ zmax = z; } - if(z face){ - if(face.size()<=3){ - return true; - } - - int k = (int)face.size()-1; - glm::vec3 a = glm::vec3(points[face[0]][0], points[face[0]][1], points[face[0]][2]) - glm::vec3(points[face[k]][0], points[face[k]][1], points[face[k]][2]); - glm::vec3 b = glm::vec3(points[face[0]][0], points[face[0]][1], points[face[0]][2]) - glm::vec3(points[face[1]][0], points[face[1]][1], points[face[1]][2]); - glm::vec3 n = glm::normalize(glm::cross(a,b)); - - for(int i=2; i EPSILON) || (abs(m[1] - n[1]) > EPSILON) || (abs(m[2] - n[2]) > EPSILON)){ - return false; - } - } - - glm::vec3 c = glm::vec3(points[face[k]][0], points[face[k]][1], points[face[k]][2]) - glm::vec3(points[face[k-1]][0], points[face[k-1]][1], points[face[k-1]][2]); - glm::vec3 d = glm::vec3(points[face[k]][0], points[face[k]][1], points[face[k]][2]) - glm::vec3(points[face[0]][0], points[face[0]][1], points[face[0]][2]); - glm::vec3 m = glm::normalize(glm::cross(c,d)); - if((abs(m[0] - n[0]) > EPSILON) || (abs(m[1] - n[1]) > EPSILON) || (abs(m[2] - n[2]) > EPSILON)){ - return false; - } - return true; -} - -void obj::addPoint(glm::vec3 point){ - - if(points.size()==0){ - xmin = point[0]; xmax = point[0]; - ymin = point[1]; ymax = point[1]; - zmin = point[2]; zmax = point[2]; - } - - points.push_back(glm::vec4(point[0], point[1], point[2], 1)); - - compareMaxMin(point[0], point[1], point[2]); -} - -void obj::addFace(vector face){ - faces.push_back(face); - float facexmax = points[face[0]][0]; float faceymax = points[face[0]][1]; float facezmax = points[face[0]][2]; - float facexmin = points[face[0]][0]; float faceymin = points[face[0]][1]; float facezmin = points[face[0]][2]; - for(int i=0; ifacexmax){ facexmax = points[face[i]][0]; } - if(points[face[i]][0]faceymax){ faceymax = points[face[i]][1]; } - if(points[face[i]][1]facezmax){ facezmax = points[face[i]][2]; } - if(points[face[i]][2] facen){ - facenormals.push_back(facen); -} - -void obj::addFaceTexture(vector facet){ - facetextures.push_back(facet); -} - -void obj::addNormal(glm::vec3 normal){ - normals.push_back(glm::vec4(normal[0], normal[1], normal[2], 1)); -} - -void obj::addTextureCoord(glm::vec3 texcoord){ - texturecoords.push_back(glm::vec4(texcoord[0], texcoord[1], texcoord[2], 1)); -} - -float* obj::getBoundingBox(){ - return boundingbox; -} - -float obj::getTop(){ - return top; -} - -void obj::recenter(){ - glm::vec3 center = glm::vec3((xmax+xmin)/2,ymin,(zmax+zmin)/2); - xmax=points[0][0]-center[0]; xmin=points[0][0]-center[0]; - ymax=points[0][1]-center[1]; ymin=points[0][1]-center[1]; - zmax=points[0][2]-center[2]; zmin=points[0][2]-center[2]; - top=0; - for(int i=0; i face = faces[i]; - - float facexmax = points[face[0]][0]; float faceymax = points[face[0]][1]; float facezmax = points[face[0]][2]; - float facexmin = points[face[0]][0]; float faceymin = points[face[0]][1]; float facezmin = points[face[0]][2]; - - for(int j=0; jfacexmax){ facexmax = points[face[j]][0]; } - if(points[face[j]][0]faceymax){ faceymax = points[face[j]][1]; } - if(points[face[j]][1]facezmax){ facezmax = points[face[j]][2]; } - if(points[face[j]][2]* obj::getPoints(){ - return &points; -} - -vector >* obj::getFaces(){ - return &faces; -} - -vector >* obj::getFaceNormals(){ - return &facenormals; -} - -vector >* obj::getFaceTextures(){ - return &facetextures; -} - -vector* obj::getNormals(){ - return &normals; -} - -vector* obj::getTextureCoords(){ - return &texturecoords; -} - -vector* obj::getFaceBoxes(){ - return &faceboxes; -} - -glm::vec3 obj::getColor(){ - return defaultColor; -} - -float* obj::getVBO(){ - return vbo; -} - -float* obj::getCBO(){ - return cbo; -} - -float* obj::getNBO(){ - return nbo; -} - -int* obj::getIBO(){ - return ibo; -} - -int obj::getVBOsize(){ - return vbosize; -} - -int obj::getNBOsize(){ - return nbosize; -} - -int obj::getIBOsize(){ - return ibosize; -} - -int obj::getCBOsize(){ - return cbosize; -} - +//OBJCORE- A Obj Mesh Library by Yining Karl Li +//This file is part of OBJCORE, Coyright (c) 2012 Yining Karl Li + +#include "obj.h" +#include "../utilities.h" +#include +#include + +#define EPSILON std::numeric_limits::epsilon() + +using namespace std; + +obj::obj(){ + pos = glm::vec3(0,0,0); + rot = glm::vec3(0,0,0); + scale = glm::vec3(1,1,1); + vbosize = 0; + nbosize = 0; + cbosize = 0; + ibosize = 0; + top = 0; + defaultColor = glm::vec3(0,0,0); + boundingbox = new float[32]; + bounds = new float[6]; + maxminSet = false; + xmax=0; xmin=0; ymax=0; ymin=0; zmax=0; zmin=0; + +} + +obj::~obj(){ + /*delete vbo; + delete nbo; + delete cbo; + delete ibo;*/ + delete boundingbox; + delete bounds; + for(int i=0; i VBOvec; + vector NBOvec; + vector IBOvec; + int index = 0; + bool genNormals = false; + if(faces.size()!=facenormals.size()){ + genNormals = true; + } + transform = getTransform(); + glm::mat4 invTransform = transform._inverse(); // This return the transpose of the inverse for the system that is used in the code by me + for(int k = 0; k face = faces[k]; + + glm::vec4 p0 = points[face[0]]; + p0 = transform*p0; + for(int i=2; i facenormal = facenormals[k]; + glm::vec4 n0 = normals[facenormal[0]]; + n0 = transform*n0; + glm::vec4 n1 = normals[facenormal[i-1]]; + n1 = transform*n1; + glm::vec4 n2 = normals[facenormal[i]]; + n2 = transform*n2; + NBOvec.push_back(n0[0]); NBOvec.push_back(n0[1]); NBOvec.push_back(n0[2]); //NBOvec.push_back(0.0f); + NBOvec.push_back(n1[0]); NBOvec.push_back(n1[1]); NBOvec.push_back(n1[2]); //NBOvec.push_back(0.0f); + NBOvec.push_back(n2[0]); NBOvec.push_back(n2[1]); NBOvec.push_back(n2[2]); //NBOvec.push_back(0.0f); + }else{ + + glm::vec3 a = glm::vec3(p1[0], p1[1], p1[2]) - glm::vec3(p0[0], p0[1], p0[2]); + glm::vec3 b = glm::vec3(p2[0], p2[1], p2[2]) - glm::vec3(p0[0], p0[1], p0[2]); + glm::vec3 n = glm::normalize(glm::cross(a,b)); + NBOvec.push_back(n[0]); NBOvec.push_back(n[1]); NBOvec.push_back(n[2]); //NBOvec.push_back(0.0f); + NBOvec.push_back(n[0]); NBOvec.push_back(n[1]); NBOvec.push_back(n[2]); //NBOvec.push_back(0.0f); + NBOvec.push_back(n[0]); NBOvec.push_back(n[1]); NBOvec.push_back(n[2]); //NBOvec.push_back(0.0f); + } + + IBOvec.push_back(index+0); IBOvec.push_back(index+1); IBOvec.push_back(index+2); + + index=index+3; + } + } + } + + vbo = new float[VBOvec.size()]; + nbo = new float[NBOvec.size()]; + ibo = new int[IBOvec.size()]; + vbosize = (int)VBOvec.size(); + nbosize = (int)NBOvec.size(); + ibosize = (int)IBOvec.size(); + for(int i=0; ixmax){ xmax = x; } + if(xymax){ ymax = y; } + if(yzmax){ zmax = z; } + if(z face){ + if(face.size()<=3){ + return true; + } + + int k = (int)face.size()-1; + glm::vec3 a = glm::vec3(points[face[0]][0], points[face[0]][1], points[face[0]][2]) - glm::vec3(points[face[k]][0], points[face[k]][1], points[face[k]][2]); + glm::vec3 b = glm::vec3(points[face[0]][0], points[face[0]][1], points[face[0]][2]) - glm::vec3(points[face[1]][0], points[face[1]][1], points[face[1]][2]); + glm::vec3 n = glm::normalize(glm::cross(a,b)); + + for(int i=2; i EPSILON) || (abs(m[1] - n[1]) > EPSILON) || (abs(m[2] - n[2]) > EPSILON)){ + return false; + } + } + + glm::vec3 c = glm::vec3(points[face[k]][0], points[face[k]][1], points[face[k]][2]) - glm::vec3(points[face[k-1]][0], points[face[k-1]][1], points[face[k-1]][2]); + glm::vec3 d = glm::vec3(points[face[k]][0], points[face[k]][1], points[face[k]][2]) - glm::vec3(points[face[0]][0], points[face[0]][1], points[face[0]][2]); + glm::vec3 m = glm::normalize(glm::cross(c,d)); + if((abs(m[0] - n[0]) > EPSILON) || (abs(m[1] - n[1]) > EPSILON) || (abs(m[2] - n[2]) > EPSILON)){ + return false; + } + return true; +} + +void obj::addPoint(glm::vec3 point){ + + if(points.size()==0){ + xmin = point[0]; xmax = point[0]; + ymin = point[1]; ymax = point[1]; + zmin = point[2]; zmax = point[2]; + } + + points.push_back(glm::vec4(point[0], point[1], point[2], 1)); + + compareMaxMin(point[0], point[1], point[2]); +} + +void obj::addFace(vector face){ + faces.push_back(face); + float facexmax = points[face[0]][0]; float faceymax = points[face[0]][1]; float facezmax = points[face[0]][2]; + float facexmin = points[face[0]][0]; float faceymin = points[face[0]][1]; float facezmin = points[face[0]][2]; + for(int i=0; ifacexmax){ facexmax = points[face[i]][0]; } + if(points[face[i]][0]faceymax){ faceymax = points[face[i]][1]; } + if(points[face[i]][1]facezmax){ facezmax = points[face[i]][2]; } + if(points[face[i]][2] facen){ + facenormals.push_back(facen); +} + +void obj::addFaceTexture(vector facet){ + facetextures.push_back(facet); +} + +void obj::addNormal(glm::vec3 normal){ + normals.push_back(glm::vec4(normal[0], normal[1], normal[2], 0)); +} + +void obj::addTextureCoord(glm::vec3 texcoord){ + texturecoords.push_back(glm::vec4(texcoord[0], texcoord[1], texcoord[2], 1)); +} + +float* obj::getBoundingBox(){ + return boundingbox; +} + +float* obj::getBounds(){ + bounds[0] = boundingbox[16]; + bounds[1] = boundingbox[17]; + bounds[2] = boundingbox[18]; + + bounds[3] = boundingbox[8]; + bounds[4] = boundingbox[9]; + bounds[5] = boundingbox[10]; + return bounds; +} + +float obj::getTop(){ + return top; +} + +void obj::recenter(){ + glm::vec3 center = glm::vec3((xmax+xmin)/2,ymin,(zmax+zmin)/2); + xmax=points[0][0]-center[0]; xmin=points[0][0]-center[0]; + ymax=points[0][1]-center[1]; ymin=points[0][1]-center[1]; + zmax=points[0][2]-center[2]; zmin=points[0][2]-center[2]; + top=0; + for(int i=0; i face = faces[i]; + + float facexmax = points[face[0]][0]; float faceymax = points[face[0]][1]; float facezmax = points[face[0]][2]; + float facexmin = points[face[0]][0]; float faceymin = points[face[0]][1]; float facezmin = points[face[0]][2]; + + for(int j=0; jfacexmax){ facexmax = points[face[j]][0]; } + if(points[face[j]][0]faceymax){ faceymax = points[face[j]][1]; } + if(points[face[j]][1]facezmax){ facezmax = points[face[j]][2]; } + if(points[face[j]][2]* obj::getPoints(){ + return &points; +} + +vector >* obj::getFaces(){ + return &faces; +} + +vector >* obj::getFaceNormals(){ + return &facenormals; +} + +vector >* obj::getFaceTextures(){ + return &facetextures; +} + +vector* obj::getNormals(){ + return &normals; +} + +vector* obj::getTextureCoords(){ + return &texturecoords; +} + +vector* obj::getFaceBoxes(){ + return &faceboxes; +} + +glm::vec3 obj::getColor(){ + return defaultColor; +} + +float* obj::getVBO(){ + return vbo; +} + +float* obj::getCBO(){ + return cbo; +} + +float* obj::getNBO(){ + return nbo; +} + +int* obj::getIBO(){ + return ibo; +} + +int obj::getVBOsize(){ + return vbosize; +} + +int obj::getNBOsize(){ + return nbosize; +} + +int obj::getIBOsize(){ + return ibosize; +} + +int obj::getCBOsize(){ + return cbosize; +} + +glm::mat4 obj::getTransform() +{ + transform = utilityCore::buildTransformationMatrix(pos, rot, scale); + return transform; +} + +cudaMat4 obj::getTransformCuda() +{ + return utilityCore::glmMat4ToCudaMat4(getTransform()); +} + +void obj::setTransforms(glm::vec3 v_pos, glm::vec3 v_rot, glm::vec3 v_scale) +{ + pos = v_pos; + rot = v_rot; + scale = v_scale; +} \ No newline at end of file diff --git a/src/ObjCore/obj.h b/src/ObjCore/obj.h index 20b28f9..fd7c9d7 100755 --- a/src/ObjCore/obj.h +++ b/src/ObjCore/obj.h @@ -1,77 +1,89 @@ -//OBJCORE- A Obj Mesh Library by Yining Karl Li -//This file is part of OBJCORE, Coyright (c) 2012 Yining Karl Li - -#ifndef OBJ -#define OBJ - -#include "../glm/glm.hpp" -#include -#include - -using namespace std; - -class obj{ -private: - vector points; - vector > faces; - vector > facenormals; - vector > facetextures; - vector faceboxes; //bounding boxes for each face are stored in vbo-format! - vector normals; - vector texturecoords; - int vbosize; - int nbosize; - int cbosize; - int ibosize; - float* vbo; - float* nbo; - float* cbo; - int* ibo; - float* boundingbox; - float top; - glm::vec3 defaultColor; - float xmax; float xmin; float ymax; float ymin; float zmax; float zmin; - bool maxminSet; -public: - obj(); - ~obj(); - - //------------------------------- - //-------Mesh Operations--------- - //------------------------------- - void buildVBOs(); - void addPoint(glm::vec3); - void addFace(vector); - void addNormal(glm::vec3); - void addTextureCoord(glm::vec3); - void addFaceNormal(vector); - void addFaceTexture(vector); - void compareMaxMin(float, float, float); - bool isConvex(vector); - void recenter(); - - //------------------------------- - //-------Get/Set Operations------ - //------------------------------- - float* getBoundingBox(); //returns vbo-formatted bounding box - float getTop(); - void setColor(glm::vec3); - glm::vec3 getColor(); - float* getVBO(); - float* getCBO(); - float* getNBO(); - int* getIBO(); - int getVBOsize(); - int getNBOsize(); - int getIBOsize(); - int getCBOsize(); - vector* getPoints(); - vector >* getFaces(); - vector >* getFaceNormals(); - vector >* getFaceTextures(); - vector* getNormals(); - vector* getTextureCoords(); - vector* getFaceBoxes(); -}; - -#endif \ No newline at end of file +//OBJCORE- A Obj Mesh Library by Yining Karl Li +//This file is part of OBJCORE, Coyright (c) 2012 Yining Karl Li + +#ifndef OBJ +#define OBJ + +#include "../glm/glm.hpp" +#include "../cudaMat4.h" +#include +#include + +using namespace std; + +class obj{ +private: + glm::vec3 pos; + glm::vec3 rot; + glm::vec3 scale; + + vector points; + vector > faces; + vector > facenormals; + vector > facetextures; + vector faceboxes; //bounding boxes for each face are stored in vbo-format! + vector normals; + vector texturecoords; + int vbosize; + int nbosize; + int cbosize; + int ibosize; + float* vbo; + float* nbo; + float* cbo; + int* ibo; + float* boundingbox; + float* bounds; + float top; + glm::vec3 defaultColor; + float xmax; float xmin; float ymax; float ymin; float zmax; float zmin; + bool maxminSet; + glm::mat4 transform; +public: + obj(); + ~obj(); + + //------------------------------- + //-------Mesh Operations--------- + //------------------------------- + void buildVBOs(); + void addPoint(glm::vec3); + void addFace(vector); + void addNormal(glm::vec3); + void addTextureCoord(glm::vec3); + void addFaceNormal(vector); + void addFaceTexture(vector); + void compareMaxMin(float, float, float); + bool isConvex(vector); + void recenter(); + + //------------------------------- + //-------Get/Set Operations------ + //------------------------------- + float* getBoundingBox(); //returns vbo-formatted bounding box + float* getBounds(); + float getTop(); + void setColor(glm::vec3); + glm::vec3 getColor(); + float* getVBO(); + float* getCBO(); + float* getNBO(); + int* getIBO(); + int getVBOsize(); + int getNBOsize(); + int getIBOsize(); + int getCBOsize(); + vector* getPoints(); + vector >* getFaces(); + vector >* getFaceNormals(); + vector >* getFaceTextures(); + vector* getNormals(); + vector* getTextureCoords(); + vector* getFaceBoxes(); + glm::mat4 getTransform(); + cudaMat4 getTransformCuda(); + + void setTransforms(glm::vec3 v_pos, glm::vec3 v_rot, glm::vec3 v_scale); +}; + +#endif diff --git a/src/main.cpp b/src/main.cpp index dfb689a..4658dae 100755 --- a/src/main.cpp +++ b/src/main.cpp @@ -3,6 +3,10 @@ #include "main.h" +int theButtonState = 0; +int theModifierState = 0; +int lastX = 0, lastY = 0; + //------------------------------- //-------------MAIN-------------- //------------------------------- @@ -18,6 +22,10 @@ int main(int argc, char** argv){ //renderScene = new scene(data); mesh = new obj(); objLoader* loader = new objLoader(data, mesh); + glm::vec3 meshPos = glm::vec3(0,0,0); + glm::vec3 meshRot = glm::vec3(0,-45,0); + glm::vec3 meshScale = glm::vec3(1.5,1.5,1.5); + mesh->setTransforms(meshPos, meshRot, meshScale); mesh->buildVBOs(); delete loader; loadedScene = true; @@ -29,6 +37,9 @@ int main(int argc, char** argv){ return 0; } + // Get eye details + eye.SetBoundariesOfView(-1, 1, -1, 1, 20, 52); + frame = 0; seconds = time (NULL); fpstracker = 0; @@ -71,7 +82,8 @@ int main(int argc, char** argv){ #else glutDisplayFunc(display); glutKeyboardFunc(keyboard); - + glutMouseFunc(onMouseCb); + glutMotionFunc(onMouseMotionCb); glutMainLoop(); #endif kernelCleanup(); @@ -90,17 +102,34 @@ void runCuda(){ vbo = mesh->getVBO(); vbosize = mesh->getVBOsize(); - float newcbo[] = {0.0, 1.0, 0.0, + /*float newcbo[] = {0.0, 1.0, 0.0, 0.0, 0.0, 1.0, - 1.0, 0.0, 0.0}; + 1.0, 0.0, 0.0};*/ + + float newcbo[] = {0.5, 0.5, 0.5, + 0.5, 0.5, 0.5, + 0.5, 0.5, 0.5}; cbo = newcbo; cbosize = 9; ibo = mesh->getIBO(); ibosize = mesh->getIBOsize(); + nbo = mesh->getNBO(); + nbosize = mesh->getNBOsize(); + + glm::mat4 modelView = eye.TransformWorldToEye(); + eye.TransformIntoPerspectiveView(); + glm::mat4 perspective = eye.transformIntoPerspective; + cudaMat4 modelViewCudaMat = utilityCore::glmMat4ToCudaMat4NoTranspose(modelView); + cudaMat4 perspectiveCudaMat = utilityCore::glmMat4ToCudaMat4NoTranspose(perspective); + + glm::mat4 finalTransform = eye.GetTransformWorldToPerspective(); + cudaMat4 finalTransformCudaMat = utilityCore::glmMat4ToCudaMat4NoTranspose(finalTransform); + cudaGLMapBufferObject((void**)&dptr, pbo); - cudaRasterizeCore(dptr, glm::vec2(width, height), frame, vbo, vbosize, cbo, cbosize, ibo, ibosize); + cudaRasterizeCore(dptr, glm::vec2(width, height), frame, vbo, vbosize, cbo, cbosize, + ibo, ibosize, nbo, nbosize, eye, finalTransformCudaMat, modelViewCudaMat, perspectiveCudaMat); cudaGLUnmapBufferObject(pbo); vbo = NULL; @@ -110,6 +139,9 @@ void runCuda(){ frame++; fpstracker++; + /*int temp; + std::cin >> temp;*/ + } #ifdef __APPLE__ @@ -150,10 +182,11 @@ void runCuda(){ void display(){ runCuda(); time_t seconds2 = time (NULL); - - if(seconds2-seconds >= 1){ + if(seconds2-seconds >= 1){ fps = fpstracker/(seconds2-seconds); + printf("Difference:%f, FPS: %f\n", ((float)seconds2-(float)seconds), (float)fpstracker/(float)(seconds2-seconds)); + fpstracker = 0; seconds = seconds2; @@ -164,7 +197,8 @@ void runCuda(){ glBindBuffer( GL_PIXEL_UNPACK_BUFFER, pbo); glBindTexture(GL_TEXTURE_2D, displayImage); - glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, width, height, + glBindTexture(GL_TEXTURE_2D, displayImage); + glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, width, height, GL_RGBA, GL_UNSIGNED_BYTE, NULL); glClear(GL_COLOR_BUFFER_BIT); @@ -352,3 +386,51 @@ void shut_down(int return_code){ #endif exit(return_code); } + +void onMouseCb(int button, int state, int x, int y) +{ + theButtonState = button; + theModifierState = glutGetModifiers(); + lastX = x; + lastY = y; +} + +void onMouseMotionCb(int x, int y) +{ + int deltaX = lastX - x; + int deltaY = lastY - y; + bool moveLeftRight = abs(deltaX) > abs(deltaY); + bool moveUpDown = !moveLeftRight; + + switch(theButtonState) + { + case GLUT_LEFT_BUTTON: + // Move Camera + if (theModifierState & GLUT_ACTIVE_ALT) + { + + if (deltaY > 0) + { + if (eye.pos.z < 100) + eye.pos.z += 5; + + } + else if (deltaY < 0) + { + if (eye.pos.z > 10) + eye.pos.z -= 5; + } + } + else + { + if (deltaX > 0) eye.rot.y -= 5; + else if (deltaX < 0) eye.rot.y += 5; + else if (deltaY > 0) eye.rot.x += 5; + else if (deltaY < 0) eye.rot.x -= 5; + } + break; + } + lastX = x; + lastY = y; + glutPostRedisplay(); +} diff --git a/src/main.h b/src/main.h index 63bf0fa..b5fb223 100755 --- a/src/main.h +++ b/src/main.h @@ -26,6 +26,7 @@ #include "rasterizeKernels.h" #include "utilities.h" #include "ObjCore/objloader.h" +#include "Eye.h" using namespace std; @@ -43,6 +44,7 @@ GLuint pbo = (GLuint)NULL; GLuint displayImage; uchar4 *dptr; +Eye eye; obj* mesh; float* vbo; @@ -51,6 +53,8 @@ float* cbo; int cbosize; int* ibo; int ibosize; +float* nbo; +int nbosize; //------------------------------- //----------CUDA STUFF----------- @@ -101,5 +105,7 @@ void cleanupCuda(); void deletePBO(GLuint* pbo); void deleteTexture(GLuint* tex); void shut_down(int return_code); +void onMouseCb(int button, int state, int x, int y); +void onMouseMotionCb(int x, int y); #endif \ No newline at end of file diff --git a/src/rasterizeKernels.cu b/src/rasterizeKernels.cu index 826ec80..bb6bf2d 100755 --- a/src/rasterizeKernels.cu +++ b/src/rasterizeKernels.cu @@ -1,6 +1,9 @@ // CIS565 CUDA Rasterizer: A simple rasterization pipeline for Patrick Cozzi's CIS565: GPU Computing at the University of Pennsylvania // Written by Yining Karl Li, Copyright (c) 2012 University of Pennsylvania +#include +#include +#include #include #include #include @@ -12,10 +15,20 @@ glm::vec3* framebuffer; fragment* depthbuffer; float* device_vbo; +float* device_nbo; float* device_cbo; int* device_ibo; triangle* primitives; +struct IsInvalidPixel +{ + __host__ __device__ + bool operator()(const fragment& fragmentVal) + { + return !fragmentVal.valid; + } +}; + void checkCUDAError(const char *msg) { cudaError_t err = cudaGetLastError(); if( cudaSuccess != err) { @@ -72,6 +85,29 @@ __host__ __device__ glm::vec3 getFromFramebuffer(int x, int y, glm::vec3* frameb } } +__host__ __device__ glm::vec3 ScreenPointInWorldCoordinates(glm::vec2 resolution, float x, float y, + glm::vec3 M, glm::vec3 A, glm::vec3 B) +{ + float sx = (x/(float)(resolution.x-1)); + float sy = (y/(float)(resolution.y-1)); + return (M - A*(resolution.x/2.0f)*(2.0f*sx - 1) - B*(resolution.y/2.0f)*(2.0f*sy - 1)); +} + +__host__ __device__ glm::vec2 WorldToScreen(glm::vec2 resolution, glm::vec3 pt, glm::vec2 eyeLeftRight, glm::vec2 eyeTopBottom) +{ + /*float sx = (((M.x - pt.x)/(A.x * resolution.x/2)) + 1)/2; + float sy = (((M.y - pt.y)/(B.y * resolution.y/2)) + 1)/2; + glm::vec2 screenPoint; + screenPoint.x = sx*(resolution.x - 1); + screenPoint.y = sy*(resolution.y - 1); + return screenPoint;*/ + + glm::vec2 screenPoint; + screenPoint.x = (pt.x - (eyeLeftRight.x))*(resolution.x/2); + screenPoint.y = (pt.y - (eyeTopBottom.x))*(resolution.y/2); + return screenPoint; +} + //Kernel that clears a given pixel buffer with a given color __global__ void clearImage(glm::vec2 resolution, glm::vec3* image, glm::vec3 color){ int x = (blockIdx.x * blockDim.x) + threadIdx.x; @@ -121,42 +157,451 @@ __global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* color.z = 255; } + int PBOIndex = resolution.x - 1 - x + ((resolution.y - 1 - y) * resolution.x); + // Each thread writes one pixel location in the texture (textel) - PBOpos[index].w = 0; - PBOpos[index].x = color.x; - PBOpos[index].y = color.y; - PBOpos[index].z = color.z; + PBOpos[PBOIndex].w = 0.5; + PBOpos[PBOIndex].x = color.x; + PBOpos[PBOIndex].y = color.y; + PBOpos[PBOIndex].z = color.z; } } + //TODO: Implement a vertex shader -__global__ void vertexShadeKernel(float* vbo, int vbosize){ +__global__ void vertexShadeKernel(float* vbo, float* vboOriginal, float* normals, int vbosize, cudaMat4 transform, cudaMat4 modelView){ int index = (blockIdx.x * blockDim.x) + threadIdx.x; if(index=yEnd; --y) + { + float xMin = 100000; + float xMax = -xMin; + if ((float)y<=y0 && (float)y>=y1) + { + float x; + if (y1 != y0) + { + x = x0 + ((float)y-y0)*(x1-x0)/(y1-y0); + if (xxMax) + xMax = x; + } + else + { + // Line parallel to x + if (x0xMax) + xMax = x0; + if (x1xMax) + xMax = x1; + } + + } + + if ((float)y<=y0 && (float)y>=y2) + { + if (y2 != y0) + { + float x = x0 + ((float)y-y0)*(x2-x0)/(y2-y0); + if (xxMax) + xMax = x; + } + else + { + // Line parallel to x + if (x0xMax) + xMax = x0; + if (x2xMax) + xMax = x2; + } + } + + if ((float)y<=y1 && (float)y>=y2) + { + if (y1 != y2) + { + float x = x1 + ((float)y-y1)*(x2-x1)/(y2-y1); + if (xxMax) + xMax = x; + } + else + { + // Line parallel to x + if (x2xMax) + xMax = x2; + if (x1xMax) + xMax = x1; + } + } + + // Check if the depth is smallest + int xStart = ceil(xMin); + int xEnd = floor(xMax); + + xStart = max(xStart, 0); + xEnd = min(xEnd, (int)resolution.x); + + //if(y==yStart-1 && (index == 1 || index ==2)) + /*{ + printf("y: %i, xStart: %i, xEnd: %i\n", y, xStart, xEnd); + }*/ + + for (int x=xStart; x<=xEnd; ++x) + { + if (x=0 && y depthbuffer[i].depthVal) + depthValue[0] = depthbuffer[i].depthVal; + + if (depthbuffer[i].valid && depthValue[1] < depthbuffer[i].depthVal) + depthValue[1] = depthbuffer[i].depthVal; + } + + depthValue[1] -= depthValue[0]; +} + //TODO: Implement a fragment shader -__global__ void fragmentShadeKernel(fragment* depthbuffer, glm::vec2 resolution){ +__global__ void fragmentShadeKernel(fragment* depthbuffer, glm::vec2 resolution, float3* lightDir, float3* lightColor, int lightCount, + bool* stencil, bool stencilPresent, float* boundZ, bool depthShade, bool toShade){ int x = (blockIdx.x * blockDim.x) + threadIdx.x; int y = (blockIdx.y * blockDim.y) + threadIdx.y; int index = x + (y * resolution.x); if(x<=resolution.x && y<=resolution.y){ + if (depthbuffer[index].valid) + { + if (!stencilPresent || (stencilPresent && stencil[depthbuffer[index].pixelIndex])) + { + if (depthShade) + { + float color = 1.f - (0.15+(0.7*(depthbuffer[index].depthVal - boundZ[0])/boundZ[1])); + depthbuffer[index].color = glm::vec3(color,color,color); + } + else if(toShade) + { + // View Direction is camera - point. Camera is at 0,0,0 + float3 viewDir = make_float3(-depthbuffer[index].position.x, -depthbuffer[index].position.y, -depthbuffer[index].position.z); + if (viewDir.x != 0 && viewDir.y != 0 && viewDir.z != 0) + viewDir = normalize(viewDir); + + float3 normal = make_float3(depthbuffer[index].normal.x, depthbuffer[index].normal.y, depthbuffer[index].normal.z); + if (normal.x != 0 && normal.y != 0 && normal.z != 0) + normal = normalize(normal); + + float3 finalColor = make_float3(0, 0, 0); + float specularity = 20; + float3 color = make_float3(depthbuffer[index].color.x, depthbuffer[index].color.y, depthbuffer[index].color.z); + float kd = 0.6f; + float ks = 0.1f; + for (int k = 0; k < lightCount; ++k) + { + float3 halfVec = normalize(viewDir + lightDir[k]); + float cosTh = dot(normal, halfVec); + float cosTi = dot(normal, lightDir[k]); + if (cosTi < 0) cosTi = 0; + if (cosTi > 1) cosTi = 1; + if (cosTh < 0) cosTh = 0; + if (cosTh > 1) cosTh = 1; + finalColor += (kd*color/*/3.14*/ + ks/*(specularity + 8)/(8*3.14)*/ * pow(cosTh, specularity)) * lightColor[k] * cosTi; + } + + depthbuffer[index].color.x = finalColor.x; + depthbuffer[index].color.y = finalColor.y; + depthbuffer[index].color.z = finalColor.z; + } + } + else + { + depthbuffer[index].color = glm::vec3(0,0,0); + } + } + else + { + depthbuffer[index].color = glm::vec3(0,0,0); + } } } @@ -168,17 +613,143 @@ __global__ void render(glm::vec2 resolution, fragment* depthbuffer, glm::vec3* f int index = x + (y * resolution.x); if(x<=resolution.x && y<=resolution.y){ - framebuffer[index] = depthbuffer[index].color; + if (depthbuffer[index].valid) + framebuffer[depthbuffer[index].pixelIndex] = depthbuffer[index].color; } } +//Writes fragment colors to the framebuffer +__global__ void renderAntialias(glm::vec2 resolution, glm::vec2 newResolution, fragment* depthbuffer, glm::vec3* framebuffer){ + + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * resolution.x); + + if(x<=resolution.x && y<=resolution.y){ + int depthbufferOffsetX = 2*x + 1; + int depthbufferOffsetY = 2*y + 1; + int depthbufferIndex = (depthbufferOffsetX) + (depthbufferOffsetY)*resolution.y*2.f; + + // X + if (depthbuffer[depthbufferIndex].valid) + { + framebuffer[index] += depthbuffer[depthbufferIndex].color; + } + if (depthbufferOffsetY-1 >= 0) + { + depthbufferIndex = (depthbufferOffsetX) + (depthbufferOffsetY-1)*resolution.y*2.f; + if (depthbuffer[depthbufferIndex].valid) + { + framebuffer[index] += depthbuffer[depthbufferIndex].color; + } + } + if (depthbufferOffsetY+1 < newResolution.y) + { + depthbufferIndex = (depthbufferOffsetX-1) + (depthbufferOffsetY+1)*resolution.y*2.f; + if (depthbuffer[depthbufferIndex].valid) + { + framebuffer[index] += depthbuffer[depthbufferIndex].color; + } + } + + // X-1 + if (depthbufferOffsetX - 1 >= 0) + { + if (depthbufferOffsetY-1 >= 0) + { + depthbufferIndex = (depthbufferOffsetX-1) + (depthbufferOffsetY-1)*resolution.y*2.f; + if (depthbuffer[depthbufferIndex].valid) + { + framebuffer[index] += depthbuffer[depthbufferIndex].color; + } + } + + depthbufferIndex = (depthbufferOffsetX-1) + (depthbufferOffsetY)*resolution.y*2.f; + if (depthbuffer[depthbufferIndex].valid) + { + framebuffer[index] += depthbuffer[depthbufferIndex].color; + } + + if (depthbufferOffsetY+1 < newResolution.y) + { + depthbufferIndex = (depthbufferOffsetX-1) + (depthbufferOffsetY+1)*resolution.y*2.f; + if (depthbuffer[depthbufferIndex].valid) + { + framebuffer[index] += depthbuffer[depthbufferIndex].color; + } + } + } + + // X+1 + if (depthbufferOffsetX + 1 < newResolution.x) + { + if (depthbufferOffsetY-1 >= 0) + { + depthbufferIndex = (depthbufferOffsetX+1) + (depthbufferOffsetY-1)*resolution.y*2.f; + if (depthbuffer[depthbufferIndex].valid) + { + framebuffer[index] += depthbuffer[depthbufferIndex].color; + } + } + + depthbufferIndex = (depthbufferOffsetX+1) + (depthbufferOffsetY)*resolution.y*2.f; + if (depthbuffer[depthbufferIndex].valid) + { + framebuffer[index] += depthbuffer[depthbufferIndex].color; + } + + if (depthbufferOffsetY+1 < newResolution.y) + { + depthbufferIndex = (depthbufferOffsetX+1) + (depthbufferOffsetY+1)*resolution.y*2.f; + if (depthbuffer[depthbufferIndex].valid) + { + framebuffer[index] += depthbuffer[depthbufferIndex].color; + } + } + } + + framebuffer[index] /= 9.f; + } +} + +__global__ void printCulled(int* culled) +{ + printf("Culled: %i\n", culled[0]); +} + // Wrapper for the __global__ call that sets up the kernel calls and does a ton of memory management -void cudaRasterizeCore(uchar4* PBOpos, glm::vec2 resolution, float frame, float* vbo, int vbosize, float* cbo, int cbosize, int* ibo, int ibosize){ +void cudaRasterizeCore(uchar4* PBOpos, glm::vec2 resolution, float frame, float* vbo, int vbosize, float* cbo, int cbosize, + int* ibo, int ibosize, float* nbo, int nbosize, Eye eye, cudaMat4 transform, cudaMat4 modelView, cudaMat4 perspective) +{ + glm::vec3 M, A, B; + float distImagePlaneFromCamera; + eye.GetParametersForScreenTransform(resolution, M, A, B, distImagePlaneFromCamera); + + //printf("M: (%f, %f, %f) A: (%f, %f, %f) B: (%f, %f, %f)", M.x, M.y, M.z, A.x, A.y, A.z, B.x, B.y, B.z); + + bool objAntialias = false; + bool colorInterp = true; + bool antialias = false; + bool depthShade = false; + bool toBeShaded = true; + bool stencilPresent = false; + + int totalPixels = (int)resolution.x*(int)resolution.y; + int newTotalPixels = totalPixels; + glm::vec2 newResolution = resolution; + if (objAntialias) + { + newResolution *= 2.f; + newTotalPixels *= 4; + } // set up crucial magic int tileSize = 8; + int tileSizeInitial = tileSize; dim3 threadsPerBlock(tileSize, tileSize); dim3 fullBlocksPerGrid((int)ceil(float(resolution.x)/float(tileSize)), (int)ceil(float(resolution.y)/float(tileSize))); + dim3 fullBlocksPerGridNew((int)ceil(float(newResolution.x)/float(tileSize)), (int)ceil(float(newResolution.y)/float(tileSize))); + dim3 fullBlocksPerGridAntialias((int)ceil(float(newResolution.x)/float(tileSize)), (int)ceil(float(newResolution.y)/float(tileSize))); //set up framebuffer framebuffer = NULL; @@ -186,17 +757,45 @@ void cudaRasterizeCore(uchar4* PBOpos, glm::vec2 resolution, float frame, float* //set up depthbuffer depthbuffer = NULL; - cudaMalloc((void**)&depthbuffer, (int)resolution.x*(int)resolution.y*sizeof(fragment)); + cudaMalloc((void**)&depthbuffer, newTotalPixels*sizeof(fragment)); //kernel launches to black out accumulated/unaccumlated pixel buffers and clear our scattering states clearImage<<>>(resolution, framebuffer, glm::vec3(0,0,0)); fragment frag; + frag.valid = false; frag.color = glm::vec3(0,0,0); frag.normal = glm::vec3(0,0,0); - frag.position = glm::vec3(0,0,-10000); - clearDepthBuffer<<>>(resolution, depthbuffer,frag); - + frag.position = glm::vec3(0,0,0); + frag.depthVal = 10000.f; + frag.pixelIndex = 0; + clearDepthBuffer<<>>(newResolution, depthbuffer, frag); + + bool* stencil = new bool[newTotalPixels]; + bool* device_stencil = NULL; + + if (stencilPresent) + { + // Create Stencil + for (int i = 0; i 150 && i < 300 && j > 300 && j < 525) + stencil[indexVal] = true; + if (i > 350 && i < 550 && j > 550 && j < 700) + stencil[indexVal] = true; + } + } + + cudaMalloc((void**)&device_stencil, newTotalPixels*sizeof(bool)); + cudaMemcpy( device_stencil, stencil, newTotalPixels*sizeof(bool), cudaMemcpyHostToDevice); + } + int* depthBufferLock = NULL; + cudaMalloc((void**)&depthBufferLock, newTotalPixels*sizeof(int)); + cudaMemset(depthBufferLock, 0, newTotalPixels*sizeof(int)); //------------------------------ //memory stuff //------------------------------ @@ -211,6 +810,14 @@ void cudaRasterizeCore(uchar4* PBOpos, glm::vec2 resolution, float frame, float* cudaMalloc((void**)&device_vbo, vbosize*sizeof(float)); cudaMemcpy( device_vbo, vbo, vbosize*sizeof(float), cudaMemcpyHostToDevice); + float* device_vboWorld = NULL; + cudaMalloc((void**)&device_vboWorld, vbosize*sizeof(float)); + cudaMemcpy( device_vboWorld, vbo, vbosize*sizeof(float), cudaMemcpyHostToDevice); + + device_nbo = NULL; + cudaMalloc((void**)&device_nbo, nbosize*sizeof(float)); + cudaMemcpy( device_nbo, nbo, nbosize*sizeof(float), cudaMemcpyHostToDevice); + device_cbo = NULL; cudaMalloc((void**)&device_cbo, cbosize*sizeof(float)); cudaMemcpy( device_cbo, cbo, cbosize*sizeof(float), cudaMemcpyHostToDevice); @@ -221,36 +828,98 @@ void cudaRasterizeCore(uchar4* PBOpos, glm::vec2 resolution, float frame, float* //------------------------------ //vertex shader //------------------------------ - vertexShadeKernel<<>>(device_vbo, vbosize); + vertexShadeKernel<<>>(device_vbo, device_vboWorld, device_nbo, vbosize, transform, modelView); cudaDeviceSynchronize(); //------------------------------ //primitive assembly //------------------------------ primitiveBlocks = ceil(((float)ibosize/3)/((float)tileSize)); - primitiveAssemblyKernel<<>>(device_vbo, vbosize, device_cbo, cbosize, device_ibo, ibosize, primitives); + primitiveAssemblyKernel<<>>(device_vbo, vbosize, device_cbo, cbosize, device_ibo, ibosize, device_nbo, nbosize, device_vboWorld, primitives); cudaDeviceSynchronize(); + //------------------------------ //rasterization //------------------------------ - rasterizationKernel<<>>(primitives, ibosize/3, depthbuffer, resolution); - + glm::vec2 eyeLeftRight = glm::vec2(eye.l, eye.r); + glm::vec2 eyeTopBottom = glm::vec2(eye.b, eye.t); + rasterizationKernel<<>>(primitives, ibosize/3, depthbuffer, newResolution, eyeLeftRight, eyeTopBottom, + depthBufferLock, antialias, colorInterp); cudaDeviceSynchronize(); + + float totalValidPixels = newTotalPixels; + if (!objAntialias) + { + // wrap raw pointer with a device_ptr for Stream compaction + thrust::device_ptr devDepthBufferPtr(depthbuffer); + thrust::device_ptr devDepthBufferEndPtr = thrust::remove_if(devDepthBufferPtr, devDepthBufferPtr+totalPixels, IsInvalidPixel()); + totalValidPixels = devDepthBufferEndPtr.get() - devDepthBufferPtr.get(); + + // Create blocks for lesser number of pixels found by stream compaction + fullBlocksPerGridNew = dim3((int)ceil(float(newResolution.x)/float(tileSizeInitial)), (int)ceil((float(totalValidPixels)/float(newResolution.y))/float(tileSizeInitial))); + } + + float* device_depthValue = NULL; + cudaMalloc((void**)&device_depthValue, 2*sizeof(float)); + FindMinMaxDepth<<<1,1>>>(depthbuffer, totalValidPixels, device_depthValue); + //------------------------------ //fragment shader //------------------------------ - fragmentShadeKernel<<>>(depthbuffer, resolution); + int numLights = 4; + float3* lightDir = new float3[numLights]; + lightDir[0] = make_float3(0, 0, 1); + lightDir[1] = make_float3(0, 0, -1); + lightDir[2] = make_float3(1, 1, 0); + lightDir[3] = make_float3(-1, -1, 0); + for (int i = 0; i < numLights; ++i) + { + lightDir[i] = normalize(lightDir[i]); + } + + float3* lightColor = new float3[numLights]; + lightColor[0] = make_float3(1, 1, 1); + lightColor[1] = make_float3(1,1,1); + lightColor[2] = make_float3(1, 1, 1); + lightColor[3] = make_float3(1, 1, 1); + + float3* device_lightDir = NULL; + cudaMalloc((void**)&device_lightDir, numLights*sizeof(float3)); + cudaMemcpy( device_lightDir, lightDir, numLights*sizeof(float3), cudaMemcpyHostToDevice); + + float3* device_lightColor = NULL; + cudaMalloc((void**)&device_lightColor, numLights*sizeof(float3)); + cudaMemcpy( device_lightColor, lightColor, numLights*sizeof(float3), cudaMemcpyHostToDevice); + + if (toBeShaded || depthShade || stencilPresent) + { + fragmentShadeKernel<<>>(depthbuffer, newResolution, device_lightDir, device_lightColor, numLights, + device_stencil, stencilPresent, device_depthValue, depthShade, toBeShaded); + } cudaDeviceSynchronize(); + delete[] lightDir; + delete[] lightColor; + cudaFree(device_lightDir); + cudaFree(device_lightColor); //------------------------------ //write fragments to framebuffer //------------------------------ - render<<>>(resolution, depthbuffer, framebuffer); + if (!objAntialias) + render<<>>(resolution, depthbuffer, framebuffer); + else + renderAntialias<<>>(resolution, newResolution, depthbuffer, framebuffer); + sendImageToPBO<<>>(PBOpos, resolution, framebuffer); cudaDeviceSynchronize(); + // delete[] boundZ; + cudaFree(device_depthValue); + cudaFree(device_stencil); + cudaFree(depthBufferLock); + cudaFree(device_vboWorld); kernelCleanup(); checkCUDAError("Kernel failed!"); @@ -261,6 +930,7 @@ void kernelCleanup(){ cudaFree( device_vbo ); cudaFree( device_cbo ); cudaFree( device_ibo ); + cudaFree( device_nbo ); cudaFree( framebuffer ); cudaFree( depthbuffer ); } diff --git a/src/rasterizeKernels.h b/src/rasterizeKernels.h index bef3653..f458fe7 100755 --- a/src/rasterizeKernels.h +++ b/src/rasterizeKernels.h @@ -10,8 +10,11 @@ #include #include #include "glm/glm.hpp" +#include "cudaMat4.h" +#include "Eye.h" void kernelCleanup(); -void cudaRasterizeCore(uchar4* pos, glm::vec2 resolution, float frame, float* vbo, int vbosize, float* cbo, int cbosize, int* ibo, int ibosize); +void cudaRasterizeCore(uchar4* PBOpos, glm::vec2 resolution, float frame, float* vbo, int vbosize, float* cbo, int cbosize, + int* ibo, int ibosize, float* nbo, int nbosize, Eye eye, cudaMat4 transform, cudaMat4 modelView, cudaMat4 perspective); #endif //RASTERIZEKERNEL_H diff --git a/src/rasterizeTools.h b/src/rasterizeTools.h index e9b5dcc..b201467 100755 --- a/src/rasterizeTools.h +++ b/src/rasterizeTools.h @@ -16,12 +16,21 @@ struct triangle { glm::vec3 c0; glm::vec3 c1; glm::vec3 c2; + glm::vec3 n0; + glm::vec3 n1; + glm::vec3 n2; + glm::vec3 p0Original; + glm::vec3 p1Original; + glm::vec3 p2Original; }; struct fragment{ glm::vec3 color; glm::vec3 normal; glm::vec3 position; + float depthVal; + bool valid; + int pixelIndex; }; //Multiplies a cudaMat4 matrix and a vec4 @@ -75,4 +84,13 @@ __host__ __device__ float getZAtCoordinate(glm::vec3 barycentricCoord, triangle return -(barycentricCoord.x*tri.p0.z + barycentricCoord.y*tri.p1.z + barycentricCoord.z*tri.p2.z); } +__host__ __device__ void printCudaMat4(cudaMat4 matrix) +{ + printf("Matrix:\n(%f, %f, %f, %f)\n(%f, %f, %f, %f)\n(%f, %f, %f, %f)\n(%f, %f, %f, %f)\n", + matrix.x.x, matrix.x.y, matrix.x.z, matrix.x.w, + matrix.y.x, matrix.y.y, matrix.y.z, matrix.y.w, + matrix.z.x, matrix.z.y, matrix.z.z, matrix.z.w, + matrix.w.x, matrix.w.y, matrix.w.z, matrix.w.w); +} + #endif \ No newline at end of file diff --git a/src/utilities.cpp b/src/utilities.cpp index 4f8f4d3..aa7ded7 100755 --- a/src/utilities.cpp +++ b/src/utilities.cpp @@ -70,7 +70,8 @@ glm::mat4 utilityCore::buildTransformationMatrix(glm::vec3 translation, glm::vec rotationMat = rotationMat*glm::rotate(glm::mat4(), rotation.y, glm::vec3(0,1,0)); rotationMat = rotationMat*glm::rotate(glm::mat4(), rotation.z, glm::vec3(0,0,1)); glm::mat4 scaleMat = glm::scale(glm::mat4(), scale); - return translationMat*rotationMat*scaleMat; + //return translationMat*rotationMat*scaleMat; + return scaleMat*rotationMat*translationMat; } cudaMat4 utilityCore::glmMat4ToCudaMat4(glm::mat4 a){ @@ -82,6 +83,15 @@ cudaMat4 utilityCore::glmMat4ToCudaMat4(glm::mat4 a){ return m; } +cudaMat4 utilityCore::glmMat4ToCudaMat4NoTranspose(glm::mat4 a){ + cudaMat4 m; + m.x = a[0]; + m.y = a[1]; + m.z = a[2]; + m.w = a[3]; + return m; +} + glm::mat4 utilityCore::cudaMat4ToGlmMat4(cudaMat4 a){ glm::mat4 m; m[0] = a.x; diff --git a/src/utilities.h b/src/utilities.h index 3e6ef6e..f421171 100755 --- a/src/utilities.h +++ b/src/utilities.h @@ -29,6 +29,7 @@ namespace utilityCore { extern bool epsilonCheck(float a, float b); extern std::vector tokenizeString(std::string str); extern cudaMat4 glmMat4ToCudaMat4(glm::mat4 a); + extern cudaMat4 glmMat4ToCudaMat4NoTranspose(glm::mat4 a); extern glm::mat4 cudaMat4ToGlmMat4(cudaMat4 a); extern glm::mat4 buildTransformationMatrix(glm::vec3 translation, glm::vec3 rotation, glm::vec3 scale); extern void printCudaMat4(cudaMat4 m);