diff --git a/README.md b/README.md index ae0896a..e6e2cac 100644 --- a/README.md +++ b/README.md @@ -3,50 +3,11 @@ CIS565: Project 4: CUDA Rasterizer ------------------------------------------------------------------------------- Fall 2014 ------------------------------------------------------------------------------- -Due Monday 10/27/2014 @ 12 PM -------------------------------------------------------------------------------- - -------------------------------------------------------------------------------- -NOTE: -------------------------------------------------------------------------------- -This project requires an NVIDIA graphics card with CUDA capability! Any card with CUDA compute capability 1.1 or higher will work fine for this project. For a full list of CUDA capable cards and their compute capability, please consult: http://developer.nvidia.com/cuda/cuda-gpus. If you do not have an NVIDIA graphics card in the machine you are working on, feel free to use any machine in the SIG Lab or in Moore100 labs. All machines in the SIG Lab and Moore100 are equipped with CUDA capable NVIDIA graphics cards. If this too proves to be a problem, please contact Patrick or Karl as soon as possible. - -------------------------------------------------------------------------------- -INTRODUCTION: -------------------------------------------------------------------------------- -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 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. - -You MAY NOT use ANY raycasting/raytracing AT ALL in this project, EXCEPT in the fragment shader step. One of the purposes of this project is to see how a rasterization pipeline can generate graphics WITHOUT the need for raycasting! Raycasting may only be used in the fragment shader effect for interesting shading results, but is absolutely not allowed in any other stages of the pipeline. - -Also, you MAY NOT use OpenGL ANYWHERE in this project, aside from the given OpenGL code for drawing Pixel Buffer Objects to the screen. Use of OpenGL for any pipeline stage instead of your own custom implementation will result in an incomplete project. - -Finally, note that while this basecode is meant to serve as a strong starting point for a CUDA rasterizer, you are not required to use this basecode if you wish, and you may also change any part of the basecode specification as you please, so long as the final rendered result is correct. +Completed Features: ------------------------------------------------------------------------------- -CONTENTS: -------------------------------------------------------------------------------- -The Project4 root directory contains the following subdirectories: - -* src/ contains the source code for the project. Both the Windows Visual Studio solution and the OSX makefile reference this folder for all source; the base source code compiles on OSX and Windows without modification. -* objs/ contains example obj test files: cow.obj, cube.obj, tri.obj. -* renders/ contains an example render of the given example cow.obj file with a z-depth fragment shader. -* windows/ contains a Windows Visual Studio 2010 project and all dependencies needed for building and running on Windows 7. - -The Windows and OSX versions of the project build and run exactly the same way as in Project0, Project1, and Project2. - -------------------------------------------------------------------------------- -REQUIREMENTS: -------------------------------------------------------------------------------- -In this project, you are given code for: - -* A library for loading/reading standard Alias/Wavefront .obj format mesh files and converting them to OpenGL style VBOs/IBOs -* A suggested order of kernels with which to implement the graphics pipeline -* Working code for CUDA-GL interop - You will need to implement the following stages of the graphics pipeline and features: - +Required: * Vertex Shading * Primitive Assembly with support for triangle VBOs/IBOs * Perspective Transformation @@ -56,129 +17,40 @@ You will need to implement the following stages of the graphics pipeline and fea * Fragment to framebuffer writing * A simple lighting/shading scheme, such as Lambert or Blinn-Phong, implemented in the fragment shader -You are also required to implement at least 3 of the following features: - -* Additional pipeline stages. Each one of these stages can count as 1 feature: - * Geometry shader - * Transformation feedback +Additional: +* Additional pipeline stages. * Back-face culling - * Scissor test - * Stencil test - * Blending - -IMPORTANT: For each of these stages implemented, you must also add a section to your README stating what the expected performance impact of that pipeline stage is, and real performance comparisons between your rasterizer with that stage and without. - * Correct color interpolation 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 primitices. * Lines - * Line strips - * Triangle fans - * Triangle strips * Points -* Anti-aliasing -* Order-independent translucency using a k-buffer * MOUSE BASED interactive camera support. Interactive camera support based only on the keyboard is not acceptable for this feature. -------------------------------------------------------------------------------- -BASE CODE TOUR: -------------------------------------------------------------------------------- -You will be working primarily in two files: rasterizeKernel.cu, and rasterizerTools.h. Within these files, areas that you need to complete are marked with a TODO comment. Areas that are useful to and serve as hints for optional features are marked with TODO (Optional). Functions that are useful for reference are marked with the comment LOOK. - -* rasterizeKernels.cu contains the core rasterization pipeline. - * A suggested sequence of kernels exists in this file, but you may choose to alter the order of this sequence or merge entire kernels if you see fit. For example, if you decide that doing has benefits, you can choose to merge the vertex shader and primitive assembly kernels, or merge the perspective transform into another kernel. There is not necessarily a right sequence of kernels (although there are wrong sequences, such as placing fragment shading before vertex shading), and you may choose any sequence you want. Please document in your README what sequence you choose and why. - * The provided kernels have had their input parameters removed beyond basic inputs such as the framebuffer. You will have to decide what inputs should go into each stage of the pipeline, and what outputs there should be. - -* rasterizeTools.h contains various useful tools, including a number of barycentric coordinate related functions that you may find useful in implementing scanline based rasterization... - * A few pre-made structs are included for you to use, such as fragment and triangle. A simple rasterizer can be implemented with these structs as is. However, as with any part of the basecode, you may choose to modify, add to, use as-is, or outright ignore them as you see fit. - * If you do choose to add to the fragment struct, be sure to include in your README a rationale for why. - -You will also want to familiarize yourself with: - -* main.cpp, which contains code that transfers VBOs/CBOs/IBOs to the rasterization pipeline. Interactive camera work will also have to be implemented in this file if you choose that feature. -* utilities.h, which serves as a kitchen-sink of useful functions - -------------------------------------------------------------------------------- -SOME RESOURCES: -------------------------------------------------------------------------------- -The following resources may be useful for this project: - -* High-Performance Software Rasterization on GPUs - * Paper (HPG 2011): http://www.tml.tkk.fi/~samuli/publications/laine2011hpg_paper.pdf - * Code: http://code.google.com/p/cudaraster/ Note that looking over this code for reference with regard to the paper is fine, but we most likely will not grant any requests to actually incorporate any of this code into your project. - * Slides: http://bps11.idav.ucdavis.edu/talks/08-gpuSoftwareRasterLaineAndPantaleoni-BPS2011.pdf -* The Direct3D 10 System (SIGGRAPH 2006) - for those interested in doing geometry shaders and transform feedback. - * http://133.11.9.3/~takeo/course/2006/media/papers/Direct3D10_siggraph2006.pdf -* Multi-Fragment Effects on the GPU using the k-Buffer - for those who want to do a k-buffer - * http://www.inf.ufrgs.br/~comba/papers/2007/kbuffer_preprint.pdf -* FreePipe: A Programmable, Parallel Rendering Architecture for Efficient Multi-Fragment Effects (I3D 2010) - * https://sites.google.com/site/hmcen0921/cudarasterizer -* Writing A Software Rasterizer In Javascript: - * Part 1: http://simonstechblog.blogspot.com/2012/04/software-rasterizer-part-1.html - * Part 2: http://simonstechblog.blogspot.com/2012/04/software-rasterizer-part-2.html - -------------------------------------------------------------------------------- -NOTES ON GLM: -------------------------------------------------------------------------------- -This project uses GLM, the GL Math library, for linear algebra. You need to know two important points on how GLM is used in this project: - -* In this project, indices in GLM vectors (such as vec3, vec4), are accessed via swizzling. So, instead of v[0], v.x is used, and instead of v[1], v.y is used, and so on and so forth. -* GLM Matrix operations work fine on NVIDIA Fermi cards and later, but pre-Fermi cards do not play nice with GLM matrices. As such, in this project, GLM matrices are replaced with a custom matrix struct, called a cudaMat4, found in cudaMat4.h. A custom function for multiplying glm::vec4s and cudaMat4s is provided as multiplyMV() in intersections.h. - -------------------------------------------------------------------------------- -README -------------------------------------------------------------------------------- -All students must replace or augment the contents of this Readme.md in a clear -manner with the following: - -* A brief description of the project and the specific features you implemented. -* At least one screenshot of your project running. -* A 30 second or longer video of your project running. To create the video you - can use http://www.microsoft.com/expression/products/Encoder4_Overview.aspx -* A performance evaluation (described in detail below). - -------------------------------------------------------------------------------- -PERFORMANCE EVALUATION -------------------------------------------------------------------------------- -The performance evaluation is where you will investigate how to make your CUDA -programs more efficient using the skills you've learned in class. You must have -performed at least one experiment on your code to investigate the positive or -negative effects on performance. - -We encourage you to get creative with your tweaks. Consider places in your code -that could be considered bottlenecks and try to improve them. - -Each student should provide no more than a one page summary of their -optimizations along with tables and or graphs to visually explain any -performance differences. - -------------------------------------------------------------------------------- -THIRD PARTY CODE POLICY -------------------------------------------------------------------------------- -* Use of any third-party code must be approved by asking on Piazza. If it is approved, all students are welcome to use it. Generally, we approve use of third-party code that is not a core part of the project. For example, for the ray tracer, we would approve using a third-party library for loading models, but would not approve copying and pasting a CUDA function for doing refraction. -* Third-party code must be credited in README.md. -* Using third-party code without its approval, including using another student's code, is an academic integrity violation, and will result in you receiving an F for the semester. - -------------------------------------------------------------------------------- -SELF-GRADING -------------------------------------------------------------------------------- -* On the submission date, email your grade, on a scale of 0 to 100, to Liam, harmoli+cis565@seas.upenn.edu, with a one paragraph explanation. Be concise and realistic. Recall that we reserve 30 points as a sanity check to adjust your grade. Your actual grade will be (0.7 * your grade) + (0.3 * our grade). We hope to only use this in extreme cases when your grade does not realistically reflect your work - it is either too high or too low. In most cases, we plan to give you the exact grade you suggest. -* Projects are not weighted evenly, e.g., Project 0 doesn't count as much as the path tracer. We will determine the weighting at the end of the semester based on the size of each project. - ---- -SUBMISSION ---- -As with the previous project, you should fork this project and work inside of -your fork. Upon completion, commit your finished project back to your fork, and -make a pull request to the master repository. You should include a README.md -file in the root directory detailing the following - -* A brief description of the project and specific features you implemented -* At least one screenshot of your project running. -* A link to a video of your raytracer running. -* Instructions for building and running your project if they differ from the - base code. -* A performance writeup as detailed above. -* A list of all third-party code used. -* This Readme file edited as described above in the README section. +#How to use: + +Use #define RENDER_MODE_0 in rasterizeKernels.cu to use different shading mode +0 for blinn-phong, 1 for wire-frame, 2 for interpolated color, 3 for vertices +Hold and drag mouse to rotate the camera. Use mouse scroll to zoom in and out. + +#Result: +![Alt text](/img/phong.PNG "Blinn-Phong shading") +Blinn-Phong shading +![Alt text](/img/interp.PNG) +Color interpolation +![Alt text](/img/wire.PNG) +Wire- Frame +![Alt text](/img/vertices.PNG) +Vertices +#Performance: +Graphic card: GTX780 + +Firstly I tried to ananysize the relationship between tile size and one frame runtime. However the result shows they are almost constant against tile size. (See the following chart) +![Alt text](/img/tilesize.PNG) + +Then I used the mouse to zoom in and out and I found that the larger one triangle covers the screen the slower the program runs. Here is the result. +![Alt text](/img/small.PNG) +![Alt text](/img/medium.PNG) +![Alt text](/img/large.PNG) +![Alt text](/img/huge.PNG) +![Alt text](/img/extra huge.PNG) diff --git a/external/include/glm/glm.hpp b/external/include/glm/glm.hpp index b510380..466b45d 100644 --- a/external/include/glm/glm.hpp +++ b/external/include/glm/glm.hpp @@ -113,5 +113,5 @@ #include "matrix.hpp" #include "vector_relational.hpp" #include "integer.hpp" - +#include "gtc\matrix_transform.hpp" #endif//GLM_INCLUDED diff --git a/img/extra huge.PNG b/img/extra huge.PNG new file mode 100644 index 0000000..e021528 Binary files /dev/null and b/img/extra huge.PNG differ diff --git a/img/huge.PNG b/img/huge.PNG new file mode 100644 index 0000000..bcb48ec Binary files /dev/null and b/img/huge.PNG differ diff --git a/img/interp.PNG b/img/interp.PNG new file mode 100644 index 0000000..4b0f021 Binary files /dev/null and b/img/interp.PNG differ diff --git a/img/large.PNG b/img/large.PNG new file mode 100644 index 0000000..f230594 Binary files /dev/null and b/img/large.PNG differ diff --git a/img/medium.PNG b/img/medium.PNG new file mode 100644 index 0000000..fb9b0d4 Binary files /dev/null and b/img/medium.PNG differ diff --git a/img/phong.PNG b/img/phong.PNG new file mode 100644 index 0000000..35cbca8 Binary files /dev/null and b/img/phong.PNG differ diff --git a/img/small.PNG b/img/small.PNG new file mode 100644 index 0000000..0306dc9 Binary files /dev/null and b/img/small.PNG differ diff --git a/img/tilesize.PNG b/img/tilesize.PNG new file mode 100644 index 0000000..4d7ca9c Binary files /dev/null and b/img/tilesize.PNG differ diff --git a/img/vertices.PNG b/img/vertices.PNG new file mode 100644 index 0000000..8bf6cce Binary files /dev/null and b/img/vertices.PNG differ diff --git a/img/wire.PNG b/img/wire.PNG new file mode 100644 index 0000000..88b6828 Binary files /dev/null and b/img/wire.PNG differ diff --git a/src/Camera.cpp b/src/Camera.cpp new file mode 100644 index 0000000..96420c3 --- /dev/null +++ b/src/Camera.cpp @@ -0,0 +1,53 @@ +#include "Camera.h" + + +Camera::Camera() +{ + float fovy = 60.0f; + float zNear = 0.1f; + float zFar = 100.0f; + width = 800; + height = 600; + lightPos = glm::vec3(5, 10, 10); + lightColor = glm::vec3(1, 1, 1); + cameraPos = glm::vec3(0,0.5f,0.5f); + lookAtPos = glm::vec3(0.0f, 0.3f, 0.0f); + cameraUp = glm::vec3(0, 1, 0); + viewDirection = glm::normalize(lookAtPos - cameraPos); + projectionMatrix = glm::perspective(fovy, float(width) / float(height), zNear, zFar); + viewMatrix = glm::lookAt(cameraPos, lookAtPos,cameraUp); +} + + +Camera::~Camera() +{ +} + +void Camera::UpdatePosition(float rotationX, float rotationY) +{ + glm::mat4 Transform = utilityCore::buildTransformationMatrix(glm::vec3(0.0f), glm::vec3(-rotationY, rotationX, 0.0f), glm::vec3(1.0f)); + glm::vec4 cameraPos4(cameraPos, 1.0f); + glm::vec4 up4(cameraUp, 0.0f); + glm::vec4 newPos = Transform * cameraPos4; + glm::vec4 newUp = Transform * cameraPos4; + cameraPos = glm::vec3(newPos.x, newPos.y, newPos.z); + cameraUp = glm::vec3(newUp.x, newUp.y, newUp.z); + viewDirection = glm::normalize(lookAtPos - cameraPos); + viewMatrix = glm::lookAt(cameraPos, lookAtPos, cameraUp); +} +void Camera::UpdatePosition(float zOffset) +{ + glm::vec3 translate = viewDirection * zOffset; + glm::mat4 Transform = utilityCore::buildTransformationMatrix(translate, glm::vec3(0), glm::vec3(1.0f)); + + glm::vec4 cameraPos4(cameraPos, 1.0f); + glm::vec4 up4(cameraUp, 0.0f); + glm::vec4 newPos = Transform * cameraPos4; + glm::vec4 newUp = Transform * cameraPos4; + if (glm::length(newPos) <= 0.5) return; + + cameraPos = glm::vec3(newPos.x, newPos.y, newPos.z); + cameraUp = glm::vec3(newUp.x, newUp.y, newUp.z); + viewDirection = glm::normalize(lookAtPos - cameraPos); + viewMatrix = glm::lookAt(cameraPos, lookAtPos, cameraUp); +} \ No newline at end of file diff --git a/src/Camera.h b/src/Camera.h new file mode 100644 index 0000000..cadaf94 --- /dev/null +++ b/src/Camera.h @@ -0,0 +1,27 @@ +#pragma once +#include "glm\glm.hpp" +#include "utilities.h" + +class Camera +{ +public: + float fovy = 60.0f; + float zNear = 0.1f; + float zFar = 100.0f; + int width = 800; int height = 800; + + glm::vec3 lightPos; + glm::vec3 lightColor; + glm::vec3 cameraPos; + glm::vec3 viewDirection; + glm::vec3 lookAtPos; + glm::vec3 cameraUp; + glm::mat4 projectionMatrix; + glm::mat4 viewMatrix; + + Camera(); + ~Camera(); + void UpdatePosition(float rotationX, float rotationY); + void UpdatePosition(float zOffset); +}; + diff --git a/src/main.cpp b/src/main.cpp index 13d8e67..2c197fc 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -6,7 +6,9 @@ //------------------------------- //-------------MAIN-------------- //------------------------------- - +void mouseCallback(GLFWwindow* window, int button, int action, int mods); +void scrollCallback(GLFWwindow* window, double xoffset, double yoffset); +void updateMouse(); int main(int argc, char** argv){ bool loadedScene = false; @@ -44,6 +46,7 @@ int main(int argc, char** argv){ void mainLoop() { while(!glfwWindowShouldClose(window)){ + updateMouse(); glfwPollEvents(); runCuda(); @@ -72,6 +75,7 @@ void mainLoop() { glfwTerminate(); } + //------------------------------- //---------RUNTIME STUFF--------- //------------------------------- @@ -89,12 +93,17 @@ void runCuda(){ 1.0, 0.0, 0.0}; cbo = newcbo; cbosize = 9; - + ibo = mesh->getIBO(); ibosize = mesh->getIBOsize(); - + nbo = mesh->getNBO(); + nbosize = mesh->getNBOsize(); + cudaGLMapBufferObject((void**)&dptr, pbo); - cudaRasterizeCore(dptr, glm::vec2(width, height), frame, vbo, vbosize, cbo, cbosize, ibo, ibosize); + + glm::mat4 modelMatrix(1); + + cudaRasterizeCore(dptr, glm::vec2(width, height), frame, vbo, vbosize, cbo, cbosize, ibo, ibosize,nbo,nbosize, mainCamera,modelMatrix); cudaGLUnmapBufferObject(pbo); vbo = NULL; @@ -118,7 +127,7 @@ bool init(int argc, char* argv[]) { } width = 800; - height = 800; + height = 600; window = glfwCreateWindow(width, height, "CIS 565 Pathtracer", NULL, NULL); if (!window){ glfwTerminate(); @@ -127,6 +136,7 @@ bool init(int argc, char* argv[]) { glfwMakeContextCurrent(window); glfwSetKeyCallback(window, keyCallback); + // Set up GL context glewExperimental = GL_TRUE; if(glewInit()!=GLEW_OK){ @@ -144,7 +154,10 @@ bool init(int argc, char* argv[]) { glUseProgram(passthroughProgram); glActiveTexture(GL_TEXTURE0); - + glfwSetScrollCallback(window, scrollCallback); + mouseX = new double; + mouseY = new double; + glfwSetMouseButtonCallback(window, mouseCallback); return true; } @@ -261,11 +274,14 @@ void deleteTexture(GLuint* tex){ } void shut_down(int return_code){ + kernelCleanup(); cudaDeviceReset(); #ifdef __APPLE__ glfwTerminate(); #endif + delete mouseX; + delete mouseY; exit(return_code); } @@ -281,4 +297,44 @@ void keyCallback(GLFWwindow* window, int key, int scancode, int action, int mods if(key == GLFW_KEY_ESCAPE && action == GLFW_PRESS){ glfwSetWindowShouldClose(window, GL_TRUE); } +} +void scrollCallback(GLFWwindow* window, double xoffset, double yoffset) +{ + mouseScrollOffset += yoffset; + mainCamera.UpdatePosition(yoffset * 0.1f); +} +void mouseCallback(GLFWwindow* window,int button, int action, int mods) +{ + glfwGetCursorPos(window, mouseX, mouseY); + if (action == GLFW_PRESS) + { + if (mouseLeftBtnDown == false) + { + mouseLeftBtnDown = true; + mouseClickedX = *mouseX; + mouseClickedY = *mouseY; + } + } + if (action == GLFW_RELEASE) + { + mouseLeftBtnDown = false; + mouseDeltaX = 0.0f; + mouseDeltaY = 0.0f; + } + +} +void updateMouse() +{ + if (glfwGetMouseButton(window, GLFW_MOUSE_BUTTON_LEFT) == GLFW_PRESS) + { + glfwGetCursorPos(window, mouseX, mouseY); + mouseDeltaX = *mouseX - mouseClickedX; + mouseDeltaY = *mouseY - mouseClickedY; + rotationX = 20 * mouseDeltaX / mainCamera.width; + rotationY = 10 * mouseDeltaY / mainCamera.height; + //rotationY = glm::clamp((float)rotationY, .0f, 180.0f); + + mainCamera.UpdatePosition(rotationX, rotationY); + + } } \ No newline at end of file diff --git a/src/main.h b/src/main.h index 8999110..39b6d40 100644 --- a/src/main.h +++ b/src/main.h @@ -22,7 +22,7 @@ #include "rasterizeKernels.h" #include "utilities.h" - +#include "Camera.h" using namespace std; //------------------------------- @@ -49,12 +49,25 @@ float* cbo; int cbosize; int* ibo; int ibosize; - +float *nbo; +int nbosize; //------------------------------- //----------CUDA STUFF----------- //------------------------------- -int width = 800; int height = 800; +int width = 800; int height = 600; + +Camera mainCamera; +//mouse control +bool mouseLeftBtnDown = false; +float mouseScrollOffset = 0.0f; +double mouseClickedX = 0.0f; +double mouseClickedY = 0.0f; +double mouseDeltaX = 0.0f; +double mouseDeltaY = 0.0f; +double rotationX = 0.0f; +double rotationY = 0.0f; +double *mouseX, *mouseY; //------------------------------- //-------------MAIN-------------- @@ -62,6 +75,7 @@ int width = 800; int height = 800; int main(int argc, char** argv); + //------------------------------- //---------RUNTIME STUFF--------- //------------------------------- diff --git a/src/rasterizeKernels.cu b/src/rasterizeKernels.cu index 10b0000..f0b60bd 100644 --- a/src/rasterizeKernels.cu +++ b/src/rasterizeKernels.cu @@ -7,11 +7,19 @@ #include #include "rasterizeKernels.h" #include "rasterizeTools.h" +#include "Camera.h" +#include + +//*************Render Mode.........0 for blinn-phong, 1 for wire frame 2 for barycentric color +// 3 for vertices +//#define DEPTH_TEST +#define RENDER_MODE_2 glm::vec3* framebuffer; fragment* depthbuffer; float* device_vbo; float* device_cbo; +float *device_nbo; int* device_ibo; triangle* primitives; @@ -90,6 +98,8 @@ __global__ void clearDepthBuffer(glm::vec2 resolution, fragment* buffer, fragmen fragment f = frag; f.position.x = x; f.position.y = y; + //f.position.z = 10000.0f; + f.color = glm::vec3(1, 1, 1); buffer[index] = f; } } @@ -129,34 +139,212 @@ __global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* } //TODO: Implement a vertex shader -__global__ void vertexShadeKernel(float* vbo, int vbosize){ - int index = (blockIdx.x * blockDim.x) + threadIdx.x; - if(indexprojectionMatrix * camera->viewMatrix * modelTransformMatrix; + glm::mat4 MV = camera->viewMatrix * modelTransformMatrix; + // to eye + vertex = MVP * vertex; + normal = modelTransformMatrix * normal; + //to clip + vertex /= vertex.w; + //transform to view port + vbo[3 * index] = camera->width * 0.5f * (vertex.x + 1.0f); + vbo[3 * index + 1] = camera->height *(1- 0.5f * (vertex.y + 1.0f)); + vbo[3 * index + 2] = vertex.z; + + nbo[3 * index] = normal.x; + nbo[3 * index + 1] = normal.y; + nbo[3 * index + 2] = normal.z; } } //TODO: Implement primative assembly -__global__ void primitiveAssemblyKernel(float* vbo, int vbosize, float* cbo, int cbosize, int* ibo, int ibosize, triangle* primitives){ - int index = (blockIdx.x * blockDim.x) + threadIdx.x; - int primitivesCount = ibosize/3; - if(index 1e-6); + } +} + +__device__ bool inBoundary(float x, float y) +{ + return x >= -1.0f && x<1.0f && y >= -1.0f && y<1.0f; +} + +__device__ int inBoundary(int x, int y, glm::vec2 resolution) +{ + return x >= 0 && y >= 0 && xdepthbuffer[targetIdx].position.z) + { + depthbuffer[targetIdx].position.z = depth; + depthbuffer[targetIdx].color = color; + depthbuffer[targetIdx].normal = normal; + } + + } + } + + } } //TODO: Implement a fragment shader -__global__ void fragmentShadeKernel(fragment* depthbuffer, glm::vec2 resolution){ - 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){ - } +__global__ void fragmentShadeKernel(fragment* depthbuffer, glm::vec2 resolution, Camera *camera){ + 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){ + + + float specular = 30.0; + float ka = 0.1; + float kd = 0.7; + float ks = 0.2; + glm::vec4 lightPos4(camera->lightPos, 1.0f); + glm::vec4 lightPosEye4 = camera->viewMatrix * lightPos4; + glm::vec3 lightPosEye(lightPosEye4.x, lightPosEye4.y, lightPosEye4.z); + glm::vec4 dbp4(depthbuffer[index].position, 1.0f); + glm::vec4 dbpEye4 = glm::inverse(camera->projectionMatrix) * dbp4; + glm::vec3 dbpEye = glm::vec3(dbpEye4.x, dbpEye4.y, dbpEye4.z); + + glm::vec3 lightVectorEye = glm::normalize(lightPosEye - dbpEye); // watch out for division by zero + glm::vec3 normal = depthbuffer[index].normal; // watch out for division by zero + glm::vec4 normal4(normal, 1.0f); + glm::vec4 normalEye4 = camera->viewMatrix * normal4; + glm::vec3 normalEye = glm::normalize( glm::vec3(normalEye4.x, normalEye4.y, normalEye4.z)); + float diffuseTerm = glm::clamp(glm::dot(normal, normalEye), 0.0f, 1.0f); + + glm::vec3 R = glm::normalize(glm::reflect(-lightVectorEye, normal)); // watch out for division by zero + glm::vec3 V = glm::normalize(-dbpEye); // watch out for division by zero + float specularTerm = pow(fmaxf(glm::dot(R, V), 0.0f), specular); +#ifdef RENDER_MODE_2 // barycentric color + glm::vec3 materialColor = depthbuffer[index].color; + depthbuffer[index].color = materialColor; + if (depthbuffer[index].position.z < -1.0f) + { + depthbuffer[index].color = glm::vec3(0.4f, 0.4f, 0.4f); + } +#endif +#ifdef RENDER_MODE_0 // blinn-phong + glm::vec3 materialColor = glm::vec3(0.2f, 0.4f, 0.7f); + glm::vec3 color = ka*materialColor + glm::vec3(1.0f) * (kd*materialColor*diffuseTerm + ks*specularTerm); + depthbuffer[index].color = color; + if (depthbuffer[index].position.z < -1.0f ) + { + depthbuffer[index].color = glm::vec3(0.4f, 0.4f, 0.4f); + } +#endif +#ifdef RENDER_MODE_1 // wire frame + float eps = 0.04f; + glm::vec3 materialColor = depthbuffer[index].color; + if (depthbuffer[index].color.x 1- eps) + { + if (depthbuffer[index].color.y < eps && depthbuffer[index].color.z < eps) + isVertex = true; + } + else if (depthbuffer[index].color.y > 1 - eps) + { + if (depthbuffer[index].color.x < eps && depthbuffer[index].color.z < eps) + isVertex = true; + } + else if (depthbuffer[index].color.z > 1 - eps) + { + if (depthbuffer[index].color.y < eps && depthbuffer[index].color.x < eps) + isVertex = true; + } + if (isVertex) + depthbuffer[index].color = glm::vec3(1); + else + { + depthbuffer[index].color = glm::vec3(0.4f, 0.4f, 0.4f); + } + + if (depthbuffer[index].position.z < -1.0f) + { + depthbuffer[index].color = glm::vec3(0.4f, 0.4f, 0.4f); + } +#endif + +#ifdef DEPTH_TEST + depthbuffer[index].color = glm::vec3(0.5 * (depthbuffer[index].position.z + 1)); +#endif +#ifdef NORMAL_TEST + depthbuffer[index].color = glm::normalize(depthbuffer[index].normal.x * glm::vec3(1, 0, 0) + depthbuffer[index].normal.y * glm::vec3(0, 1, 0) + depthbuffer[index].normal.z * glm::vec3(0, 0, 1)); + if (depthbuffer[index].position.z < -1.0f) + depthbuffer[index].color = glm::normalize(-(camera->viewDirection.x * glm::vec3(1, 0, 0) + camera->viewDirection.y * glm::vec3(0, 1, 0) + camera->viewDirection.z * glm::vec3(0, 0, 1))); +#endif + } } //Writes fragment colors to the framebuffer @@ -172,13 +360,19 @@ __global__ void render(glm::vec2 resolution, fragment* depthbuffer, glm::vec3* f } // 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, Camera camera, glm::mat4 modelTransformMatrix){ + clock_t timeBegin, timeEnd; + timeBegin = clock(); // set up crucial magic int tileSize = 8; dim3 threadsPerBlock(tileSize, tileSize); dim3 fullBlocksPerGrid((int)ceil(float(resolution.x)/float(tileSize)), (int)ceil(float(resolution.y)/float(tileSize))); + //set up camera + Camera *cudaCamera = NULL; + cudaMalloc((void**)&cudaCamera, sizeof(Camera)); + cudaMemcpy(cudaCamera, &camera, sizeof(Camera), cudaMemcpyHostToDevice); + //set up framebuffer framebuffer = NULL; cudaMalloc((void**)&framebuffer, (int)resolution.x*(int)resolution.y*sizeof(glm::vec3)); @@ -214,32 +408,37 @@ void cudaRasterizeCore(uchar4* PBOpos, glm::vec2 resolution, float frame, float* cudaMalloc((void**)&device_cbo, cbosize*sizeof(float)); cudaMemcpy( device_cbo, cbo, cbosize*sizeof(float), cudaMemcpyHostToDevice); - tileSize = 32; + device_nbo = NULL; + cudaMalloc((void**)&device_nbo, nbosize*sizeof(float)); + cudaMemcpy(device_nbo, nbo, nbosize*sizeof(float), cudaMemcpyHostToDevice); + + tileSize = 8; int primitiveBlocks = ceil(((float)vbosize/3)/((float)tileSize)); //------------------------------ //vertex shader //------------------------------ - vertexShadeKernel<<>>(device_vbo, vbosize); + vertexShadeKernel<<>>(device_vbo, vbosize, device_nbo, nbosize, cudaCamera, modelTransformMatrix); 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, primitives, cudaCamera); cudaDeviceSynchronize(); //------------------------------ //rasterization //------------------------------ - rasterizationKernel<<>>(primitives, ibosize/3, depthbuffer, resolution); + rasterizationKernel<<>>(primitives, ibosize/3, depthbuffer, resolution, cudaCamera); cudaDeviceSynchronize(); //------------------------------ //fragment shader //------------------------------ - fragmentShadeKernel<<>>(depthbuffer, resolution); + // glm::vec3 lightPosition(0, 0, 0); + fragmentShadeKernel << > >(depthbuffer, resolution, cudaCamera); cudaDeviceSynchronize(); //------------------------------ @@ -251,8 +450,11 @@ void cudaRasterizeCore(uchar4* PBOpos, glm::vec2 resolution, float frame, float* cudaDeviceSynchronize(); kernelCleanup(); - + cudaFree(cudaCamera); + timeEnd = clock(); checkCUDAError("Kernel failed!"); + + printf("One frame completed in %d ms\n", timeEnd - timeBegin); } void kernelCleanup(){ @@ -262,5 +464,6 @@ void kernelCleanup(){ cudaFree( device_ibo ); cudaFree( framebuffer ); cudaFree( depthbuffer ); + cudaFree(device_nbo); } diff --git a/src/rasterizeKernels.h b/src/rasterizeKernels.h index 784be17..1b8c7a0 100644 --- a/src/rasterizeKernels.h +++ b/src/rasterizeKernels.h @@ -9,8 +9,8 @@ #include #include #include "glm/glm.hpp" - +#include "Camera.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, Camera camera, glm::mat4 modelTransformMatrix); #endif //RASTERIZEKERNEL_H diff --git a/src/rasterizeTools.h b/src/rasterizeTools.h index e9b5dcc..ba52572 100644 --- a/src/rasterizeTools.h +++ b/src/rasterizeTools.h @@ -9,6 +9,7 @@ #include "utilities.h" #include "cudaMat4.h" + struct triangle { glm::vec3 p0; glm::vec3 p1; @@ -16,6 +17,11 @@ struct triangle { glm::vec3 c0; glm::vec3 c1; glm::vec3 c2; + glm::vec3 n0; + glm::vec3 n1; + glm::vec3 n2; + + bool visible; }; struct fragment{ diff --git a/windows/PROJ4_Rasterizer/PROJ4_Rasterizer/PROJ4_Rasterizer.vcxproj b/windows/PROJ4_Rasterizer/PROJ4_Rasterizer/PROJ4_Rasterizer.vcxproj index f640485..d711f27 100644 --- a/windows/PROJ4_Rasterizer/PROJ4_Rasterizer/PROJ4_Rasterizer.vcxproj +++ b/windows/PROJ4_Rasterizer/PROJ4_Rasterizer/PROJ4_Rasterizer.vcxproj @@ -1,5 +1,5 @@  - + Debug @@ -19,16 +19,18 @@ Application true MultiByte + v120 Application false true MultiByte + v120 - + @@ -67,6 +69,7 @@ + @@ -77,6 +80,7 @@ + @@ -87,6 +91,6 @@ - + \ No newline at end of file diff --git a/windows/PROJ4_Rasterizer/PROJ4_Rasterizer/PROJ4_Rasterizer.vcxproj.filters b/windows/PROJ4_Rasterizer/PROJ4_Rasterizer/PROJ4_Rasterizer.vcxproj.filters index 6a1d8cf..2a16476 100644 --- a/windows/PROJ4_Rasterizer/PROJ4_Rasterizer/PROJ4_Rasterizer.vcxproj.filters +++ b/windows/PROJ4_Rasterizer/PROJ4_Rasterizer/PROJ4_Rasterizer.vcxproj.filters @@ -30,6 +30,9 @@ Header Files + + Header Files + @@ -47,6 +50,9 @@ Source Files + + Source Files +