diff --git a/README.md b/README.md index 41b91f0..f9e1bd1 100644 --- a/README.md +++ b/README.md @@ -1,18 +1,95 @@ CUDA Rasterizer =============== -[CLICK ME FOR INSTRUCTION OF THIS PROJECT](./INSTRUCTION.md) +**University of Pennsylvania, CIS 565: GPU Programming and Architecture** -**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 4** +**Anantha Srinivas** +[LinkedIn](https://www.linkedin.com/in/anantha-srinivas-00198958/), [Twitter](https://twitter.com/an2tha) -* (TODO) YOUR NAME HERE - * (TODO) [LinkedIn](), [personal website](), [twitter](), etc. -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +**Tested on:** +* Windows 10, i7-8700 @ 3.20GHz 16GB, GTX 1080 8097MB (Personal) +* Built for Visual Studio 2017 using the v140 toolkit -### (TODO: Your README) +# Introduction -*DO NOT* leave the README to the last minute! It is a crucial part of the -project, and we will not be able to grade you without a good README. +![](renders/pipeline.png) + +This implementation of Rasterization is done on CUDA and split into these five important steps: + +* Vertex Transform and Assembly - This acts as a vertex shader and transforms the input vertices from world space to NDC space and then into screen space. Screen space transformation is done because we are deadling with pixels on the framebuffer. + +* Primitive Assembly - From the list of vertices, we construct primitives containing 1/2/3 vertices depending upon the geometry(point, line or triangle). To select the vertices, index buffer is used. + +* Core Rasterization - This performs a scan of pixel in the small rectangular boundary around each primitive. Then we check if the pixel lies inside or outside this primitive and set initialize the fragment accordingly. We also perform a depth check to ensure that the fragment is valid. + +* Pixel Shading - This acts as the pixel shader, where we calculate the color from the fragment. + +* Updating frame buffer - The frame buffer for each frame is updated to the screen. Normally in modern graphics API, there might be two or more framebuffer, so that while one is updating other one is shown. Also, if any super sampling is used here, the framebuffer is downscaled accordingly. + + +# Features + +The current version of the rasterizer supprots the following features: + +* Vertex Shading and primitive assembly with depth testing + +| Color | Normal | Depth | +| ------------- | ------------- | ------------- | +|![](renders/color.PNG) | ![](renders/normals.PNG) | ![](renders/depth.PNG) + +Simple, basic implementations of vertex transform, normal calculation and transform and depth checking. + +* Instancing + +![](renders/instancing.PNG) + +This is implemented similar to modern graphics APIs. Each vertex shader of each instance has the same pointer to the input vertex buffer. The output vertices are modified here based on the instance ID by offsetting in screen space. + +* Texture Mapping (2D) with bilinear filtering + +![](renders/render_cesium.PNG) + +Similar to vertex and normals, texture coordinates is calculated per pixel. It is then corrected to account for depth by performing perspective correction. + +In bilinear filtering, we sample not only the current texture coordinate but also the neighbours. This is then accompanied by mixing all the sampled points in a ratio determined by how far off is the texture coordinate from the sampled point. + +* SuperSample Antialiasing + +| No SSAA | SSAA 2 | +| ------------- | ------------- | +|![](renders/ssaa_1.PNG) | ![](renders/ssaa_2.PNG) | + +SuperSample AntiAliasing is the most basic implementation of AA. Here, the framebuffer size is incrased by the level of AA. We perform all the calculation with this enlarged framebuffer and then downscale it when updating to screen. It can be expensive both computationally and memory. + +* Color interpolation between points on a primitive + +![](renders/render_di.PNG) + +Interpolating color between vertices of a primitive using the bary centric coordinates. + + +* Points and Lines + +| Triangles | Lines | Points | +| ------------- | ------------- | ------------- | +| ![](renders/triangles.PNG) | ![](renders/lines.PNG) |![](renders/points.PNG) + +# Performance Analysis + +![](renders/performance_analysis.png) + +**Testing Conditions** +* Launched in Release Mode +* NVIDIA V-Sync turned off +* Run with NSight Profiler + +**General Analysis** +* From the performance graph is is fairly evident that the most expensive component of the rasterizer (in terms of time) is the raterizer. This is because of the fact that for each primitive, we need to scan pixel by pixel (computationally O(n^2)). The obvious way to minimize this would be to have some hardware implementations for some of the sub-functions. Other than that, we could load up the framebuffer and depth buffer into shared memory. Although not implemented, this could potentially speed up the kernel in magnitudes of 10X. +* We also see that the Render step, where we perform pixel shading is fairly same across all models. This is because the kernel is launched per pixel (number which remains same across all models). The calculation inside this pixel is fairly short. + +# Build Command + +`cmake .. -G "Visual Studio 15 2017 Win64" -DCUDA_TOOLKIT_ROOT_DIR="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.2" -T v140,cuda=9.2` ### Credits diff --git a/Report/New Microsoft PowerPoint Presentation.pptx b/Report/New Microsoft PowerPoint Presentation.pptx new file mode 100644 index 0000000..052a8cf Binary files /dev/null and b/Report/New Microsoft PowerPoint Presentation.pptx differ diff --git a/gltfs/BoxTextured.gltf b/gltfs/BoxTextured.gltf new file mode 100644 index 0000000..4a6ab9c --- /dev/null +++ b/gltfs/BoxTextured.gltf @@ -0,0 +1,181 @@ +{ + "asset": { + "generator": "COLLADA2GLTF", + "version": "2.0" + }, + "scene": 0, + "scenes": [ + { + "nodes": [ + 0 + ] + } + ], + "nodes": [ + { + "children": [ + 1 + ], + "matrix": [ + 1.0, + 0.0, + 0.0, + 0.0, + 0.0, + 0.0, + -1.0, + 0.0, + 0.0, + 1.0, + 0.0, + 0.0, + 0.0, + 0.0, + 0.0, + 1.0 + ] + }, + { + "mesh": 0 + } + ], + "meshes": [ + { + "primitives": [ + { + "attributes": { + "NORMAL": 1, + "POSITION": 2, + "TEXCOORD_0": 3 + }, + "indices": 0, + "mode": 4, + "material": 0 + } + ], + "name": "Mesh" + } + ], + "accessors": [ + { + "bufferView": 0, + "byteOffset": 0, + "componentType": 5123, + "count": 36, + "max": [ + 23 + ], + "min": [ + 0 + ], + "type": "SCALAR" + }, + { + "bufferView": 1, + "byteOffset": 0, + "componentType": 5126, + "count": 24, + "max": [ + 1.0, + 1.0, + 1.0 + ], + "min": [ + -1.0, + -1.0, + -1.0 + ], + "type": "VEC3" + }, + { + "bufferView": 1, + "byteOffset": 288, + "componentType": 5126, + "count": 24, + "max": [ + 0.5, + 0.5, + 0.5 + ], + "min": [ + -0.5, + -0.5, + -0.5 + ], + "type": "VEC3" + }, + { + "bufferView": 2, + "byteOffset": 0, + "componentType": 5126, + "count": 24, + "max": [ + 6.0, + 1.0 + ], + "min": [ + 0.0, + 0.0 + ], + "type": "VEC2" + } + ], + "materials": [ + { + "pbrMetallicRoughness": { + "baseColorTexture": { + "index": 0 + }, + "metallicFactor": 0.0 + }, + "name": "Texture" + } + ], + "textures": [ + { + "sampler": 0, + "source": 0 + } + ], + "images": [ + { + "uri": "data:image/png;base64,iVBORw0KGgoAAAANSUhEUgAAAQAAAAEACAYAAABccqhmAAAEDWlDQ1BJQ0MgUHJvZmlsZQAAOI2NVV1oHFUUPrtzZyMkzlNsNIV0qD8NJQ2TVjShtLp/3d02bpZJNtoi6GT27s6Yyc44M7v9oU9FUHwx6psUxL+3gCAo9Q/bPrQvlQol2tQgKD60+INQ6Ium65k7M5lpurHeZe58853vnnvuuWfvBei5qliWkRQBFpquLRcy4nOHj4g9K5CEh6AXBqFXUR0rXalMAjZPC3e1W99Dwntf2dXd/p+tt0YdFSBxH2Kz5qgLiI8B8KdVy3YBevqRHz/qWh72Yui3MUDEL3q44WPXw3M+fo1pZuQs4tOIBVVTaoiXEI/MxfhGDPsxsNZfoE1q66ro5aJim3XdoLFw72H+n23BaIXzbcOnz5mfPoTvYVz7KzUl5+FRxEuqkp9G/Ajia219thzg25abkRE/BpDc3pqvphHvRFys2weqvp+krbWKIX7nhDbzLOItiM8358pTwdirqpPFnMF2xLc1WvLyOwTAibpbmvHHcvttU57y5+XqNZrLe3lE/Pq8eUj2fXKfOe3pfOjzhJYtB/yll5SDFcSDiH+hRkH25+L+sdxKEAMZahrlSX8ukqMOWy/jXW2m6M9LDBc31B9LFuv6gVKg/0Szi3KAr1kGq1GMjU/aLbnq6/lRxc4XfJ98hTargX++DbMJBSiYMIe9Ck1YAxFkKEAG3xbYaKmDDgYyFK0UGYpfoWYXG+fAPPI6tJnNwb7ClP7IyF+D+bjOtCpkhz6CFrIa/I6sFtNl8auFXGMTP34sNwI/JhkgEtmDz14ySfaRcTIBInmKPE32kxyyE2Tv+thKbEVePDfW/byMM1Kmm0XdObS7oGD/MypMXFPXrCwOtoYjyyn7BV29/MZfsVzpLDdRtuIZnbpXzvlf+ev8MvYr/Gqk4H/kV/G3csdazLuyTMPsbFhzd1UabQbjFvDRmcWJxR3zcfHkVw9GfpbJmeev9F08WW8uDkaslwX6avlWGU6NRKz0g/SHtCy9J30o/ca9zX3Kfc19zn3BXQKRO8ud477hLnAfc1/G9mrzGlrfexZ5GLdn6ZZrrEohI2wVHhZywjbhUWEy8icMCGNCUdiBlq3r+xafL549HQ5jH+an+1y+LlYBifuxAvRN/lVVVOlwlCkdVm9NOL5BE4wkQ2SMlDZU97hX86EilU/lUmkQUztTE6mx1EEPh7OmdqBtAvv8HdWpbrJS6tJj3n0CWdM6busNzRV3S9KTYhqvNiqWmuroiKgYhshMjmhTh9ptWhsF7970j/SbMrsPE1suR5z7DMC+P/Hs+y7ijrQAlhyAgccjbhjPygfeBTjzhNqy28EdkUh8C+DU9+z2v/oyeH791OncxHOs5y2AtTc7nb/f73TWPkD/qwBnjX8BoJ98VVBg/m8AAAAJcEhZcwAACxMAAAsTAQCanBgAAAFZaVRYdFhNTDpjb20uYWRvYmUueG1wAAAAAAA8eDp4bXBtZXRhIHhtbG5zOng9ImFkb2JlOm5zOm1ldGEvIiB4OnhtcHRrPSJYTVAgQ29yZSA1LjQuMCI+CiAgIDxyZGY6UkRGIHhtbG5zOnJkZj0iaHR0cDovL3d3dy53My5vcmcvMTk5OS8wMi8yMi1yZGYtc3ludGF4LW5zIyI+CiAgICAgIDxyZGY6RGVzY3JpcHRpb24gcmRmOmFib3V0PSIiCiAgICAgICAgICAgIHhtbG5zOnRpZmY9Imh0dHA6Ly9ucy5hZG9iZS5jb20vdGlmZi8xLjAvIj4KICAgICAgICAgPHRpZmY6T3JpZW50YXRpb24+MTwvdGlmZjpPcmllbnRhdGlvbj4KICAgICAgPC9yZGY6RGVzY3JpcHRpb24+CiAgIDwvcmRmOlJERj4KPC94OnhtcG1ldGE+CkzCJ1kAAEAASURBVHgB7b13eGRXdif2q1wooAo5NzrnHBmazTjkcDgckjPiiErjTzMre72W/cl/eL2fLa/Xtj5bK8sray3tSJrIGc4ORQ7JYQ7dzSabbHYgO+cckDNQQAGVg3/nFYqNbqKBKtSriHvJalR47757z73n3JOP4cyZMzGopiCgIDAnIWCck7NWk1YQUBDQIKAIgNoICgJzGAKKAMzhxVdTVxBQBEDtAQWBOQwBRQDm8OKrqSsIKAKg9oCCwByGgCIAc3jx1dQVBMwKBIUPAQOnYDDE3TnkvbTE38SbLz/Hf550Ab+YwhPky6/kDW+OTXyR+D7+2TDVrYknqL8FAAFFAApgkWSIcQSOxf9OwuZozIAQsTIUMSISNSDC9/JdlFfGiKXRqHyOI7B8TjT5Lo7V7Cz+P4lI/Ff5a+A/Rv41kkeU20z8ayKRMU18J+/NfG82xmAkGdB65j/aX3l24kHqb15DQBGAvF6eOEIJAgrCRgXVJt7H/8YQjAAjISPcQRO8YQN88oqYEIiZEAxHEODvgbC8ouDHL5HdG4ySYJCgENGJx7AQs+UlzWwCbGYjX/zL96QhcJpjKDNHYeffEn7vMEXhtPA7SxQWEodEP0YSBiEgQjwmutP6VP/kJwQUAcjLdZGTXv6LEZGN8BDBRwIGDAaI6GEbBsciGPVHMU7s9od41hPZYjFB00QjxrNNdwoTr29B0BiJgxATaUG+vNq7m//EScOkz9ozhViYUEYqUe4wwWUD6kvCJBJAjT2KalsENkOU41AcwU3I5dc7RQByuB6CFnJSCgLLaeknCz/oN6HPZyKymzDijcHtDcNHHj/CIz9EJI/E/Ajx5I5orD1Zb8Fyvnj7zXbLh/jXU3x18/o7vJtMQLT3t30hH6OUOdy+CDwkRjKHa+QGZE4WcgJ2EgdniRU1TqNGDCptURKGCDkH4T7iXI3Gydw6+juMRn2dCQgoApAJqN6hzwTCC6JEifXesAm943LCmzA8FtVO9TEe3p5wDOPk1wXx/fwscrxGKOQsFRab/Wsv/iPEI1PtK13f9kXiY3SCIMk4xieIRIIfsfqiKBuLoVREB+oLXJQXKh1mVJQaSQgiqCKnIATBzHkJUVAEIVOrOXW/igBMDRfdvhUFmSCpvPwRA0Z5snuCBowFgGF/DD0+A9whA/rGDBjn93K9mWyzILy8RNFmEH49j5s2v4nxyZhvNiohhUOgPDFEwqApKEm6HFYD6soMqLAYUM33VSUUIewGOG1xnYKDugbhahQxuAnJTL1TBCADkBUckJM6Qm28j9p5PxVzoogbJtK3jpnQzlc3T/wxHu82KtNMVLOJlr2sCFcjQchuglmIQgRdwxG0E8nFgmGh1rGJYsJCVwyN1CHUlRhQYgGsphjhQ4JIzkGIQZz3udmTepc+BIpwy6UPlHR64D6Ns7I86cZ4sneMm3HJbcZ1dwwef4gnfIQnO1+8rtyazpMK916NKJCrIY7DzhdJJQY8EfSOagc/rBYLVlQbSBCimOcIa2KCzFZEJ+E2VNMPAgaVECQ9YArCy55MsPrC5l8cseAqkb5zJExNvWjBJxRe6T1qTt2tITtnbKfOoNFpwsqaCBY44wpEgbX4OqiWPgQUBzBrGMadYoQ9HaUSr9tnQ9uoiaxtkDZ5UeKFyfbLRhXioE6uVMFMtwWNcAYJS787jIFx4Dh9ExrINi2uiKKlNKQpD+U6EbVUmx0EFAFIEW5y+oi8Lpuulwq9Aa+JJ30U3TTZDfgjNInFuQHxlBN2VU4y1VKHwE3FIp2dqD/pC8WViL3kqHo81BXQ76DBaeErghr6G4gVQawlEfJjCuTJw1sRgCRhlWDxvVTqjXiNGCKi3xgzotVj1DZkmCe+eM3ZblWDJ9m7umw6CCQUiaIz8NPf4LIvhituE2poSlxWEaOegI5HpTFU2mOah6JwXUphOB1Eb/6mCMBNWEz5Tk4TkeFFmy+ecqLBPz1owqVharB5NFmMETq9GKi4UufOlADU+UvhvkpIEcTzcWQ8igMePoAOR6uqgfU1IhqIG7O4MMuqqTYTBJQScBoIaVuI/9DbFufcFpzuJ7s/GkGYgqcEySiUnwZ4Wf5JnIgMXJQmlxGb6qNYUR7SOLKbokSWB1Qgj1MEYIqFEsS2iHIvbMRFQfxeuuh6GVhDpZ44swhhUMg/BeBy+FViTURcsFPx4rQZsbkJWFkZhssc0Yi4Egu+ukBKBJgEE3HVlcg2DxH/7JAVl4egOewMMxBHIuf4k6bYU8g/CWh58jaxJsIJeLhWY+TSDrbT8WrYjGVVZiypoNWAHobCzSlCcHPRFAEQpOaWkJPDHzXiKm347e4oLruBLjq2iyurldyA0ubf3DT5/E6zvHCAkvugl6bDfrpb99Pdun/chJYKC1rKQuQQ6JtBQiA5E+Z6m/MEQOLXg3TeGWaobRe1+seo4Gsd5gaJiFafJqU898Of6xv4TvMXQiBWGVEWto/E0MG1XUAisKkqioYyo+ZdKIRgrjsUzWkCIHKjJNDo5MY43mfCyT6KAIyGZ3wKDBKFo1rBQ0AIgZ1rKRxBx1AQlwcM9Cq04K76CBa5QprIJ9zfXG1zmgCMEfm/6LHiVB8TYDBap0SEfNWKEgJCCISmOynOtQ+H0Etz7upqG+5uCtGRiPLAHG1zjgCIh16AiTdOD9lwrpcBKJp2X2n159L+lwM/RCXhuX5ggKLB+iYL1lUF6MQVj+BMWBTmAkzmDAEQxBelT5fXjFP9Ep0XRT/j1IU1FBZQTgjV5gYE4mvNg4Auxh1jzJfYGWWsgQ3ra8NocDCzEZXCYbp6z4UtUfQEQDT8jCFh4kzG4Y8acG7ASKceIMykmGLym8vy39xA9zvPUrwKJc9At4eHAZXAY3ytpkdhS3kM5cxWJCbDYucNi54ASNCO2ycOPSYc6QGJAFNU0THEqux6d8aMOfSLcANWWgui4RD1QQYGdRmxlYFHKyqoL7DGMzMVMziK2hNQQkUluebeLhuuDYY0F17hBlRTELgTBMJ0IjIzQ9HiagseagqgipGGwikWayvKqUnknrjsdtC89+ZVC1oHmeiasaIK+Yt1G+s3L7PIhNwrsmdevWjGlREzglQai79IMbaiIgCyROK1NxY14VifBe9fMWBoPAyh6qopCKQCgTBPkBHmd9hzHThBpfEYi63I3iq2VjQ6ANHYSgBPp9dCRZ8BZ/rpCsokHXaSOKXhL7Ztm/n5yJ4RC1E/8z4c7qLDGJXIa2oZbVgSKqp4gqIgAKLpFxNf+5hF8+g70y+FKuIlrDK/VdQTihUCQgTsVBD2MiaE9VmY5s2EDTWsflQaz0BUDPMuCgIgMlo/U3N9TK++GwNBinAK+Ythc+bLHKQWop/5y4/1RJj70UrlYBDzyiIkDhNFUfNloLMYR4HrAOIOG60eC167ZkVrfwAGsm1K2TeLnaBumRYCsqdEudzrDuD1q2a0eUysaSCCZ2HrBQqWAMhihMn2nx00Y+cNOnF4g/FYfVkT1RQEMgAB2VoiFgSDIexqNWl7T/ag7MVCbQVJAETmDxHwJ3rMONjJUF4WpxSPLtUUBLICAe412XOy9070mrkXjdyNhbkBC0oHkDjcA0zcIfn5jjJVVw8TQ1ppo1Wa/qxsffUQQkD2mpEnjuy9I91E/RiVg3WSETpOBAqJFBQMARDkF9dsKbd1ddiIzzqZnps54hXyK5zMBQSECNh46vcwoGxfh4EmaCOWVkZRaomnjisUIlBAIgCr7RD5LwyZ8Q6L0I8xP7zUoFcnfy62v3qmBgEhAjz1fYwaept78hyzSUnl50JBfplDQRAAAamw/VcZ0LOnzaBV01VRfAoJ8wUCshfN5E/fu8b08SQCXiaVLRTX4bwnAAltv8j8n3Swtm6Y1TlUUxDIQwgIR/ppO+gvYISXrsOFQATynAAwYSf1q8ep7T/CUE23X+VxzcN9r4Y0AQHhBBg+wNgB7tlu44SfQH6DJ68JgGRlOTtgwZFeI7rpjmmmKKBk/vzeUHN9dBIwNMjahcf6DMw1adYyC+UzTPKSAIjML7H8bfTwE23/gFdp+/N5E6mx3YSA+KOIKDDEegQisl4foccg93K+qgbzkgAIo9/vM+GjTgs8/rDmaaVO/pubTL3LbwjIXpW8gj5Wk33zmpnBREYtWC0f3YbzjgAIC9XLxJ0fd9vQNxLQPPwU8uf3hlej+yoEEns2GKDbcIedyUfNeRmjklcEQJBf4vmPU+N/nVF9JKKqKQgUNgTIDXSPBHGk34o2r51JRcSdLX9a3hAAMfe5GW+tJfPoY8kmhvQmqGj+gEuNREEgNQgQ/1mHMIoLA2FcYJIad9icV8FDeUMAJLTy3AAjrAgkD118VUhvahtNXZ2/EDDRPuhnEYILAyxGwj0eDyPOj/HmnAAIhYwn8DTjWDfQxxRMkoBBNQWBYoKA1CeUvS17vJ1lyWTP54OMm3MCIHAQH/9POiys0MtS3DShqKYgUIwQkL09SPPg7laLtufFOJjrllMCYCZARoNG7O0uQe9oCBEl9+d6P6jnZxACotOKcY+7mWBQ9rwnZKSom9u0YjljtsVPepz+0jeYWul8Xyiu9Msg8FXX+Q8BcaKJ8p8IFeWiK5esvAZijbjYSiEnkaULvQkRkDT1F7jnG0uYaZilyJysVCUVrHLRckgAWKhzzKAlVPDTYULKM6k2NyEgnJ8/TMTn9MuJFI1lVjisRg3pBVnGWcexzxPCMB3trcQTkacLnRj4OafDjBdwcS4rq3OD/LLbckIAhPXv95txgaGTV4ajcFrm5saf67OWEz9AbVgZc28vq7NiQYUVdS4LykkA7KzHZeSJL8RBPOrcniCGxsJoo029dTiIURIDO9mCQjUVW3jgSZ3Kc4NGVJQAzaVSwCb7hCAnBEDMIKcHWKZ7CHCYhNnL/sTnOvLlev6i6hXknVduwYrGUty12IW1zQ4m2LjzXiAtwKmucXx+ZRRnO8cZdBPWfMXufEeuZzn988vMMZxldqtyVq+pL8lNmHtOCMClEQuuDkXgo0bUpqr0Tr9LivBXQX4Tsd9Fdv8HOxqxqtGRVLVmIRgbm0uxrqkUp9vG8KP9PRibIAKFCCYjJxQIhdE6YsBFcj4rylmWOMstq1YAWfignP4s090zLqW8sjxb9bicQ0BOcWlLam348280EflLkkJ+uSdx0suZsXqeA//60XrMq7JonESiX7mukBpVHegeA453MfdFDhSBWUFBWXPR+ot290ifDZ2csMh203B7hbSGaqwpQMBPmX9VXQn+YFsdGisF+We3BeW+BTWl+Bfb69DM0zNUoBRAuBrBBTkQBTckDF5wZYJOpgDZ2V06O+in+Czx8w9RwdHLEN+jPayzxjJLWXlwiuNUl2cWApI8c1GVDduXlWN5g4MneuJMn91zzczEu6y+DA+tqCQRsDI5Z7bQZnbjvdNdwtGwzICGG4IjkggnW8VGsoKHctJ76PBzkn7QI+OhCfvuncChvi9GCGg2fu6DLS2l2LzQqeMUDbh3WQXWUpRgsn4tfFzHzrPSlXAB4iAkuHGSkbBuPwuN8LtstKwQgAgpfd+4gXbPuDNHtiaXDQCqZyQHgQgpQCNP6RX1JaiQkrs6NtGiL2uwo5kWhUIWBcS34WBnDB0eI4LMgp2NLEIZJwDi6tg1bsaZYRocIrkxdei411RXs4RAiLL/Gmr7q1y2WfYw/W3N1Q4srXUgGI+ymf7iPP7VQneo0yPUk3mtmptwpoeacQIQoOzf5o7hPEMhldY/08uZv/0HyOKu5unfWG7NyCAXVVuxvJpBNvQoLORmJhdwpT+ITncYgVjG0TOzujgLT//2MSuujZpo72SmvyzJNYW8AYp17OLrXltm0dx4MzFHk4EONSUWmKlRK2QSIDgSIad8gQ5Cl4ctPDQzO5uMkRgZti9iwCV6+7XT7GfTV+zLxB5SfWYQAjb6vMsrk81iNsEpGy2zOJPJKWh9i4uz4MzFQSYSIQ5lcjoZIwAy6rYxC9qGIxin37ay+Wd83+T1A4S1zXSTR4hnaSYRJtNzkP6FC4jSIaB/LMJoWQbKZHBCGSMAoos5zaIebsY8q9M/G9smv5/hpQgoDi+ZbFrgEBEn86Qmk7OI9y3RsQM0nR/pleIimXue7gRA6K94M4lDQyejtvxc+CwQ/8xBSPWsDwRoBvSwtFsG9zJCdDWVfJLFoGsSnAkRkSQM+jo56SCV6ZngbXQlALK42sA52JODVlZJvem/rc8uUr0UKgQs3BgSyjskmyIDTRC/m1mlsuVBl4EpfKVLQU7xbjzSw0SitKCL17TeBFRXAiCdSajvoM+gZUAVp4xioMZfWRn1RcoQsJIAXOrzYpBx/Zlo7cMBXB0Mwi4nUJE0zSJAjG8fiaLHa6JCkDkSdCYB+hIABjGIy++lETM8PnH5VRxAkezFtKdhoXLuQq8X3cP+tPuaqoMbJC4X+30ZtzRM9exMfxemWfAcHelGAsySpDMXoCsBEIo16KPPPzP9CMunTv9Mb43C6V8O5t5R2re7fWTV9eUC+iknX+4mcRkJ0XuueDiAxOryXGXmIGCInLWIAHrOUDcCIEMbDxvRzzLe/VxghfyJ5VN/ExCoZAzAASby+PQCnUN0YmUFId452YdT5C5cElxfhE1wSTjq7jEGDAVNuooBukFMfP57/DZc99pY/0yWRTUFgVshIOyrlgyz1YsPzrlv/XEWn8Sq+MG5YZzq9MFLJaD0X6xNagpcH7fQumbWdZ66gYyuPuhyR3BtOKyx/8W6EGpe6UFAUoH1kmXfe8GNd0/1M3hndgFiYyy19fapQew+O4SB8TA5znSzC6Q3r0zfLaJNK/GrhwrBqI7stS45AYX9HwpIHXSKAbT10uVbNQWBKSEgIro47LQNBfD2GdaDkAxBzU60VNs1L74pb5r0pWQRbh/w4XirB3uuejBEYiLIUYSi/6RZx70Dxaemh2H1vV4zam1UsuugDdCFAAjwbzDgp8/H0l6K/b9l4dSHr0JA9otYiIbp6vqLI0O4uy+Aexe5sKiG4iPzZVsZMyDZfoSlFzY/SIcY7UUEaB0MYN/VUXxOXYLEFcq1c6Uxhyq6vLSyUYVS1xSHYbpz14UASKRX63C8sKcK+U13SebG/cLFimmQmbHx+XUPDt0YQxNDhdfNK0Wd06olDZHiIBLfLyKDWA4u9njJAtOCwHtKea+OnHBBAF0qZg/6ooyvIVFsYHEUwjBd8pc2ARAKPRAwYZhJzcJ08lJ+/wWxl/JmkILEJROnhpspsfZfEuXghDw/sbuFWxCsl78SKTeXm4jbYyym20aF4ILSkEYE0oFHWkpAGYwQgIsjVowy6GeOr00666DuJQQEzyVztJQDEy9SySIkL/mcqBc41wElUfVehgi3Mz4g7minUcdZgyUtAiBPFQJwfYiZfhmypAjArNdB3UgIyNkuHIHoCG5/yffy+1xvAhfJqn1tMDKRKyA9qMyaAMjpL+mLB+meOEQzTISUeq7JZHN9M6r5Zx8CgmPCEQ0S54YoegsOCi7Ots2eAHAgEpxwjQkLCjUT62yBpu5TEMglBOTMF5RvG2OqPYoD6Ry8sycAHACLtjJjiYliwITSJpdQKYBny6JNfhXAkNUQ8wgCib0jBIAMN/MGmhgunJ5oNHsrAEcRoCzSTs+/GLUR6VChPIKxbkMRBY0sWJRvRE8i7+Uf7e/EU0SekxaXd+NEVMExDpO5/q+2Z7hZtP0jwLht7xh4+F4b8GO0Eah2SJiwJFuZ2FApAG9WBEBql41R69/JGOUAMxUUaQxGCmDk+nCBRDYTu7XUv7Mzp1NDmQnzKq2od9lRZjNqLtJ2GnOlAERgInuNZEvupp273xNAF6PZRmhOFRu32MhVRGVKS1CwF8uhEOXeEYuH7J0YnaCqSoxaoZOGCjvq6Fpr5X6QvSNej4nUZ2563dosbprfg0kXWL0dSLMjAOzFEzajy8fKrAjc3uec+izmKcl3Z6FHmuS8X1Rtw2LWv2tgFRxB+hI6RthtZq0Ippz0gvzyV+4L8h9ZTC9lqQBfPvq39zKrzcX+AG4M+VlPgfwdf7drnnGpU/c5tRAFOFk5NOTACPFNmd2MFY02rXbiwkobKhwm2HiyltgsKOFfKSWe2DuJw0YzkfoN8I0OI+gLwEDCkWqbHQHgXvQS76XK71w1/QkCi/Kz1GbAynoHltXaNX/2ugobmrmAZUm5RIpV99YWYHDMUma26R32oXMoSHNPgAFWQbiZSks4LdkEqhU2BASBNcU5kbq5wopldIFeVGNHQ5UdTdw/Dc7kg2lirkp0hvzwjfs09+lUITMrAkCdn5bgsYunlX2O7UdZPGHZhCovLDNz4ay4Z2kFVjWVwaIDLGwmE1awfLa85DnnmOji8PVRXOj0ot8Xhi8Yz3qrdAWpbvXcXy/rKftHaHg92frmCgvWz3dqxVIbZhlBR2dqmM1W6uBmt/lSJgBic/RQ/neTAwhLCubkiVXuV0CnEYiP+pomB57eVEOWza4tqE5df6Ubqae3kqW0b1Dh8+5Zxr4zCs4vaZdVKzgICIpK5SKpkPT1dVXYvtiFMmHr0mw2mxU2mw2hUChlQpAyARDq5RbnnyBpj+QqmiNNKLebSSc2zHPgiTWV2Da/lJSX2VlmR3iTglqiaxGzFlK38IN76nBugRNvnhjAJea/s2hyYVJdqYtyDAEvObdKsvYPLy/HN9e4KDpaNb2QHsOyMtOSrdSEIEXFVDmBWRGAQb8ZfcztaDZkJsWzHkDRqw9BfIr68PP1uxurcM8SF5HRDntSMr5eo4jL/k4u9HoSoJrSBnzIJHEfMyyWesOiTISpH+Ry39MALTsbyDE+vroSWxeWwUWFn57NanXAbndiNOZNuduUR6JxAN4YBryilUz5eQV1gyC/VGURbf5zPPXvWlJOJQ3lrRzOQiLnltTZqYOoRlWpGR9c9GCUUXRiJpqlGJjD2RT3o+XgGKevzOM89XfwtbqpFKU6sPy3Q81sNNNsbNX8cW7/babPKRMAyfsvtf7GAwZw/xVt005+zq7JZcYDS114fF01RPbPl9ZE+/DX19EMS6XhJwyhHRoLa5xKJkWSfJl7IYyD1j0tw9GWllJ8a2M1llKpKwcHv9b9AJF+tQQq3AuptpRReJjyv0fj/GUqxdtkAeucZjyywoWnNtZyork896eGczlFgqfWV1EXYcDHTI4pufGEcClOYGp4Zetb2Tviu7GGyts/2t5wi1kvU7vIYqW/gKMEfi9l8xQekvKR1seaf2NhUQBmC5zZf06USnYjJ/j1lYL8NRxA/k7WRu+wb62twsOUL6vJkqnArOzvl8lPFAIszj1r6BvyX+xo0Mx9k3/P1HsT94GdZmn6EcbZjCQflBIBEDQYoAJwjPn/i9kBaJiefb+3oRoPrKzmjFMCUZJg1/cycQ99Yk0V7lpQqvkniIeYarmBwHAggm0tZfgmDw6x9WeLGzPRC9BqtmmT1ohAktNPeXdLaTfOMY/PxCRnPsVlorQZI/L/LtnqbbTRushiF0qzkSI/ub4ajyx1YpwcjJxEqmUXAmLqW09F34Mry+m7UZLVh5uMFsaflKb8zJQIgFCWUS8zkVCzWWzKJnHtFTfbTfPK8PCqCjRQ219orbLUiodWVuLrS5x01lLOQtlaP6G1cng47EZ8jTqjDeQAso0fYv+30CNQInNTaUkTAOlWapR7aAGQIJZssTapTGa21yYUZw1U+j3J07+FvvxJA2a2D83QfQtqSvA4dQKLOAfZlCnuhwyNqvi7lejNR5Y5sWVBGRxZ9hFJQNdIqmM2mYmbyeuskt7nsplGWPlXvFCT7z4xtPz+KyJzFaOvti0sxdb5ZVrIZX6PePrRNdNR6Tuba+LZdottsaafetZ/FbwQ5F/MYJ5v0zW8ilxYrpogfkkpzY38mywnkBQBEP9/OUncIclBVlwEQBZQjvtN9LD7NpGmGJr4l2+m/Xk7nU/EcUjyFKiWGQgIbOtKLfije+vhKskd8svsDCREJXQ3NqQgfyRFAKRzSfs15JfNxHh2+aJImvj337/QSaeaGnrTpewWkSIUmEjFP6S9UtHUpvgQ7XJxG35mXSUWVnFT8kRQNGA2UJz+nnEq/ZqZA+IxBvYsoMgl1phMtnBMokE9xMGpy6sbyZvbLPaURICkd7yolKQ0cThWPCYAWcA11NbeRY2/yP36twg6B67jasd5DLqH4RnzMI/CqPYYh82JcmclqirKMb9hCRbUr+T3+m0gYQNryqz4zqZaRI/04yjDip05kk31h2vuexRTa1mJmTK/E/dy/4gYoHcb943gQttx9A31YsTjhc9P5I8EKOdb4SqtQFmZAy0NLVjctI6xKaUwGk1E/tSO56QIAPeSJgKwOpOWjFD/qeoNuun70xhi/lNB1u0xavxXMlBD79Y33EnEP4cTV/bj8OW3cLnnOE9+Whr4IHm+EFRmCsPypq3YsPBRbFr2IFYu3ITq8npdh7KmuRTdLKfVz4QiXfybqMKj60PmWGciDgf5eog6owcoZukR0jsZhHLKt/dcx5krx3Dwwpu40PUGelgwidn3NP8b2TsShzO/xomNi57F1mWPY/WiLaivaYBFOAA5SGSTJYGohjNnzsil0zYtByAVgC9fsaOP+esKvQaALKBQbDH3PUPbeYVDPKiSgte0cJIfI5EwBkb68d6hX+CdQ/+AG/0dTNpYRRONY4I1u7kq0SjLqUW8GA0Mo9wFfO+Bv8HXtj6L2gou5IRTx4wPTOKCIUajfXhmEL89OaSJbxnmVJMYUWFfInJ/C9O+/eCeWqxmZWM92zhP+RNXPsWbn/0M7554HfV0J7CZG4nwPKvlJE40buJwNMSYnB4wOz8eXf+HeGrHn2Bx40b0d/UhHIqXTE9cfqe/pj/90z/93+/0Y+J7ea6EnR7uMWo+APL95LEkriuEv4L8kpRhPhfwX2yvI/LHM5pMAu2spxEmQnf23cD/88q/xIEzb2lZk10lVTCRZYubZm59irBrJpONQUblMETKcOTKq+gZGEB9dT2JwDzekxo7d6eBy6nvoE7AQw+udpblVmnF7gSp5L6X3Hvfp9Jv9TynrnK/PxjAyx/9Lf55z9/iSufnqC1t5kFQrrH2X0E4IqCw/HYLFb0mF270HcKF6xfhYJ6BMkstx5Vcpp6kd5iYyrzcQJKmuFCRX5ZX6svXkvX/HhewlAkX9WzXOy/gb175b9HadQ2xqJnsWlISljYEIzdVibkJJ699jF/v/DGu9ZwmrMnz6dQW0TT4LYY0O2nulCaEULXUICDOYtKeopv4Mvr6i/dl+i2+EL6AFz9793/Dx0ffYni3m4gtAWhJNiKk1dSIgdEOvPrx83CP9ZGjlb0z8/hmJADShWj+pSJpoWuSE1rbR9dWYjERQjgBvdqFttP49e7/gIsdu1jLnsK+AD8lSinx/AzmYYrni+0H8cK7f4fr3Wepc2FmYJ3afM75T+6pJ9dh1HQ5OnU7J7qR7M2S+XklXX0fWubScjGkO/G4JciAkfEhvH3geew79RYGPO3sVvbCjKh5y+Pl+gitA8PjHUwZN8b1jSSB/klFupCWyOnPAKCkerxlWPnzIcxJlFJrKxlZtjOxh55a27a+q9j9xW/w8dlfUMnXHF+8lJD/JpzMphISEB8+O/c8Xv/keXT3y4bQp0kWI0lC+U1GDkoosXBDqs0MAeGW6P2OFrqHf3dTlZbyXY+jQ5R1/tAYjl/ah9/u+48Y9gzwm9Q1+YkZGAwmsv5W4urwxMEx8yhnJDPSRYQ+AH5yAZqmLPG0Avor21y0tvcwj9+OZeVwMcOPXs0XHMNHR17BnhN/Dyuq4vJamp0bKb+VWFrwm4P/EQdO7SZLN5Bmjzdvl6Qmj1AUkGSjEkosJ5tq00NACGUL7f0PUuO/siH1gJvpej9//Sg++OLXaOu7Qg5DvPhmvzeFoMhqjgUHua7JJQid+WnEeyEAwgFI5zPTlOmmm5vfJL5/Piv0PLzcxcIdtL3p1MKRIA6cfRefn/sYQ54Rym36aYRFaVhprcNvDvyPOHT2Q51GHF+/WuYN2L6inKnH7Rp3p/QBdwavEEgbOafti50MD6+884Wz+KV76Dr2ndiFD0+9gqqyBbPoYapb6GzmG6V1iVaAJLB1RgIgCC+HBGOACraRfuE5uvkuaSjTbQ7hSAgdA1fw5t4Xcb33AJz2FhJIIZH6NYvZjv4hDz459ha+OL9bv47Z01ZGrO1YUkZZ1qS4gGkg66Xv+w4i/zaKjXYddUaiT3zns5ew++Rfo9bBAn+67B3BVlaaCri5pjTbJXFaz0gA4rCZTdnB+J25/DfB3n6Ltv7lTM8kaZr0aqPecfzivf8bnYMXGPrpJOumX9+Tx1hqbcSJG+/i0+O7mfeva/JPab/ftrgC31xViREJ8FDtKxAQvdEi5vK7VzJB02ysZ3vnwM9x5OIeBAIlcRu/Tp3LSgZCPlp5hMTMvCeTIgASB+BnHfJCanGtrZH11kqZ16+c6bL0MPnFEUVk8vcPvYCjl9+jEmecBECPvqeGrugDwuEQjl3Zgzc/+RWCIf1qMZaxbuGmhS48Q6XgOPMHKFHg5hrISkd54v/ehkqsor5EryaFPE9dPog9X7zOw+MslXb0ANPz8OAiirt5hHJvMhibFAGQjSGmwEJpMl7R2s6j1vY7zOXfRAWOPjA2aBr6k5f34/X9/ye9raQ4KkGYYdDYzLUY9FzEziMv4JPj75DCi5lRmmzT9FoTYyAeXlmB1dzkAiOlE5S9HmMdRgN+n/Z+Se4hylK9Ws9QB9749Hlc7T5NOT2qi9L49rGJAlDXcGARAJKzKt4+lNx8TmhtxV67lnZbPdu5G0ew++hv0DXQT3s6T359KMv0QyRiWoxVGPd58JOd/xUutp0ihU/O0WP6juO+5YtY2PQJRrTVTOR5n8ucgCC/lOHewBgKLRU8Kzvr1QY9fVTo7sb7J39C0h2iH4p+CunJY9T2hixiEgfTjKRN+hBmIhDTDxCTB6v3eznBJDPKA5TbHlsrST31a33uNnxGWXzXiReptZ2vX8fJ9ERCI2s65A7jt5/+EFc7zyRzV1LXSFGR+xjRtnm+g/4BxjnLBQh8Zf8sFYLI5B561oEQifzU5aP40c5/haqSFp4b+uNTQusvpumonjqApHZRnlw0Snflr9Feu5kEIAkCmNKo39n/Mj46/VNUl+gbsZf0IEgE7OYKfHL6BZygGOL29iR964wXcuM/s4n+7VSWCm8xF7mAANNdzWcNxnvpK7KCREBP5u7g6ffx8t6/hCFcwX713pm3ru5UcSe3XnHz04wcgFwq8kQwUZE2s2O/ObIU33H/MlwyhnVM6rmdlXwkQYOe7fV9P8XBs++TDR+j1jaHmV+4eSxUHL118Ef44vRnuk1R9mR5iQlfp8VkG7MJDRWy3XcWUIny6I/RSrSdIb73c//o2a52ncL+k7twufMzeqDqZ4q+0xiDE1aAZFB1ZgLAXuKskaBYfrbEyGz08HtmXQWWSwylTk3s/aeuHMCuL/6ZWtvTDPDRWWs7i3FaTRVoHziFD4++isPnP55FD3e+ZTnNXneTe1rGv5L9ea40D02hT9PRZ/vSCt2q9oqSVjJAffjFWzh8aSfVxYz6TMPTL7m1iJH9jyTNwc1MALgH5HT4Mt1RHu4JMfnZGeDyNF1c188r1VFrG0Nn/zXI6d9OV81YjOXAGYKZ88b1KLU00z/gZbx14BfMOnRD49L0GJecGqL5fpTprc0UU4tdFJD5Cee4iUo/KeDZUK4f5+gP+mm1eRefnXmHQT7n6SlaoccS6drHzASAjxOZwpKnpYAlRFPMNKuZ2usbzOvnsOqnXBkY6WFc/y7sPP48aXmErL9+myPdVdQIUbQSF1qPMWjonxhRNphul9r9Qt8lUOieRU5sXyRppuL+5bp0nmedCPLL4SZVlr9JK0iTjs4+EtV5ue0MfrnrL9A/ep5VgedzD2Xj9KRhmkFBMq9kWlIEIN5RNgafzJBvXiMLGOY/yya0tnqXXj595ST+6f0/Q6Wmtc2Dk//m1LV3VrOTLGYv9hx9nXkEPmPV5pHbrkj9Y2Lf1Lls+IO7alHPLLMm7iaBdbE1yW3hpDPUQ/SGXN/sYD6GFNBhBmB0D7Tipx/8G/j94QmxcYYbdPzZaokHFSWzZDPOWDox0RBoN2glgXUcZvpdidZ2HpV9d9NPezXl/mSpXjJPPnR2F519/pYRfolUXsnclf1rTAY7HVcC+Kc3/xwdvW26DsBZQiJwT51WIq3YQoc1czHNn4tqbHiap7/Dqh+B7x1uw8fHXyd3dkQzx2Va63/7osedgJJB/6TyAdzefX58lgUUV80dZFPvZz08PduNnnP47OROnGvbRYVQlZ5d69+XnM4k0H3u8/jN3r+nf8B53Z4RF60ceIQnpHACviIqNyal09ZQ0fnc1jpdE6WGmOPxxKVDePvgT2iGFhdiQcQEX6Xb0kzZUULEKLFKaTI522cmAjNyAPIkSSJpM83c2ZSjytCXsoBPrqhgoEYFlSt6UW/R2g5Ta/sGDl/cRfiVsshCUiDK0CyT7Zb+AaZ52HvmJ9h77E30Dncke+O018m2lXyCD9FVeCMLp1j4PpEWa9ob8/zHMaaDX0eO8WFq/Reyoo+e7ej5/dh5+CV0uxnfz+Qu2UL+yXMwiaJa2OEkUDap3W0wMB2SMYneJo8iQ+81uZ/H/2a6+IoPu/iy69X8QR8OnN6JT069jb7RM9z8hVMpSAhVLGxnYpIXGGP+jhazoBdcqugf8ABzKWxtpmmQGvNCbuLq62IW6B1Ly7GRBWH0bK09l7Dzi1dx8PLr1BstIP7lBlZSKThZsSMpAiBULDtMzPTLIeAUwlbO1F7P0FWzuUo/6i0Rd9c7L+Jn7/1bJle8mEWt7fRzTuVXh62e/gHnsZ9E7NilvRPxAqn0MPW1AvcVzKXwMDmuBhejE0X+KtAm+2cHqydvYA3IEh3Dw73+Mbz6yX/Cietvo9zawI2aIxhxgg6bixareAKfmZZpRgIg06CoTWDN1FXmfxd7v4Sw3s+NuKLBTvOffmSpe6gNP9v5P7P6SoQiDwOI9Os684D58gk83ezzcL7jPfx27wsYHOmG1B5ItyVAsbypjNGV1VqkZbp9Zvt+4RwF+RdUW/EEw5+bdSz/Lsgv4eEnLh2G1xev3JPt+SWeJ2tlm7ACJCMDzEwACDgTqZnDFI8vzhFd04I0xCa9iAv47Q36am17iPx7Dr+Cs9cO5URrm1g8Pf6Kp5mRuQlv9J7G8+//e8aG+/ToVutDTGbrJUqOCCQIVUiMgATjmOnL8r1tDajX0dlHErhe6jyGdw6+CPd4F5+hX+6A1BdOsNOAUntlPMlIEsg6IwGQDqUykKYETBwFqY8s7TsSWttnt9SxFBMdHdLuMd5BOOpnlNbnePeLn/O0FDYnDkSdus9JN5KgZJzKzM/Pv4uPjr0KD2vMxVsSO2KGEde5rPgGKyrNJyEW5XAhEAE/3XzLSbyepdi4mJp/iX7Uq3Uxa/PLe/4BPcOXCAsGzWfc1XfmkQsHEBcBZl7vGQmAdCH+EWVm0tCZ+5t5dLO4QrS2q2tLNG30Ui6gnu34xYP0qX8NHUOXYTVLoIZ+m0PPcabUl/C65AN8/gDdmH+Ic9cOa74CesxNkF7KYj27kWx0OfUBZAWEG8jXFqTS0kUlpqSDl4zQeob49gy148PDr2HfuZfJJUv0oF7WqNlCUxZCIkYdJM7JjWVGAiBD0UQAc0xzC53t0GZ7n5idpHbfg0zrJemr9GztvVfx3oGX8dnFl1m/L3daWz3nlOhLTiIpSXa+/Qg9BV/FlQ798geIhnnrogo8QIRqpH9AME/ZACFMghLrGnl40JehgoRArxZg8Zd9Jz/AS/v/nFWXm4h3uT844uhvIIdcTQ7ASsIs30zfkiIA0oXmC0BWwMiJztzt9A9N5VcxpUg+9k0LyuCw6AdkL0sx/Wr3X+PotVeptWVW1lxpbVMBxiyurWHikjeO/Aj7T3ysW7xAfBgsrrqmGnfTlCa6mXykAWKtmF9pYVxDGZbqbO/fd+Itevu9gjBTNJpMeaAh1xaFmEn8dNEaZCHxT8YMmTQBkI5LWFRCI3RZoABCvATd51VZ8Aht0I2UPfVqorV9e//PcP76KQSDrMWSNwuo1wwn9UMgVtnqGND0Q3I7L076If23DlphhCvbSpOaN8+8BLXTn4TpG6uqsIXcil5NTtW2/nPYf+ojXOzczfqS8/TqOs1+JHFfjEpIM0UAmgG1RLUzI2rSBEA4gOpSo6YPmLnbNOfC28XfRGrYfW9bPRp11NqGIn5c6TqBDz5/mVrbbrJK4kugH2eR/sz178HMKEa3p5MJTXbz1HpN1wesYhTmw+TQWmpKmOQymTNH18dP2ZkgvyiNv0uT5YYFTl2VfkEmZH3to1/gTOsndI6r5YGYNApNOVa9vpQ5G1lCvMwR58iS7Tep0UvncvKX20hh+JcfM9p81NpWUl57bksN7f2lVM4lNcykxtTZ14bX9v4EXUNnybZKDXX9+k5qALm4iItnMtawgMkxfHDoFZoIz9HioV9w1+omB35nHfUzOmrXZwsmEUVkr26lyHgXQ5rrqKNItyXImniKvrX/l/jiwk5aVnp42uqrkE5vnHL6W+EqFT1Z8tGbSe9+ubDCGuVGIgHIIAWQ5AxOxqNv5gLuoMOPiB16tb7hLvrKv41dJ18g9ZZiHvophfQaY6b6kXrxkirqfOsXeG//rzE02q/Lo2QrSCTdZp60T9A8aCGrmKt4gcS+rC0z46n1VXT20cdNXPxg5eQ/eeUgXvrkzzHqHSCt09eNON3FkEIgFmYZrilvSklZnzR2iS9AlZ0P4d9M4b8soPitidb2Afr5u3RE/jBPvH0n38eL+/41lX75obVNd9FTuV/WzEIzZyQWwK8/+Uscv3hgkn9AKj3dem1CeHKV0MOO+QRXMsjGSjYxF0RAChzVstTZQ8ud2MTckHpWgL7ccRqv7v07jHpkh4ouLDHzW+GRq09RlpGXNOP1lfNIAExJKQBlrEkRAFEviA6gwhrJqA5AFrC53KxpbVfQ7q9n+/TEm9h74rdgliZNUaJn34XUl3A9FVRc/eP738Xpq0d1G7qgQw2rLz2+oUYrxCr+AdlsWiUoiiAbmRLumU21uj66x92q5V7ce+ZNutlKVt+k0EbXMczUWYzirMVYgsrSBuIqx5ck/JOeiRA8p4UcgHaH/tRPxhvg63fIut3LSC29mmhtO4Yu4rMTe2gTfy+PtLZ6zTDVfsjQkpqPeoD3P/85Tl09kGoHd7xeDolNzKyzZZELNU4rIxKzRwRG6Cx2F8u/f41p4ZI81+44j8k/SJWdfUd345V9/w41ZVLEMz+bVmiEtn9XSXWcACQ5zOQJADsUBWAJY+8lQWiSBCapYUhfY9TaPscyXqvmUYupI4UV2e3V3c/j1I29lNuq8pJ6JwUkXS+SHI+NzHnwGhOffIDu4au69v516gIeYKIWiULQc5/caZDiKbpF0sHTMWm+jkE+8ryPmV9h78lX4AsEqP9KX6F4pzmk+30kNsbDuQx1ZYs1a0CygnrSBEAGKHbGaqeZSh9x/NCHumuyopwczEQrVVhreXLo1cTZR8x9n519g/JuhyYD69V3ofcTVwr6sf/sO/j4iOQP8HJK6a+p9CBKQeHiHmN+/VEiZyabaP2dtBg9zCzGa5v1zbkvZeD2HPktrvTsYjRsvtj7vwpN4XLN1Gc77JLXUAriJM+hp0QABNh1hLH45IidPt0mNEQS7jQxxvxJsv7zqLVNfujTP11O/lNXDuGFj/4MfpZKMhv1cwaZ/smF8assX6ltATqHjuOjo2/ii3MfMT22fqHDC+kXIAlbVtRJzkKeRzrsl9shq/XJDfPEKir9WMxEyr/r8RjRqEuE6Ouf/hgX2r9gxFMN92lKqHL7UDP6WczZZfYFqK6U9HWpYVDSsxLASkBwjS2IMoYGR1gyPN0miiIpSLljmRPb6E2mp9b2YtsJ5sj7DxgZ4Thj+ae1TRd2etwvbGIZ01X3uE/j5x/8BWsftFJ7r9+JvZiK3CdpGXDRrCs6JD2JgBxGsl9W0erwBGtAVpbGOcf0dyUg5d/3nXwTh87thC/oJufo0IWw6LFmU/URYUSry8Ew5yrqKFKkgEkTAAGs+Hm4LBHmn5M0lOk1WUChquvpRPI0Pbb0bJ0Dl3HwzE4cuvA+7Fae/HlmstFzrmn3pSGmDYPuIbyw839FHyPc9GqSpn0b/TnuJycg72XN9WhCSKSrJrKif3xfAxVf+omN4UgQV9rP4/ldf6Y9w2zMZXx/ctCKRN3k5qpQ5axLFf+TMwNOHka5NQatijSRKp31HPJH8cDiMjy5ke6UTDyuV5MCHodOfYp3D/+QJ0+dXt0WdT808vLkD+GL87tx6PwHGB7r1W2+gvhPr63Emqa4f4A4eqXbPNQrLGJI8u/eVYemciuVc3qc+/FRXWg9jV/u/j8Qi4gpkWMtgMMjRMpabm9kVaMl8TGnAOCkOYBEn0aS34oyE8NnWS1gFmyAUO8+bwSPLXPhYQZqSI45PdueI2/goxNMguHt1xwj9Oy7aPvSNnmMvvxhvPHZT3GW/gFCSPVo4jBTXmLB72+rxT0LyUpzx6VDBAZ9EU2v8PjaKqwl9/hlyTodBtvWd4la/9dxtu1jWqIKJEZEEIr0z2mrYdlxOril2FImAMLGlRM2FYwLSKVYBG/TlEFe2oYfpXZYFlDP5B7S/6X2Uyzi+TIudX9AbWhLiqCY65cb6MFXjotdR7CL6dHOXDuiK0DqXHY8wQIcjywpQ5S7TrL0yJol22Tfycm/lsFHT3DvbGUYsp46o/GAm9mU32dG6BdhRiUPfv24imTnOJvrolRYOukzV+agDo2OQKm2lAOZ5dCvsYexyGVG37iZ9tGwRoWFC5sKZkKgIvIPWwUzi66l0uYxKoaWUkEkcQV6NNlK/cOdeO2Tf9Ai/YwxihV5EJiix9yy2wfNvCXz8fG5XzCopBYN1YtQW66PGCU7YGFNKes3xnU/p7p86B9jMBbXTnJMTMXFy7YRxI/yGilScld9KR5gYo9NVBjrWQYuQjfxz89+yBDf96kQvY7KAkoOE476ML9qG6rKKyjGpc61pUwARPtfZwthUxNTIDkcrJ4zConek9RLUmN9cpNPEshnp7Kvkll9NtJU8x3m9NNz8eR5I9Ta7j/9Dgt5vk/20ketbeqUcPK45/J78cEqMVTTSehNuMrs+P4T/ws5g/SDauQ8lf0gROAP7y1B3flhHLw0gpFAhGXIoxA5VrtgAvhyrRwQdhJyp82spSH7Lqv4zKfsr+fZHKHY095/FW98+itc7N7JQypbRTwnJprmn2BkiOnxn0JdVTOJZRYIgABf6sYvqLVh/cpadNIB45PLozjTMYZ+T4hUSCh6fC0lhqCO+oItjBS7Z3GFVsRTFlbPFqEN9FrHJfxi15+zWycJjn61AvQcZyH1ZbU4qAjsweFzB7Bq4Xu4e+UTsFrSh2sCccuoGPwWxYEdi8txrHMMB6+O4lKvj4cIo01l73CTML0m6pl/YiMdxO7m3lnDADHhFBJ96AVPj8+DH7/9b9ExcI6JbyVITK+es9NPiDhfW7oIVY5GHsCpK+VS5gC0afGYCDEdMqIeLKlhTXXGXD/LVN1jVNAMesMYC0aYv9+IOpoL7FxsG92HE+W79IbvRWpt//Oev4I/IOGQ3DlTySHZWYsieoqBClQnet1X8dLuHzPZxwrMb1hOBJzddpkKMKK8q2LY7o7FLtxFznCM9RgGx8MaRyDlyCSqT2pAWGXv0LNQKhTr3YY9/XiTmaEutp1EgKHSySbS1Hscs+1PPABLaaWsKm3U8gCGWNwm1TarFRUFSZCyf2CcTkF0PpJKPdKqyyxopCggWl5R0OiZfnmqiYnWds+R13D82jtwaEo//TfJVM+dC98JMkjV4Svd+xkG+2N87/F/g4aq1LXM08FKOEXJ0iuvCkYS1lfERUkhDnoWfZlqDIHwOM5cPYR3D/6IeqwQkZ/yRgaIzFTP1uu7GFn+purVKC0pI+MipvTUCQBnnXoTOIWCYfi8zIh4WxOkFxk/08jvC3rwyfF38NHJn1FrK0E+CvlvW4q0P0p9AQtdqN858v9hL2E95BlIu8/pOpBDQ/ZOppFfxnDhxilGQ77IQKgbcTffgts/NNtSAbi08S7qaspnxf4LHGZJAJjwgX7jgWCAg0hd8SAPTqeJ1vaLcx/i4OkP0TfSQ/Eiv7KzpDO3vLuXiGEz1OHFvf81Dp3+YFaKpnybk/j5f3rsA7xz4iVN419wgv8EQKOxYRY62UTuqWZWFgDpZlYEQG6UE1eUDj5m2BVZJFstRK1t90ArXtnzc2pt3+cCFpbWNltw0vM5FrMd7lHg05PvMjHGR3p2nfW+gjy4frv3ebx99C8wz9lMw0P29q5ukyW+iTrUwdj/WscSir8VnMXsDuK0CEA4RKXNEDNLZLGNjnvwn17/n9A1fIGsor4yaRanUXCPctqbceTqS9h79H1yXa0FN/7EgN/c9484enkPzQwVcbk/8UMB/RWiZWRugpXz74HdZk8ryCotAiBigNfrJfuRuvkhFXgnqPTgaB9e2vM3ON/2+YTWdlY6zFQera6dgICkEjMzocqxqx/ipV0/odY+UW+wMEAUCgc1X5G9x95jHb8rVFIzrr3g5P44rCX812SwYv38B1kKvDQtDnzWBCCx7JougCYUiaHOVBN2Z8jThw+PvIqdR/8vhOg4IgEsqmUXAhaaBt3jrfj09Ct469MXWIBUkojkc4uz9xEmzLzUfpys/8/R2ndGc1jLx7x+yUMyTNOoAwurN5ELLk1LL5M2Fon8P+bxZJQL8HiH8cmxt/D6/r/B+DgdQpj1tFCpd/KLnH9XCidmMZbTbObBrz7+M6YT28lyY+78G+iXI5LsxCFcbD2B91i++/CV1zVkkbp5hdqknoPd6kJL3SItACheAWj2s9GFAIwMj0H0AZlo/qCXwSkvE/n/EV0D18jy6OObnomxzok+yTbTTYiRXY34f9/8Hew7/i4jL/NTHIjQTn6p/QTePfCf8eJnf8e8AfN1dWbKxXqHo1763TRgzaIt2uPT5bzTIgAJ23sg4GfmlCCpq74a1QArL76y9+/xxr6foXvwiraACX1ALoCvnjkBARIBk8kCY6gR//De9/D+wZcQnoUXWqbhefbGAby85x/x9hd/h2bn/Ew/Liv9R2Ju5rloxLL6bYyVSD7//50Gp4sWTc4Ef3AEpRErA0fS9xmXwfYOdeCXO/8Kp68cptzZx371LQ1+J4Co75OHgIm16KLhOryx/58wONqDZx/5PurKFyTfQYaulLiiXYd/jd2fv4HLDG8uszFVVsH7ifHo47xK7RWoraxlmb7GtJR/CdDrQgBEHh8eGGVaonJYnekRgBC9m05ePoQPDr6Go1ffRIBumgZ6pBW20iYB7iL7y3Wn3x4GPVeZOvtVjHhG8MCmR7Ft9aPUsmdTzhbOM47h4uTz3oF/xpGLH7OK7xkmOQkyrqEYokMNtHwNYlnd3VjasopwN9Hyn77YrQsB4D5gxZ0Q/H4fnE7xS56dZHGt6xwTURxmWO9OfHbhn1mIpJFsDjeSPEC1/IQA18ZiqmT0YAdLkP8tBjyt6O7vweZV92JBw1LK3PpmfJoaCAaNA7nEwLDjFw/hnWP/DqGAneOq4qukEF19ppymLzLOeodrsLR+U1qa/8md60IApEORR3w+L7x8SXBCsk3Sd/e7uzEw3I+dR17A3lM/JMvPCCcmZQDrEKpWGBAQEU1s60ev/BYH+PpOz3+PhzbhiPdxAAAWBElEQVQ9g/n1S1HurKK7dqnuE5FEIW76hoyOjzGhx268f+THOENzX62jSYtAFa6gWHaQWDMaKpowr2YJK2c30Q2fNe50aIYzZ87oAiMxB5pMJtTU1KC2dvrabMKWhcNh6g38uNF9EW8f+hE+PfNLypPQvPtEttRlUDoASHWROgREVvX425hVyIDta36Ar2/9IyxuXIcSOq2YTWaYzRTpJlj2VHuPMvYkxP0T5P4ZGh1iFp/f4oMjP9fYfauxirqi5A+fVJ+dy+u9gTY8tOH7uG/VM5jnWs0YnNQj/6Yav24cgFgEQjQFjnvH4Aw6aKucmuIH6D14rfMEDl/4kH7ln6C15zKRnSmfjNTSToiNCvmnWqrC+U4kNldJC/NFGBiw9RETi3yGefXzsWXZw9i64hEsmUdicIf9MdMse0facPrafm3vnGY4ry/go2NPiPHwfN4sicpMz8yH38OE6aKqrcy9sYi+DenL/ok56UYApENZ+CCVdkODwzDYezHuc8M77qV81oduVljtHejC0Mgohj29cHu7WK6rl9TcTznRyntFzpeXasUBAVlLpomjF16AYl5bz0UMud344swhVLnqUF9dyzRWDaxn34jaigY4yyq0Q0MIg1w/7hth5qlxsvdu9Ax0slJPh3b/4Egvhsd7Mcq94w0M8RkmTSFWvHoi0f5HsbJlC6qZn9HMxJ/hiD6nv+wzXQmAJFWIMq1TR08rPjn/KiMFfdTi++k3Poxh1uZzj9+g+ygXjUe82VitaWdNKoWXrEORNjL6jCGQV5j7YsBzgwfBERh7mMW2pB7lpS1a7QaXo4YcAbPa0oRsZY37cDSoIX8o7IeXeR/cRPgRbydGvdeYiUoK1JRpGYvEH754ET++JUS09od6sHXRv0KNi6a/2/JuprtxdCUAcoDLgCVPQOfQeZy8upN53ijX03PXTE2xidTLzsw9CQeidAev7i8cCIgZ12Jy8MXwbdkjzCjV7buC9tgREod4rUnZ2/ISZlByA0rGIPlrMVF3YHTxNQ8WC7/QuMXCmftsR6rp1agPa6nbgCWM+y+1MfMvA4H0bPoSAI5MFEBORwW+tuEPMDg8Qqp/hYvIxGVcN9UUBAQCcgCIJ6G8WNE+SaBwY82xTSTik8NSjq3LH0WFo5bCjoWWf/rF6AiH2Rnsp1kyqRpope11UeUWNNcsojmmRDeb5TSPVT8VPQTm2AnCk9RgiDDbTz3WtzyspWaX2AY9kV+2jO4EQDoVpYVo9u9d+wgaK1cjEO6Ur1VTEFAQSBICogdxldZh+YKVrMTVMOFQJVyQvi0jBECGKJRqQdVGNFUtho11JWIZThqiL1hUbwoCuYWAL9yLRuca3Lv0aYrQIqnrj/wyw8wRAMp5dqMTKxesw6K6r1Gb25FbiKqnKwgUCATE+lFfsYj+EstRU5ZZpXnGCIDAWpKGLq5fj+VMXWymxUa0mqopCCgITA+BQKQPq5ofwZoFdxOJMoqimeMAZIqSqbTC2oQlJAIrmr9Gp5+x6WeuflUQmOMQEJ//mvJ5WNl0F+ZTf6an199UoM0seaEeQLy6WuqW4t5l32Ft+OGpxqC+UxBQEJiAQDDcjXtWPot5dfPp9ESbf4aNHxkmAMIFROnn76Q1YCE2L/uWJgakm8ZI7RYFgaKDgCYex1BeDqyo287aiAu1fIaZnmfGCYBMQHQBNSzCcP+y7zGPuXhy0fVLqQMyvbaq/wKCgKT6DrFe4ZMb/woNldkLbMoOAaAuwEIf76bKJdix7lmUWKpYDz7fU0oX0O5RQy1oCEiIs8RBrJy/DSuattHl15Vx2T8BsKwQAPEJEC5A0kRtW/oYK5ouZlwAnYUyWEsgMUH1V0EgryFATjiKcVSWNWLHmmfgtLPctqabzw6LnBUCEF8AljOiQ0O1bTHWLtyGuopljPpSXEBeb041uIxDIMIS7KX2aixr2oQ1TQ8w4MkycTBmWPs3MbMsEgB5osSHh3H30iewsvFeWgWGFBeQ8S2mHpDPEPCGepnj7x7cv+4b9JbN/kizTADiE7QaS7Fl1X24i0rB4fGOTFs6sg9V9UQFgSQg4A+6saRpM4t8bGIew8VJ3KH/JTkhAOIR2OBcjrXzd2DF/PXM7yYJDrMj8+gPQtWjgkDqEBAcMFpHsWXhU1jTvJ3bPyeomFlPwDuBRXwDzCjB4tp12L78OXo79t3pUvW9gkARQiAGX6idgT4/wOqWu5gVqTZrWv/bgZkbsiNWgVhQS3Kwquk+bFz4LBWERuoD6B+gmoJAEUNAs3yR2W2u2YS7l3+DyvAW2v+DWpKUXEw7RwQgPlWpJ1deWokntnwfVc55WqYnZRrMxTZQz8wKBCa8/UrtlXh44++gztUCY8xK4Td34m9OCYBMXNIcVZY04oGN32D2k2ZSQ09W1kI9REEg2xCQXP4OuxPrl23B+sbHtHRfYhXLZcspAZCJCxGwGG3Y2Pw41i+hPOQsRzCk/ANyuSnUs/WHgDjCGc09WNywHo+s+mOtUlIuT/7EDHNOAGQgkiTSYXVh+7KnsWnhkzDbBhBhARHVFASKAwIGjDEhztK6p7Fj1dOodNTnTOa/HZ55QQBkUBL3XF++GFuXPopNC/6IuU+ZRzB3otHtcFKfFQRmDQFvoBVLG7fj7hWPMUHOOu71HHj83GH0uqcFv8Nzkvo6GApgYc0qhFY8xXTinWjrO0cdARMKzpE88EkBSV1UQBCgwZuBPpVly3E//fxXLdgMU8ymS1lvvYCQNxxAfELMIUTOf37tcnxzy5/AZmVdd3GJ1rSnek1Z9aMgkD0IyNn16JbfZYqvu+Ew1iDMXP/51PKMAAjXL05CEjq8GM/c919qYcShKOuFq6YgUEAQkIKlEtizbfWDWNf0CFw2cfbRt6iHHuDIOwIgkxLtaInFhTUND+Hr236P5sF5GGd5ZL2LIugBQNWHgsBkCAjDKll9LRYT1izegkdWfl9zeBN9Vj5o/SePVd7nlQ7g5uBYFYVeymXWamxe+ChCwTAOXoihd/g0kyUsyEtA3hy7ejeXIRCKBJkB20cX34fx4KrnUOuUWohSKCealwdYnhIA2UJCL2OotDVjy9JHyD5FsP+8Dx6Wi7aYWWtQNQWBPIOA2PoNpmGsbXkG21c+RYX2Os3NNz7M7MT3pwqSPCYA8amEon7UOZZgy5IIQjEf9p16g5pVqZsmv+cnUFNdBHV9EUBAFNWGAJY03o/7Vj5Npd89rJItUa753fJSB3AryAwsMe5jltTF2LboKWxd+YCG/KrIyK1QUp9yCAEN+YGqinI8seFfshDOloJAfoFYARCA+MKKz3StswkPrfg+Vi9eTyIQYV15v+IBcrjv1aMl43WE0n0A5S4bvnffX6OlZjmikfxx9JlpjQqGAEjqRKPBhGomT3xi3Z/inlVPMMW4Fd5AZ14qV2YCvPq9sCEgwmeYCr+YcYBZfTbgD7f/pVYI12zMbXRfqlDNex3AzQkZtHwBkjegqXIZ7lvxHZRYK3Hk8m70DJ+idWD+zUvVOwWBDENAsliJtl8UfiLzL2lYp538hRbOXkAEQFaURID/CeVtYd00q9VGRyEL9p8D3MwtaGGuQeU2nOGdP+e7F2/VILlPi2bqE21/oSj8plq6AiMAN6cQCHtRW7JIy6pisZqx6/DLLD4aoPVQtLHKOnATUuqdfhCIR6dZLWVMbb8ZD65+TjP1FYK2/04wKFgCINxAiDnVnZY6bFnwOCmyHW/u/xmCQS81mxJAdKcpq+8VBGYHgSgjVg3UQ21f/zXct+Q5VJc2TbLzz67PXN9VQErAqUElIoHDWoG1dBv+wwf/B9Syrpo33Kb5Ckx9h/pWQSBFCPDg9wbb4Chx4lv3/jF2LP19VJTUk9mMcwQp9pZXlxvOnDlT+LOYsAMEoh6091/DZ+ffwLHW38AYrYXVzIhC1RQEZgkB8e4bD3ewkM1DWumuZY2btMAenjv8X8x9hc1qFrAIMHlFhQ8A7KZyrGjeSKS3wVlSjZNt72HU4+Znl9ILTAaXej8zBHi6Sw4/g6mHDmjPUdf0GJbPW89S9xVaVF+8g8JGfplDkRCA+HJEmVacYho9sTbD5aiEzW7EqauHMTwWzy4k8ptqCgIzQSBuyovB6ajAovodeGTdc1hct0bLVSEhvcXUikQEuH1JmG3YYIE3NIoTnTvx8dF3aSbsIpeQkHYKn3LfPmP1WS8IyB4x0K+kktl7N2sJPCWHnwSjFWMrUgIQXypB+FA4gB7PBew+8RLOtL3DPOwNMLNMuWoKArdDQJR6/nA7FtTdg0c2fRer6u6HjZGnkrS2WFsRE4A4JZckIqGYHx7vCHUCe/HphZ+io7eLSRqkEEnBG0GKdV9mfV6+QD9Az75HVv43uHfFkyxYU4MSs3OCZ0xwjlkfVsYfWFQ6gFuhFafawgVYDCWoK3di6+JHUVlah3Odn+OLy7+CIVLFtE1C4RUhuBV2c+STpujzs05fPxN23o+NLY/Tu+9u7pVmluqW8nXFyfZPXt0iJgA3pynmGsk4XFHSgA0LK1BbVY8yczUudB9E38g1KncMFAvsN29Q74oeAnFlnh8VZU3YWvdtrF26gaa+e1BqqWS1al/Rzz8xwTlBAGSyIseF6TloYKLxec41qNu6EFVXqnDq2kF0Dl7loo/yKrPiBhI7o1j/8tSPMjOvyPa15Suwav423Lf6Sbis9ZobeYBVqYpZ5r99WYtYB3D7VL/62WQ0o23kND6/+D6OXPoAkTDNhEWs8PkqBObmN+HoGNYvfgA7Vj+DFbXbNXv/3IQEt3txeALObvlEQRiOBhlEFEK3+yr2Xf41Dp9/h3qBCkYZOufUSTA7CBbIXdThhaM+BML9WNi8Bt9Y+99hYd1a1ucrYwRpYcXv6w3xOU0ABJhCBIxGI9OO+bVqRF1DN3Doyuu42LkHrF7OOIMWRQj03nVZ7C9AX5BQzI35NRvw4JrfQ3PlSpagm6chP6UBaoci2h7I4pDy6lFzRgdwJ6iLlSAcCcNssKGpYimqXQ00EVZjWf1dON95EFd69nKXOGA1VTIjEa0FSkS4Eyjz5nvx5AtFxhGIDPOk34yVTfdhacNGLG5cwVTzNVqgWELDP9drTcx5AiC7VpQ+YimQoo0WOLQyTk21LWionoda1nXrHblCa8FlBEI+GGJWcgwKbHmD7ZMGIogfjQVhMhrQVL0C9Vy7FQvWYGXLZtSVLqJTWEgj9pNumfNv57wIMPUOkMIkJpiYbcgb8uBy/2c4enE/ugZaMeYf5OniYygozw7FDUwNvmx/S15eCLiJpbgc1nLUuOZh08q7aNd/DE5bDc28jBERV97ideibNcQVAZgBdLK14idLBK3DJ+lAtJOmw11MPGKmslBVLp4BfJn/WTPrhcnud2Np0w7cs/IbWNP0IOz04jMaTXNavk8G+IoAJAMlXiOHfZApxwJBH7rpPHS8/QMcu/IiRsdAT8NShhxXaBsuye7UZelAgMo74cKCkX5YSIO3LP02Nsx7jMlil6PMXs61sBPxlXdnMiBWBCAZKE1cI5tKLAYBKpiGx3vhHnWjy3MRl7s/x6WuD+D2AKXWKpoQlXtxCmBN7lKe9JLz0R/u06pZLKhdR6/OxzC/cj2qyqs1F+8Si5POPHPDhTc5oM18lSIAM8PoK1cIIRAnIiOVTe5gF7oGW9E90IUemhBbh46gk69YxMpryvmykXtQp9FXgJjEFxKdJ2W2wzFSVoyjrnINmivWoN61GE01C9FSv4gVo5gOPmbSZHzR7IvINtc1+0mA9stLFAH4EhSzeyPFSgzMPWAwRjDs7cDZjgO41n4dg0xCMurvwziVhuJeapRrtGtFE6W0UVNDW/QtkmpLlHassmsywWGvRnlJoxbHsbB5IeX8tWimC28JRa5QOIgIK0apNnsIKAIwe9h95U45eeS0jxnCuD50Ald7TuBK+zl09rdr30UYi6AllpBdrtEARQjiQNSwXoOJEFTR5ktm5wpnJZbOW0kz3jYsrFrPlG9M7UbAJWz4X1kA9UXKEFAEIGWQJXdDhLnJhIUVD8NhbxeuDR7H+bYDuNJ5BD7/KPUE1Zp4IBueVCO5TovtKmHxJ2z3otCzUaPXUreeSL8Z65oeQU3ZfMLJpjlgxTmtOQqnDK67IgAZA27CT0CSSwbhD43DH/DDS+TvGbuK6wPHcZXKw46Bq2RlGYfIvW1lUlOLqZQbno5GRbjXxRYvXFAoMsCaDnSw5BybqmtYVmsHC2xsQINrGST9ltVqpT3fqSG/BggSCvlPNf0hoAiA/jD9So8J0UAsCFLleDzkxsj4IEbHRzDmHcN4wI1Bbys6h8+hY3A/hjwRuiGRKFB3aDKIeVFOQVE6ijIx/ylDwm9C5Phw1M0XKzkTf10OoL6CiF6xjEk3FqGGlZ1KS0rhLHWhnElcS20MwmKCFu1+3qOQ/itbSfcvFAHQHaQzdyjsrPaiFSHGMuc+ehsOeXvQ5+7AwHA/RjwkDMEBjAeHMBYY0hSJ4/5++IKDRH/5j04uBislB+Yv4H9iFxOvRO19xsWJuKJOVHWxCa27/BUX3BhGNbHHYnHwBK9Fqb0GrpJauOz1TLRRwyy7Zagsr0B1RR2qnQ2otDdzzGZm3xFtv/ShkH7m3aPvFYoA6AvPWfUmiMs8xtQJ0LxokpJnfhKENgyN95Ab6MXgSD+GR4dJGDxMXOLl7+IEM645wwjiiIghGW7kJaen2MLjeoU4eYgPajLnIO+FpZ78XWLok1ntxBk88d2E8jJuAjWTRbdrCjspiS2+D1aThNfaeKo7UO5yobq8FnUVTahzzUeVg/K80U5k55M1uV/89os/5VYCqvn6VxGAfF2ZiXF9icI82eWkHfZ3YyTQjV73DQx6eqhXGKc4MURxwq2JFP5ggEhGLNMQPIHA0tmtiD3R/Zff3iQFN98JgYh/kn+FSDG/IuMjykqdPM1dGnI77E6y9lX0v5+PurIFmu+9RFYm7ow/dfKzE09Wf/MBAooA5MMqpDAGOTW1l8YyRzWWW6wN8Vf8dPVHxuiD0E8OYYx1Eke16yUgJsBcd1LXXtxoBaGDIT+j4wI8xc2wWUX2FgS3a6604k5rs5TQ3u5i+iwnbFRO2sx0eZZTX1NSSh4FormIHjR9Sqi0xhloTk+TiUgKk1OXZh0CKq416yBP74GiDNQQkKLCzdOZffKDfJaz1hkrRwXTWovyjTaI+PdEbvFBEPZbEF2amODkvdwngTNywsdzHohOwUglJJFalI/y4u8J3YU8eaIL3jnBZfCLiXfStWoFAgFFAApkoW4OczLC3fyW2DepGWGlNj1uRYiTifiPclpPuoyInCAak0UE6uTY4oRBe9rEZyEWQjRUKx4IKAJQPGs5aSbx01iQO47Mt/w06cP0b2+hKdNfqn4tUAioKJUCXTg1bAUBPSCgCIAeUFR9KAgUKAQUASjQhVPDVhDQAwKKAOgBRdWHgkCBQkARgAJdODVsBQE9IKAIgB5QVH0oCBQoBBQBKNCFU8NWENADAv8/ntz2JfaoXvEAAAAASUVORK5CYII=" + } + ], + "samplers": [ + { + "magFilter": 9729, + "minFilter": 9986, + "wrapS": 10497, + "wrapT": 10497 + } + ], + "bufferViews": [ + { + "buffer": 0, + "byteOffset": 768, + "byteLength": 72, + "target": 34963 + }, + { + "buffer": 0, + "byteOffset": 0, + "byteLength": 576, + "byteStride": 12, + "target": 34962 + }, + { + "buffer": 0, + "byteOffset": 576, + "byteLength": 192, + "byteStride": 8, + "target": 34962 + } + ], + "buffers": [ + { + "byteLength": 840, + "uri": "data:application/octet-stream;base64,AAAAAAAAAAAAAIA/AAAAAAAAAAAAAIA/AAAAAAAAAAAAAIA/AAAAAAAAAAAAAIA/AACAPwAAAAAAAAAAAACAPwAAAAAAAAAAAACAPwAAAAAAAAAAAACAPwAAAAAAAAAAAAAAAAAAgD8AAAAAAAAAAAAAgD8AAAAAAAAAAAAAgD8AAAAAAAAAAAAAgD8AAAAAAAAAAAAAgL8AAAAAAAAAAAAAgL8AAAAAAAAAAAAAgL8AAAAAAAAAAAAAgL8AAAAAAACAvwAAAAAAAAAAAACAvwAAAAAAAAAAAACAvwAAAAAAAAAAAACAvwAAAAAAAAAAAAAAAAAAAAAAAIC/AAAAAAAAAAAAAIC/AAAAAAAAAAAAAIC/AAAAAAAAAAAAAIC/AAAAvwAAAL8AAAA/AAAAPwAAAL8AAAA/AAAAvwAAAD8AAAA/AAAAPwAAAD8AAAA/AAAAPwAAAD8AAAA/AAAAPwAAAL8AAAA/AAAAPwAAAD8AAAC/AAAAPwAAAL8AAAC/AAAAvwAAAD8AAAA/AAAAPwAAAD8AAAA/AAAAvwAAAD8AAAC/AAAAPwAAAD8AAAC/AAAAPwAAAL8AAAA/AAAAvwAAAL8AAAA/AAAAPwAAAL8AAAC/AAAAvwAAAL8AAAC/AAAAvwAAAL8AAAA/AAAAvwAAAD8AAAA/AAAAvwAAAL8AAAC/AAAAvwAAAD8AAAC/AAAAvwAAAL8AAAC/AAAAvwAAAD8AAAC/AAAAPwAAAL8AAAC/AAAAPwAAAD8AAAC/AADAQAAAAAAAAKBAAAAAAAAAwED+/38/AACgQP7/fz8AAIBAAAAAAAAAoEAAAAAAAACAQAAAgD8AAKBAAACAPwAAAEAAAAAAAACAPwAAAAAAAABAAACAPwAAgD8AAIA/AABAQAAAAAAAAIBAAAAAAAAAQEAAAIA/AACAQAAAgD8AAEBAAAAAAAAAAEAAAAAAAABAQAAAgD8AAABAAACAPwAAAAAAAAAAAAAAAP7/fz8AAIA/AAAAAAAAgD/+/38/AAABAAIAAwACAAEABAAFAAYABwAGAAUACAAJAAoACwAKAAkADAANAA4ADwAOAA0AEAARABIAEwASABEAFAAVABYAFwAWABUA" + } + ] +} diff --git a/gltfs/TextureTest.gltf b/gltfs/TextureTest.gltf new file mode 100644 index 0000000..891b3c5 --- /dev/null +++ b/gltfs/TextureTest.gltf @@ -0,0 +1,610 @@ +{ + "accessors" : [ + { + "bufferView" : 0, + "componentType" : 5121, + "count" : 6, + "max" : [ + 3 + ], + "min" : [ + 0 + ], + "type" : "SCALAR" + }, + { + "bufferView" : 1, + "componentType" : 5126, + "count" : 4, + "max" : [ + 1.2000000476837158, + 1.2000001668930054, + -5.205485820169997e-08 + ], + "min" : [ + 0.19999980926513672, + 0.20000004768371582, + -1.5933926533762133e-07 + ], + "type" : "VEC3" + }, + { + "bufferView" : 2, + "componentType" : 5126, + "count" : 4, + "max" : [ + -2.1316282072803006e-14, + 1.0728441424134871e-07, + 1.0 + ], + "min" : [ + -2.1316282072803006e-14, + 1.0728441424134871e-07, + 1.0 + ], + "type" : "VEC3" + }, + { + "name": "TopRight_TEXCOORD_0", + "bufferView" : 3, + "componentType" : 5126, + "count" : 4, + "max" : [ + 1.0, + 0.3999999761581421 + ], + "min" : [ + 0.6000000238418579, + 0.0 + ], + "type" : "VEC2" + }, + { + "bufferView" : 4, + "componentType" : 5121, + "count" : 6, + "max" : [ + 3 + ], + "min" : [ + 0 + ], + "type" : "SCALAR" + }, + { + "bufferView" : 5, + "componentType" : 5126, + "count" : 4, + "max" : [ + -0.20000006258487701, + 1.2000000476837158, + 1.2601539367551595e-07 + ], + "min" : [ + -1.2000001668930054, + 0.19999974966049194, + -3.3740951721483725e-07 + ], + "type" : "VEC3" + }, + { + "bufferView" : 6, + "componentType" : 5126, + "count" : 4, + "max" : [ + 1.7807019503379706e-07, + 2.853547016457014e-07, + 1.0 + ], + "min" : [ + 1.7807019503379706e-07, + 2.853547016457014e-07, + 1.0 + ], + "type" : "VEC3" + }, + { + "name": "TopLeft_TEXCOORD_0", + "bufferView" : 7, + "componentType" : 5126, + "count" : 4, + "max" : [ + 0.3999999463558197, + 0.3999999761581421 + ], + "min" : [ + 7.915305388905836e-08, + 0.0 + ], + "type" : "VEC2" + }, + { + "bufferView" : 8, + "componentType" : 5121, + "count" : 6, + "max" : [ + 3 + ], + "min" : [ + 0 + ], + "type" : "SCALAR" + }, + { + "bufferView" : 9, + "componentType" : 5126, + "count" : 4, + "max" : [ + 1.2000001668930054, + -0.1999996304512024, + 5.255118367131217e-07 + ], + "min" : [ + 0.2000000923871994, + -1.2000000476837158, + 6.20869826661874e-08 + ], + "type" : "VEC3" + }, + { + "bufferView" : 10, + "componentType" : 5126, + "count" : 4, + "max" : [ + -1.7807025187721592e-07, + 2.853545879588637e-07, + 1.0 + ], + "min" : [ + -1.7807025187721592e-07, + 2.853545879588637e-07, + 1.0 + ], + "type" : "VEC3" + }, + { + "name": "BottomRight_TEXCOORD_0", + "bufferView" : 11, + "componentType" : 5126, + "count" : 4, + "max" : [ + 0.9999998807907104, + 1.0 + ], + "min" : [ + 0.6000000834465027, + 0.599999874830246 + ], + "type" : "VEC2" + }, + { + "bufferView" : 12, + "componentType" : 5121, + "count" : 6, + "max" : [ + 3 + ], + "min" : [ + 0 + ], + "type" : "SCALAR" + }, + { + "bufferView" : 13, + "componentType" : 5126, + "count" : 4, + "max" : [ + 1.0000001192092896, + 1.0000001192092896, + -0.052591003477573395 + ], + "min" : [ + -1.0000001192092896, + -1.0000001192092896, + -0.05259115248918533 + ], + "type" : "VEC3" + }, + { + "bufferView" : 14, + "componentType" : 5126, + "count" : 4, + "max" : [ + 1.0658142730467397e-14, + -7.450581307466564e-08, + -1.0 + ], + "min" : [ + 1.0658142730467397e-14, + -7.450581307466564e-08, + -1.0 + ], + "type" : "VEC3" + }, + { + "bufferView" : 15, + "componentType" : 5121, + "count" : 6, + "max" : [ + 3 + ], + "min" : [ + 0 + ], + "type" : "SCALAR" + }, + { + "bufferView" : 16, + "componentType" : 5126, + "count" : 4, + "max" : [ + -0.19999969005584717, + -0.2000000774860382, + 5.255118367131217e-07 + ], + "min" : [ + -1.2000000476837158, + -1.2000001668930054, + 6.208701108789683e-08 + ], + "type" : "VEC3" + }, + { + "bufferView" : 17, + "componentType" : 5126, + "count" : 4, + "max" : [ + -8.526512829121202e-14, + 4.6342486825778906e-07, + 1.0 + ], + "min" : [ + -8.526512829121202e-14, + 4.6342486825778906e-07, + 1.0 + ], + "type" : "VEC3" + }, + { + "name": "BottomLeft_TEXCOORD_0", + "bufferView" : 18, + "componentType" : 5126, + "count" : 4, + "max" : [ + 0.40000009536743164, + 0.9999999208469248 + ], + "min" : [ + 0.0, + 0.6000000536441803 + ], + "type" : "VEC2" + } + ], + "asset" : { + "copyright" : "Copyright 2017-2018 Analytical Graphics, Inc., CC-BY 4.0 https://creativecommons.org/licenses/by/4.0/ - Mesh and textures by Ed Mackey.", + "generator" : "Khronos Blender glTF 2.0 exporter, plus hand-edits", + "version" : "2.0" + }, + "bufferViews" : [ + { + "buffer" : 0, + "byteLength" : 6, + "byteOffset" : 0, + "target" : 34963 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 8, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 56, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 32, + "byteOffset" : 104, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 6, + "byteOffset" : 136, + "target" : 34963 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 144, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 192, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 32, + "byteOffset" : 240, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 6, + "byteOffset" : 272, + "target" : 34963 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 280, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 328, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 32, + "byteOffset" : 376, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 6, + "byteOffset" : 408, + "target" : 34963 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 416, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 464, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 6, + "byteOffset" : 512, + "target" : 34963 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 520, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 48, + "byteOffset" : 568, + "target" : 34962 + }, + { + "buffer" : 0, + "byteLength" : 32, + "byteOffset" : 616, + "target" : 34962 + } + ], + "buffers" : [ + { + "byteLength" : 648, + "uri" : "data:application/octet-stream;base64,AAECAwEAAACamZk/2sxMPuySX7PAzEw+mZmZP9gWK7TKzEw+0MxMPuySX7OYmZk/m5mZP9gWK7QAAMCoO2TmMwAAgD8AAMCoO2TmMwAAgD8AAMCoO2TmMwAAgD8AAMCoO2TmMwAAgD8AAIA/zMzMPpqZGT8AAAAAmpkZP8zMzD4AAIA/AAAAAAABAgMBAAAA0cxMvsnMTD7okl+zm5mZv5iZmT/aFiu0mZmZv7zMTD7ZTgc03sxMvpqZmT82JbW0kDM/NNoymTQAAIA/kDM/NNoymTQAAIA/kDM/NNoymTQAAIA/kDM/NNoymTQAAIA/y8zMPszMzD7a+qkzAAAAANr6qTPMzMw+y8zMPgAAAAAAAQIDAQAAAJuZmT+YmZm/5hANNdPMTD7HzEy+rlSFM+bMTD6amZm/Boi6NJmZmT+0zEy+8u6ANJQzP7TWMpk0AACAP5QzP7TWMpk0AACAP5QzP7TWMpk0AACAP5QzP7TWMpk0AACAP/3/fz8AAIA/m5kZP5iZGT+cmRk/AACAP/7/fz+YmRk/AAECAAMBAAABAIC//f9/P9JpV70BAIA//f9/v6ppV739/3+/AQCAv6ppV739/38/AQCAP9JpV70CAEAoAQCgswAAgL8CAEAoAQCgswAAgL8CAEAoAQCgswAAgL8CAEAoAQCgswAAgL8AAQIDAQAAALjMTL6ZmZm/5hANNZqZmb/izEy+uFSFM5iZmb+bmZm/5hANNcjMTL7SzEy+slSFMwAAwKmhzPg0AACAPwAAwKmhzPg0AACAPwAAwKmhzPg0AACAPwAAwKmhzPg0AACAP9DMzD7//38/AACAM5qZGT8AAAAA//9/P83MzD6amRk/" + } + ], + "images" : [ + { + "uri" : "data:image/png;base64,iVBORw0KGgoAAAANSUhEUgAAAgAAAAIACAYAAAD0eNT6AAAACXBIWXMAAAsTAAALEwEAmpwYAAAAB3RJTUUH4QYcDhkW4QlOQwAAHBNJREFUeNrt3XuMlNX9P/DP4u7OLkshVEGkgKyVWzTNsloTbYtGaIgmRGM1tRq1TbVSaqm2qdREwbSJoQlQGm9NiNSmCSyRRhNtm0YTTUyq1jWhtFDF2lKrqcKuCsJegfP74/vbDbC3mZ2Zvcy8Xslkk5lnn3l2zjPn8z7nuWxFSikFAFBWJvgIAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABg9FRmu2BFRUVWy7mxIAAUXqHrcOVY/kOECQBKvWiPVq2rLPUPV4gAQH0pwQBQzAYWHgD08aWq0u5R2J1KaADQ7woADLrj2mkBfaO+UQCg3y+DLwKgP0MA4JQvlC8ToE9iuNwIqIS+hNleIwqgn0EAKPFU7ssK6EPot41TgedsCnmnokLcQCGf6ziHc/lILmepZnPyy2DH0Pqbeut5LqUUFRUVpuSAvPrznr6kv35msD4pm/4sl75xOP1uPn18IepLIWtdtusyA0BRdhZAH0LpcBKgLyyAfkUAwJcSQN8jAOCLBqC/EgDwZQHQ5wkA2MEB9KMCgJ0QAH21AGDHAEB9EAA0DgACxDgNAIW66xEAkFvdLHQNdidAAChDAgAACAAAgAAAAAgAAIAAAAAIAACAAAAAjFk53QgomxsMuFkQAORmNGqnGQAAKEMCAAAIAACAAAAACAAAgAAAAAgAAIAAAACMWZW5/sJQNytwkx8AyE0hbrSXa/01AwAAZUgAAAABAAAQAAAAAQAASklFRUVW/4lPAGDEHTt2LNatWxf19fWRyWRizpw5PhQg74KX7WO0ty+TycTnPve5uP7666O5uXlUC3iphYWKlON1A6N9GWCuH/5IX5bYs32Fet/169fHfffd1+dvKvT7AOUTAMZDH9qfqqqqeOaZZ+Lqq68uWJ+by++Pdr876pcBppQGfYy1NDnePfnkkxERsXXr1ujo6FDwgbz0129n89pobWdHR0f8/e9/j69//evR3d3d74CoVPrFoernUPU3VxPsvGPb/v37o6qqKr75zW9GJpPRewEjpru7OzZt2hSNjY1RV1cXdXV10djYGJs3b45jx471W7w+/PDDWLlyZcyYMSNqa2vjS1/6Urz22mvD3oZMJhMXXHBBPPLIIxERsW/fvn7f92Tt7e2xdu3aqK+vj6qqqpg7d2488MADMWnSpAEHou+++27ceuutMXXq1JgxY0bcdddd0dHR0e/ou2QGtqnAImLQRzEMtO6urq60cePGtHjx4jRx4sQ0ceLEtHjx4vSLX/widXd397uODz74IN15553p7LPPTjU1Nemyyy5Lr776at7bcrqdO3emZcuWpalTp6aqqqp07rnnptWrV6eWlpZBP8uhPmOAQvShXV1dadmyZQP2NcuXLz+lH+15fsqUKX2WraurS2+//XZefeg//vGPFBGpvr5+0OW7urrSV77ylaz7yZ7nzjzzzD7LrVu3LqvaNl7rZ8kGgLG2857sjjvuGHC75s+fnz755BMBABjVALBhw4YUEWn69OmpqakptbS0pJaWlrRt27Y0bdq0FBFp06ZNfdZx/fXXp1dffTW1t7en3bt3p8bGxhQR6dvf/vawtuXQoUPpxRdf7F3P2rVrB13+l7/8Ze92P/300+nw4cPp4MGDaefOnYMGgMsvvzzt2rUrHTp0KK1ZsyZFRJo3b96wBngCwCgHgLGy857uySefTBGRZs2alXbs2JEOHDiQ2traUnNzc7riiitSRKT77rtvyPUp+EAx+9CGhoYUEempp57qs3xTU1OKiLR48eIh+6SXX36535H7cGrIddddlzo6Ogbd9osvvjhFRNq+fXtWf2fPc8ePH+997sMPP0wRkTKZjAAwHgPAaO+8A7n00ktTRKTm5uY+r/33v/9NEZEWLVokAACjGgBqa2tTRKTW1tY+y7e0tKSISBMnThyyT/r0009TRKTq6uq8asj69euz2va6uroUEengwYM5BYCTdXd3Z73seK6fJXsfgLfeeisiIq688so+ry1btuyUZQbT0NAQERHvv/9+QbZr9+7dERFx8cUX9zmRZPbs2RER8e9//9vZR8CYOfE6m+cG0nMiXVVV1bBO+H700UcjImLz5s2xf//+IX+vq6srIiJqa2uH/TdXVlaWRdtOsPMWZ+cdyIkTJ7J+T4DRMn/+/IiIeP755/u89sILL0RExIIFC4Zcz0svvRQREYsWLRrWdqxatSruueee+OCDD+Kqq66Kjz76aNDlZ86cecpgqxg6OzvLNwCM5HWK433nPd3ChQsjIuKdd94pyOdUKjsiMLbccsstERGxevXqaGpqitbW1mhtbY2mpqZYvXr1Kcuc7Pe//3188sknceTIkXjuued6l73pppuGvS0bNmyIa6+9Nt5888245pprBu33emZ977777tizZ08cPXo0mpub47bbbsv7M+kJFw899FC0tbUVZcCahnmfneG+YUGPU4y1M1inTZuWtm/f3nsS4Pbt2wc9CfC5555LH3/8cfr000/Ts88+m84555w+y+ZzDsBjjz2WIiLNnj07PfHEE+k///lPam9vT21tbWnPnj3p0UcfTRdddNGQ65s5c2bvGbFHjx51EBMoaB/a1dWVli5dmvOVVP09LrvsstTZ2ZlXH3r06NHeE/xuuOGGdOLEiX6X37dvX5o0aVKfbaiurs7puH5/z69cuXJU616ha2/JBoDR3nkHehw/fjx961vfyvpkj4E+19HeEYHSDgA9/ejGjRtTQ0NDqq2tTbW1tamhoSFt2rRpwHuprFmzJn3+859PmUwmzZo1K/34xz9OR44cyXtbUkrpf//7X5ozZ06KiHT33XcPuPzrr7+elixZkmpqatLkyZPTihUr0htvvJEiItXU1Aw7ABw5ciTde++96bzzzkuZTCZNmTIlXXLJJeM2AOT8vwB6DgHkcsy92Aa6P3N3d3c8/PDD8dvf/rb3hL8FCxbErbfeGt///vdPOdGjZx1r1qyJnTt3xnvvvRfTpk2Lb3zjG7Fu3bqoq6vL+7M5eRv/+Mc/xhNPPBGvvfZaHDhwICoqKqK+vj6WLl0aN954Y3z5y18e9G87evRo/PSnP42dO3fG+++/HzU1NbFgwYK87rgFUOh+eKzYs2dPXHjhhTFv3rw+dxMcb59xoWpvSQSActh5AfShuevu7o69e/fGqlWr4s9//nPcdddd8fDDDwsAAoAAAFBqfehANWrGjBnxxhtv9J7MV+4BYILdFoBSsmLFipg9e3bU1tZGdXV1zJ07N+688854/fXXx23xL0qgSAWOa+N1dsAMAAClNsofTKWPe+yHEwAoNIcAAEAAAAAEAABAAAAABAAAYJwa9lUAA12q4Gx6ABiewWpooeuuGQAAKEMCAAAIAACAAAAACAAAQGkY9lUAA5116OoAABi+kaqjZgAAoAwJAAAgAAAAAgAAIAAAAKWhMp9f7u9MRWf7A8Dw5XKVXT411wwAAJQhAQAABAAAQAAAAEpSXicB9nfygRMDASA/I1FLzQAAQBkSAABAAAAABAAAoCRV5ruC009UcMIfAOQnm5Ps8623ZgAAoAwJAAAgAAAA5SDvcwBOPwbhnAAAyF+x66kZAAAoQwIAAAgAAEA5qCzESk4+TuGYPwDkb7Bz7ApRa80AAEAZEgAAoAwV5BDAyVMRDgcAQGEUs6aaAQCAMiQAAEAZqizUinqmKUz7A0Bh9HeIvVB11gwAAJgByD+lmAkAgMIpVl01AwAAZgDyTylG/gBQOCfPsBeyxk4o9Eae/u8LAYCxN8CeMNY3EADMAhR+gF2RVGwAKDtOAgQAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAAiqTSRwBDq6io6Pf5lJIPR7uAGQAolwKT6zJoFxAAYBwVmWxGkiklxUa7gAAApSLbApJtQUK7gAAARppoFxAAwEgT7QICABhpol1AAAAjTbQLCABgpKldtAsIAGCkqV20CwgAYKSpXbQLCABgpKldAAEAjDS1CyAAgJGmdgEBAMpg5Jjrz1xHmrn+RLuAAABFLjI9BSHXn6O5bu2iXaBo379kjgyyHqX2N8pEu4AAAOOoSIwV5foV1C4wuhwCAAABAAAQAAAAAQAAEAAAgHHKVQCQzRfF5WbaBcwAwPgtFMW4M1wx161dtAsIAJCH4d4RLpdbzhZj3dpFu0DRArhDADD4KHKw28MqGtoFzABAiRaZbEeaaBcQAKAE+Lez2gUEADDSNNLULiAAgJGmkaZ2AQEAjDSNNLULCABgpIl2AQEAjDTRLiAAgJEm2gUEADDSRLuAAABGmmgXEADASBPtAgIAGGmiXUAAgBEpNoVYBu0CY43/BggAZgAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAAAQAAEAAAAAEAABAAAAABAAAQAAAAAQAAEAAAAAEAABAAAAARkaljwCGVlFR0e/zKSUfjnYBMwBQLgUm12XQLiAAwDgqMtmMJFNKio12AQEASkW2BSTbgoR2AQEAjDTRLiAAgJEm2gUEADDSRLuAAABGmmgXEADASFO7aBcQAMBIU7toFxAAwEhTu2gXEADASFO7AAIAGGlqF0AAACNN7QICAJTByDHXn7mONHP9iXYBAQCKXGR6CkKuP0dz3dpFu0DRvn/JHBlkPUrtb5SJdgEBAMZRkRgryvUrqF1gdDkEAAACAAAgAAAAAgAAIAAAAOOUqwAgmy+Ky820C5gBgPFbKIpxZ7hirlu7aBcQACAPw70jXC63nC3GurWLdoGiBXCHAGDwUeRgt4dVNLQLmAGAEi0y2Y400S4gAEAJ8G9ntQsIAGCkaaSpXUAAACNNI03tAgIAGGkaaWoXEADASBPtAgIAGGmiXUAAACNNtAsIAGCkiXYBAQCMNNEuIACAkSbaBQQAMNJEu4AAACNSbAqxDNoFxhr/DRAAzAAAAAIAACAAAAACAAAgAAAAAgAAIAAAAAIAACAAAAACAAAgAAAAAgAAIAAAAAIAACAAAAACAAAgAACAAAAACAAAgAAAAAgAAIAAAAAIAACAAAAACAAAgAAAAAgAAIAAAAAIAACAAAAACAAAgAAAAAgAAIAAAAACAAAgAAAAAgAAIAAAAAIAACAAAAACAAAgAAAAAgAAIAAAAAIAACAAAAACAAAgAAAAAgAAIAAAAAIAAAgAAIAAAAAIAACAAAAACAAAgAAAAAgAAIAAAAAIAACAAAAACAAAgAAAAIzJAFBRUeETBYAiKHSNLWgASCkJAQBQhOKfUhrbMwCF3kAAKHfFGGBXFmsDBQEAKMzguhh11UmAAFCGCjYDYOQPAIXXU1cLXWfNAACAGYDCJRSzAQCQn2LWVDMAAFCGBAAAKEMFOQRg2h8ACu/kmlroWmsGAADKkAAAAGWoIIcABpqiOP01ACA7xa6nZgAAoAwJAAAgAAAA5SDvcwAc8weAwju9nha63poBAIAyJAAAgAAAAJSDvM8BGOoYRX/LAAADG4laagYAAMqQAAAAAgAAIAAAACUpr5MAnfAHAIXXXy0tdM01AwAAZUgAAAABAAAQAACAkpTXSYDZnqQw0LIAwOjUUTMAAFCGBAAAEAAAAAFgDKqoqDjlUVNTE+eff3786Ec/ikOHDuW93lxfA0AfPS7bKhX4rIJin7wwWCNfc8018cwzz+S13sFObHQiI4A+erQ+wzFzEuDpKa/nkVLq91FoKaU4ceJEfPrpp/H4449HRMSf/vQn3z6AMUAfnd9n199joLo74gFgrKSkSZMmxc033xwREdOnTz/l9e7u7ti0aVM0NjZGXV1d1NXVRWNjY2zevDmOHTvWb9oa7MPt7/ls3+P0dezevTuuvfbamDx5ckyePDluuumm+Pjjj+Oll16K5cuXx2c+85mYMWNGrFy5Mtrb230jAH20PrrgSWNYIqLfR7Gd/j6HDx9O999/f4qItGXLlt7nu7q60rJlywbczuXLl6fu7u5B/5ahHrm8x+nbP2HChD7Lz507t9/13H///QlgPNBHF/+zLVTdHbcB4PRHY2Njev/993uX27BhQ4qINH369NTU1JRaWlpSS0tL2rZtW5o2bVqKiLRp06YBd9rBdujhvsfJ61qyZEnatWtX6ujoSA8++GC/z69duzZFRDr//PP1KsC4CgD6aAFgxD6AiEhTp05Nu3btSiml1NDQkCIiPfXUU33W0dTUlCIiLV68OK+dK9f3OHldx48f733uwIEDKSJSVVXVKc9/+OGHKSJSJpPRqwDjOgDoowWAgk8vdXV1pf3796fbb7+9d0onpZRqa2tTRKTW1tY+62hpaUkRkSZOnJjXzpXrewy0rmPHjuX0PMBYDwD66LEfACbkce5ATmcpFusazaqqqjj33HNj48aNERHx6quvDnnZRCrwVQn5vscZZ5yR0/MA44U+OvcTJ0fqKruSuRPg4cOHT2nU+fPnR0TE888/32fZF154ISIiFixY0Oe1zs7OAd/j9NeG+x4A5UYfPfaM+wDQ1dUVu3fvjttvvz0iIpYsWRIREbfccktERKxevTqampqitbU1Wltbo6mpKVavXn3KMhERM2fOjIiIhx56KNra2k55j4Fey/U9AMqNPnoMG8kTQIq9/smTJ6e//e1vvcedli5dmvXlHytXrhxwmwd6Ldf3GOxYVa7PA4y3kwD10aNfQwtyEuBYCQBnnHFGOuecc9LNN9+c3n777VOW7erqShs3bkwNDQ2ptrY21dbWpoaGhrRp06Y+jX7kyJF07733pvPOOy9lMpk0ZcqUdMkllwz5Wi7vIQAA5RYA9NFjNwCM2P8COPnYDwAwujXUvwMGgDJUaZQPAGPDYDW00LXXDAAAlCEBAAAEAABAAAAABAAAQAAoCceOHYt169ZFfX19ZDKZmDNnjr0CoAwV8x/Xjcm/dyRvBBSR/2WCPesv1GavX78+7rvvvj7bWOj3ASC/elJdXR1nnXVWXHrppfGTn/wkLr744oLWh1x+vxg1otj1UwA4zcKFC+Ott96KrVu3xk033RSZTKZojQtAYepJVVVVPPPMM3H11VcLAALA8NTU1MSJEyeis7PzlG0XAADGRgDo6Yc7Ozvjn//8Z/zsZz+LHTt2xBe+8IX461//WrT3K/UAULLnAPzud7+Lr371q/HZz342qqurY+7cufGDH/wgWltbT/mwOzs7o7u7OyZMmND74Z8eBE5+ADA6MplMXHDBBfHII49ERMS+ffv6FNDT++n29vZYu3Zt1NfXR1VVVcydOzceeOCBmDRp0oB9+rvvvhu33nprTJ06NWbMmBF33XVXdHR09Fuox3ONKMkZgO985zuxZcuWfl+bP39+/OUvf4kpU6b0u60nH/8fiQQGQG79/ZtvvhmLFi2K+vr6+Ne//jXg8t3d3bF06dJ4+eWXs+rTe37/zDPPPGWwGBGxbt26ePDBB4esc/nUCDMAefrNb34TW7ZsiVmzZsWOHTviwIED0dbWFs3NzXHFFVfEvn374uc//3mfD/P//2vkAZ8/+XUARt7hw4fjpZdeiptvvjkiIm655ZZBl3/88cfj5ZdfjunTp8fTTz8dhw8fjoMHD8bOnTsH/b0LL7wwdu3aFYcOHYo1a9ZERMS2bdv6LcTjuUaU3AzAZZddFq+88ko0NzfHRRdddMpr7733XsyePTsWLVoUe/fuHXR9zgEAGBszAP257rrrYtu2bb0nbvfXb3/xi1+M5ubm2L59e9x4441D9vE9zx0/fjwmTPi/8fGBAwfi7LPPjkwm0+9hACcBjqEAMGnSpDh69Oig66ipqYn29nYBAGAcBoD169f3jswH67d76sHBgwfjrLPOyjoAnPzcsWPHoqqqKqtlx1sAmDCcDRzscfp0yEhPj5w4cWLIZU5OcQCMbT3149FHH42IiM2bN8f+/fuH/L2urq6IiKitrR32e1dWVo743znQY6j6W/QAMNYtXLgwIiLeeeedggSRzs5O3z6AMWDVqlVxzz33xAcffBBXXXVVfPTRR4MuP3PmzIiI2L17d9G2aTzXiJILAHfccUdERFxxxRWxdevWePfdd6OjoyPa29tj79698dhjj/W5e9RgO85DDz0UbW1tvnkAY8CGDRvi2muvjTfffDOuueaaQQvwlVdeGRERd999d+zZsyeOHj0azc3Ncdttt+W9HaVQI3I+B2Ckj1Hk+v7Hjx+P22+/PX79619ntZ0DHcf57ne/G7/61a9G/O8DIAbtn9va2uLyyy+P5ubmuOGGG2LHjh2nTIP3LP/2229HY2NjHDly5JTfr66u7j08kM1x/f6eH40aUej6W3IzABMmTIitW7fGH/7wh/ja174Ws2bNiurq6shkMrFw4cL43ve+N+g1oSenzHvvvTfOO++8yGQyMWXKlLjkkkt8IwFG2cSJE+PZZ5+NOXPmxFNPPRU//OEP+11u3rx58eKLL8aSJUuipqYmJk+eHCtWrIhXXnklIv7vhPB8ZiLGe40YdzMAAJCPPXv2xIUXXhjz5s3rczfB8TAjUqj6W2lXAKAcdHd3x969e2PVqlUREbF8+fKy/jzMAABQukVugJo1Y8aMeOONN3pP5ivHGYCczwHI9zpF/1AHgJGyYsWKmD17dtTW1vb+Y7g777wzXn/99TFV/LOpnYW+z86I3wnQLAEAjH7tnOBjB4DyIwAAgAAAAAgAAIAAAAAIAACAAAAAjBc53QrYNf4AUHjZ1M5C12AzAABQhgQAABAAAAABAAAQAAAAAQAAEAAAAAEAABizcroRUKFuVJDtugCg1I1W3awc7x+KIAGAGlViAWAkGkiAAFDAy7E+VNoFhrdzCA4A+moBgCF3RDshgH5UAGDAndwODujzEADo84XxZQH0VwgAvmy+aIC+h2FxI6AS/GJme00pgH5FAKDEUzqAPoSTVRZ6R8h2OqhQN1DI9zrO4V4+kutZqtme/DLUMbSBpt56nu95rqKiIlJKvT8Bci3+p/chp/czg/VJ2fRnufSNw+l38+njC1FfClnrilGHzQCUqJ4dQPEH9CHkNQPA+PiyAuhnEAB8AQH0SQgAvkwA+jMEAF8EAH2jAICdFkC/KwDYoQDQxwsAGhcA4UEA8MECwLiudRVJJQaAsuNOgAAgAAAAAgAAIAAAAAIAACAAAAACAAAgAAAAAgAAMIr+HyP6AtxN3EzKAAAAAElFTkSuQmCC" + } + ], + "materials" : [ + { + "name" : "BackPlaneMat", + "doubleSided": true, + "pbrMetallicRoughness" : { + "baseColorFactor" : [ + 0.16000001668930075, + 0.16000001668930075, + 0.16000001668930075, + 1.0 + ], + "metallicFactor" : 0.0 + } + }, + { + "name" : "BottomLeftMat", + "doubleSided": true, + "pbrMetallicRoughness" : { + "baseColorFactor" : [ + 0.0, + 0.16000000476837162, + 0.800000011920929, + 1.0 + ], + "baseColorTexture" : { + "index" : 0 + }, + "metallicFactor" : 0.0 + } + }, + { + "name" : "BottomRightMat", + "doubleSided": true, + "pbrMetallicRoughness" : { + "baseColorFactor" : [ + 0.0, + 0.800000011920929, + 0.0, + 1.0 + ], + "baseColorTexture" : { + "index" : 0 + }, + "metallicFactor" : 0.0 + } + }, + { + "name" : "TopLeftMat", + "doubleSided": true, + "pbrMetallicRoughness" : { + "baseColorFactor" : [ + 0.800000011920929, + 0.800000011920929, + 0.0, + 1.0 + ], + "baseColorTexture" : { + "index" : 0 + }, + "metallicFactor" : 0.0 + } + }, + { + "name" : "TopRightMat", + "doubleSided": true, + "pbrMetallicRoughness" : { + "baseColorFactor" : [ + 0.800000011920929, + 0.08000000238418581, + 0.0, + 1.0 + ], + "baseColorTexture" : { + "index" : 0 + }, + "metallicFactor" : 0.0 + } + } + ], + "meshes" : [ + { + "name" : "TopRightMesh", + "primitives" : [ + { + "attributes" : { + "NORMAL" : 2, + "POSITION" : 1, + "TEXCOORD_0" : 3 + }, + "indices" : 0, + "material" : 4 + } + ] + }, + { + "name" : "TopLeftMesh", + "primitives" : [ + { + "attributes" : { + "NORMAL" : 6, + "POSITION" : 5, + "TEXCOORD_0" : 7 + }, + "indices" : 4, + "material" : 3 + } + ] + }, + { + "name" : "BottomRightMesh", + "primitives" : [ + { + "attributes" : { + "NORMAL" : 10, + "POSITION" : 9, + "TEXCOORD_0" : 11 + }, + "indices" : 8, + "material" : 2 + } + ] + }, + { + "name" : "BackPlaneMesh", + "primitives" : [ + { + "attributes" : { + "NORMAL" : 14, + "POSITION" : 13 + }, + "indices" : 12, + "material" : 0 + } + ] + }, + { + "name" : "BottomLeftMesh", + "primitives" : [ + { + "attributes" : { + "NORMAL" : 17, + "POSITION" : 16, + "TEXCOORD_0" : 18 + }, + "indices" : 15, + "material" : 1 + } + ] + } + ], + "nodes" : [ + { + "mesh" : 3, + "name" : "BackPlane" + }, + { + "mesh" : 4, + "name" : "BottomLeftObj" + }, + { + "mesh" : 2, + "name" : "BottomRightObj" + }, + { + "mesh" : 1, + "name" : "TopLeftObj" + }, + { + "mesh" : 0, + "name" : "TopRightObj" + } + ], + "samplers" : [ + {} + ], + "scene" : 0, + "scenes" : [ + { + "name" : "Scene", + "nodes" : [ + 0, + 2, + 1, + 4, + 3 + ] + } + ], + "textures" : [ + { + "sampler" : 0, + "source" : 0 + } + ] +} diff --git a/renders/color.PNG b/renders/color.PNG new file mode 100644 index 0000000..35100d9 Binary files /dev/null and b/renders/color.PNG differ diff --git a/renders/depth.PNG b/renders/depth.PNG new file mode 100644 index 0000000..5e22ccf Binary files /dev/null and b/renders/depth.PNG differ diff --git a/renders/instancing.PNG b/renders/instancing.PNG new file mode 100644 index 0000000..e0e3c20 Binary files /dev/null and b/renders/instancing.PNG differ diff --git a/renders/lines.PNG b/renders/lines.PNG new file mode 100644 index 0000000..75a8824 Binary files /dev/null and b/renders/lines.PNG differ diff --git a/renders/lines_2.PNG b/renders/lines_2.PNG new file mode 100644 index 0000000..347704e Binary files /dev/null and b/renders/lines_2.PNG differ diff --git a/renders/normals.PNG b/renders/normals.PNG new file mode 100644 index 0000000..90656da Binary files /dev/null and b/renders/normals.PNG differ diff --git a/renders/performance_analysis.png b/renders/performance_analysis.png new file mode 100644 index 0000000..a002e15 Binary files /dev/null and b/renders/performance_analysis.png differ diff --git a/renders/pipeline.png b/renders/pipeline.png new file mode 100644 index 0000000..b55787c Binary files /dev/null and b/renders/pipeline.png differ diff --git a/renders/points.PNG b/renders/points.PNG new file mode 100644 index 0000000..d5fda06 Binary files /dev/null and b/renders/points.PNG differ diff --git a/renders/render_cesium.PNG b/renders/render_cesium.PNG new file mode 100644 index 0000000..518a296 Binary files /dev/null and b/renders/render_cesium.PNG differ diff --git a/renders/render_di.PNG b/renders/render_di.PNG new file mode 100644 index 0000000..65af634 Binary files /dev/null and b/renders/render_di.PNG differ diff --git a/renders/render_duck.PNG b/renders/render_duck.PNG new file mode 100644 index 0000000..706aa54 Binary files /dev/null and b/renders/render_duck.PNG differ diff --git a/renders/ssaa_1.PNG b/renders/ssaa_1.PNG new file mode 100644 index 0000000..d11a59b Binary files /dev/null and b/renders/ssaa_1.PNG differ diff --git a/renders/ssaa_2.PNG b/renders/ssaa_2.PNG new file mode 100644 index 0000000..b6c5339 Binary files /dev/null and b/renders/ssaa_2.PNG differ diff --git a/renders/triangles.PNG b/renders/triangles.PNG new file mode 100644 index 0000000..28774a4 Binary files /dev/null and b/renders/triangles.PNG differ diff --git a/renders/triangles_2.PNG b/renders/triangles_2.PNG new file mode 100644 index 0000000..1cb1e72 Binary files /dev/null and b/renders/triangles_2.PNG differ diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index a57f69f..779ca22 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -1,4 +1,6 @@ set(SOURCE_FILES + "CameraControls.h" + "CameraControls.cpp" "rasterize.cu" "rasterize.h" "rasterizeTools.h" @@ -6,5 +8,5 @@ set(SOURCE_FILES cuda_add_library(src ${SOURCE_FILES} - OPTIONS -arch=sm_20 + OPTIONS -arch=sm_60 ) diff --git a/src/main.cpp b/src/main.cpp index 7986959..cb4fae1 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -14,6 +14,7 @@ #define TINYGLTF_LOADER_IMPLEMENTATION #include + //------------------------------- //-------------MAIN-------------- //------------------------------- @@ -45,11 +46,12 @@ int main(int argc, char **argv) { } if (!ret) { + printf("Failed to parse glTF\n"); + getchar(); return -1; } - frame = 0; seconds = time(NULL); fpstracker = 0; @@ -116,12 +118,16 @@ void runCuda() { * glm::rotate(y_angle, glm::vec3(0.0f, 1.0f, 0.0f)); glm::mat3 MV_normal = glm::transpose(glm::inverse(glm::mat3(V) * glm::mat3(M))); + glm::mat3 M_inverseTranspose = glm::inverse(glm::transpose(glm::mat3(M))); glm::mat4 MV = V * M; glm::mat4 MVP = P * MV; - cudaGLMapBufferObject((void **)&dptr, pbo); - rasterize(dptr, MVP, MV, MV_normal); - cudaGLUnmapBufferObject(pbo); + + glm::vec3 eyePos = glm::vec3(x_trans, y_trans, z_trans); + + cudaGLMapBufferObject((void **)&dptr, pbo); + rasterize(dptr, MVP, MV, MV_normal, M_inverseTranspose, eyePos); + cudaGLUnmapBufferObject(pbo); frame++; fpstracker++; diff --git a/src/rasterize.cu b/src/rasterize.cu index 1262a09..c9907ce 100644 --- a/src/rasterize.cu +++ b/src/rasterize.cu @@ -20,101 +20,129 @@ namespace { - typedef unsigned short VertexIndex; - typedef glm::vec3 VertexAttributePosition; - typedef glm::vec3 VertexAttributeNormal; - typedef glm::vec2 VertexAttributeTexcoord; - typedef unsigned char TextureData; - - typedef unsigned char BufferByte; - - enum PrimitiveType{ - Point = 1, - Line = 2, - Triangle = 3 - }; - - struct VertexOut { - glm::vec4 pos; - - // TODO: add new attributes to your VertexOut - // The attributes listed below might be useful, - // but always feel free to modify on your own - - glm::vec3 eyePos; // eye space position used for shading - glm::vec3 eyeNor; // eye space normal used for shading, cuz normal will go wrong after perspective transformation - // glm::vec3 col; - glm::vec2 texcoord0; - TextureData* dev_diffuseTex = NULL; - // int texWidth, texHeight; - // ... - }; - - struct Primitive { - PrimitiveType primitiveType = Triangle; // C++ 11 init - VertexOut v[3]; - }; - - struct Fragment { - glm::vec3 color; - - // TODO: add new attributes to your Fragment - // The attributes listed below might be useful, - // but always feel free to modify on your own - - // glm::vec3 eyePos; // eye space position used for shading - // glm::vec3 eyeNor; - // VertexAttributeTexcoord texcoord0; - // TextureData* dev_diffuseTex; - // ... - }; - - struct PrimitiveDevBufPointers { - int primitiveMode; //from tinygltfloader macro - PrimitiveType primitiveType; - int numPrimitives; - int numIndices; - int numVertices; - - // Vertex In, const after loaded - VertexIndex* dev_indices; - VertexAttributePosition* dev_position; - VertexAttributeNormal* dev_normal; - VertexAttributeTexcoord* dev_texcoord0; - - // Materials, add more attributes when needed - TextureData* dev_diffuseTex; - int diffuseTexWidth; - int diffuseTexHeight; - // TextureData* dev_specularTex; - // TextureData* dev_normalTex; - // ... - - // Vertex Out, vertex used for rasterization, this is changing every frame - VertexOut* dev_verticesOut; - - // TODO: add more attributes when needed - }; + typedef unsigned short VertexIndex; + typedef glm::vec3 VertexAttributePosition; + typedef glm::vec3 VertexAttributeNormal; + typedef glm::vec2 VertexAttributeTexcoord; + typedef unsigned char TextureData; + + typedef unsigned char BufferByte; + +#define BYTES_PER_PIXEL 3 + +#define NUM_INSTANCES 1 + +#define AA_SUPER_SAMPLE + //#define AA_MULTI_SAMPLE + +#define SSAA_LEVEL 1 +#define MSAA_LEVEL 2 + +#define NO_INTERPOLATION +#define BI_LINEAR_INTERPOLATION + + enum PrimitiveType + { + Point = 1, + Line = 2, + Triangle = 3 + }; + + struct VertexOut + { + glm::vec4 pos; + glm::vec3 eyePos; // eye space position used for shading + glm::vec3 eyeNor; // eye space normal used for shading, cuz normal will go wrong after perspective transformation + + glm::vec3 color; + glm::vec2 texcoord0; + }; + + struct Primitive + { + int instanceId; + PrimitiveType primitiveType = Triangle; // C++ 11 init + VertexOut v[3]; + TextureData* dev_diffuseTex; + int diffuseTexWidth; + int diffuseTexHeight; + }; + + struct Fragment + { + glm::vec3 color; + glm::vec3 eyePos; + glm::vec3 eyeNor; + VertexAttributeTexcoord texcoord0; + TextureData* dev_diffuseTex; + int diffuseTexWidth; + int diffuseTexHeight; + }; + + struct PrimitiveDevBufPointers + { + int primitiveMode; //from tinygltfloader macro + + PrimitiveType primitiveType; + int numPrimitives; + int numIndices; + int numVertices; + int numInstances; + + int vertexOutStartIndex; + + + // Vertex In, const after loaded + VertexIndex* dev_indices; + VertexAttributePosition* dev_position; + VertexAttributeNormal* dev_normal; + VertexAttributeTexcoord* dev_texcoord0; + + // Materials, add more attributes when needed + TextureData* dev_diffuseTex; + int diffuseTexWidth; + int diffuseTexHeight; + // TextureData* dev_specularTex; + // TextureData* dev_normalTex; + // ... + + // Vertex Out, vertex used for rasterization, this is changing every frame + VertexOut* dev_verticesOut; + + // TODO: add more attributes when needed + }; } static std::map> mesh2PrimitivesMap; +static int ssaaWidth = 0; +static int ssaaHeight = 0; + +static int msaaWidth = 0; +static int msaaHeight = 0; static int width = 0; static int height = 0; +static int originalWidth = 0; +static int originalHeight = 0; + static int totalNumPrimitives = 0; static Primitive *dev_primitives = NULL; + static Fragment *dev_fragmentBuffer = NULL; + static glm::vec3 *dev_framebuffer = NULL; +static glm::vec3 *dev_AAFrameBuffer = NULL; -static int * dev_depth = NULL; // you might need this buffer when doing depth test +static float * dev_depth = NULL; // you might need this buffer when doing depth test +static unsigned int* dev_mutex = NULL; /** * Kernel that writes the image to the OpenGL PBO directly. */ -__global__ +__global__ void sendImageToPBO(uchar4 *pbo, int w, int h, glm::vec3 *image) { int x = (blockIdx.x * blockDim.x) + threadIdx.x; int y = (blockIdx.y * blockDim.y) + threadIdx.y; @@ -133,53 +161,65 @@ void sendImageToPBO(uchar4 *pbo, int w, int h, glm::vec3 *image) { } } -/** -* Writes fragment colors to the framebuffer -*/ -__global__ -void render(int w, int h, Fragment *fragmentBuffer, glm::vec3 *framebuffer) { - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - int index = x + (y * w); - if (x < w && y < h) { - framebuffer[index] = fragmentBuffer[index].color; - - // TODO: add your fragment shader code here - - } -} /** * Called once at the beginning of the program to allocate memory. */ -void rasterizeInit(int w, int h) { - width = w; - height = h; - cudaFree(dev_fragmentBuffer); - cudaMalloc(&dev_fragmentBuffer, width * height * sizeof(Fragment)); - cudaMemset(dev_fragmentBuffer, 0, width * height * sizeof(Fragment)); +void rasterizeInit(int w, int h) +{ + originalWidth = w; + originalHeight = h; + + ssaaWidth = w * SSAA_LEVEL; + ssaaHeight = h * SSAA_LEVEL; + + msaaWidth = w * MSAA_LEVEL; + msaaHeight = h * MSAA_LEVEL; + +#ifdef AA_SUPER_SAMPLE + width = ssaaWidth; + height = ssaaHeight; +#endif // AA_SUPER_SAMPLE + +#ifdef AA_MULTI_SAMPLE + width = msaaWidth; + height = msaaHeight; +#endif // AA_MULTI_SAMPLE + + cudaFree(dev_fragmentBuffer); + cudaMalloc(&dev_fragmentBuffer, width * height * sizeof(Fragment)); + cudaMemset(dev_fragmentBuffer, 0, width * height * sizeof(Fragment)); + cudaFree(dev_framebuffer); - cudaMalloc(&dev_framebuffer, width * height * sizeof(glm::vec3)); + cudaMalloc(&dev_framebuffer, width * height * sizeof(glm::vec3)); cudaMemset(dev_framebuffer, 0, width * height * sizeof(glm::vec3)); - - cudaFree(dev_depth); - cudaMalloc(&dev_depth, width * height * sizeof(int)); - checkCUDAError("rasterizeInit"); + cudaFree(dev_AAFrameBuffer); + cudaMalloc(&dev_AAFrameBuffer, originalWidth * originalHeight * sizeof(glm::vec3)); + cudaMemset(dev_AAFrameBuffer, 0, originalWidth * originalHeight * sizeof(glm::vec3)); + + cudaFree(dev_depth); + cudaMalloc(&dev_depth, width * height * sizeof(float)); + + cudaFree(dev_mutex); + cudaMalloc(&dev_mutex, width * height * sizeof(unsigned int)); + cudaMemset(dev_mutex, 0, width * height * sizeof(unsigned int)); + + checkCUDAError("rasterizeInit"); } __global__ -void initDepth(int w, int h, int * depth) +void initDepth(int w, int h, float * depth) { - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - - if (x < w && y < h) - { - int index = x + (y * w); - depth[index] = INT_MAX; - } + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + + if (x < w && y < h) + { + int index = x + (y * w); + depth[index] = FLT_MAX; + } } @@ -187,550 +227,1098 @@ void initDepth(int w, int h, int * depth) * kern function with support for stride to sometimes replace cudaMemcpy * One thread is responsible for copying one component */ -__global__ +__global__ void _deviceBufferCopy(int N, BufferByte* dev_dst, const BufferByte* dev_src, int n, int byteStride, int byteOffset, int componentTypeByteSize) { - - // Attribute (vec3 position) - // component (3 * float) - // byte (4 * byte) - - // id of component - int i = (blockIdx.x * blockDim.x) + threadIdx.x; - - if (i < N) { - int count = i / n; - int offset = i - count * n; // which component of the attribute - - for (int j = 0; j < componentTypeByteSize; j++) { - - dev_dst[count * componentTypeByteSize * n - + offset * componentTypeByteSize - + j] - - = - - dev_src[byteOffset - + count * (byteStride == 0 ? componentTypeByteSize * n : byteStride) - + offset * componentTypeByteSize - + j]; - } - } - + + // Attribute (vec3 position) + // component (3 * float) + // byte (4 * byte) + + // id of component + int i = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (i < N) { + int count = i / n; + int offset = i - count * n; // which component of the attribute + + for (int j = 0; j < componentTypeByteSize; j++) { + + dev_dst[count * componentTypeByteSize * n + + offset * componentTypeByteSize + + j] + + = + + dev_src[byteOffset + + count * (byteStride == 0 ? componentTypeByteSize * n : byteStride) + + offset * componentTypeByteSize + + j]; + } + } + } __global__ void _nodeMatrixTransform( - int numVertices, - VertexAttributePosition* position, - VertexAttributeNormal* normal, - glm::mat4 MV, glm::mat3 MV_normal) { - - // vertex id - int vid = (blockIdx.x * blockDim.x) + threadIdx.x; - if (vid < numVertices) { - position[vid] = glm::vec3(MV * glm::vec4(position[vid], 1.0f)); - normal[vid] = glm::normalize(MV_normal * normal[vid]); - } + int numVertices, + VertexAttributePosition* position, + VertexAttributeNormal* normal, + glm::mat4 MV, glm::mat3 MV_normal) { + + // vertex id + int vid = (blockIdx.x * blockDim.x) + threadIdx.x; + if (vid < numVertices) { + position[vid] = glm::vec3(MV * glm::vec4(position[vid], 1.0f)); + normal[vid] = glm::normalize(MV_normal * normal[vid]); + } } glm::mat4 getMatrixFromNodeMatrixVector(const tinygltf::Node & n) { - - glm::mat4 curMatrix(1.0); - - const std::vector &m = n.matrix; - if (m.size() > 0) { - // matrix, copy it - - for (int i = 0; i < 4; i++) { - for (int j = 0; j < 4; j++) { - curMatrix[i][j] = (float)m.at(4 * i + j); - } - } - } else { - // no matrix, use rotation, scale, translation - - if (n.translation.size() > 0) { - curMatrix[3][0] = n.translation[0]; - curMatrix[3][1] = n.translation[1]; - curMatrix[3][2] = n.translation[2]; - } - - if (n.rotation.size() > 0) { - glm::mat4 R; - glm::quat q; - q[0] = n.rotation[0]; - q[1] = n.rotation[1]; - q[2] = n.rotation[2]; - - R = glm::mat4_cast(q); - curMatrix = curMatrix * R; - } - - if (n.scale.size() > 0) { - curMatrix = curMatrix * glm::scale(glm::vec3(n.scale[0], n.scale[1], n.scale[2])); - } - } - - return curMatrix; + + glm::mat4 curMatrix(1.0); + + const std::vector &m = n.matrix; + if (m.size() > 0) { + // matrix, copy it + + for (int i = 0; i < 4; i++) { + for (int j = 0; j < 4; j++) { + curMatrix[i][j] = (float)m.at(4 * i + j); + } + } + } + else { + // no matrix, use rotation, scale, translation + + if (n.translation.size() > 0) { + curMatrix[3][0] = n.translation[0]; + curMatrix[3][1] = n.translation[1]; + curMatrix[3][2] = n.translation[2]; + } + + if (n.rotation.size() > 0) { + glm::mat4 R; + glm::quat q; + q[0] = n.rotation[0]; + q[1] = n.rotation[1]; + q[2] = n.rotation[2]; + + R = glm::mat4_cast(q); + curMatrix = curMatrix * R; + } + + if (n.scale.size() > 0) { + curMatrix = curMatrix * glm::scale(glm::vec3(n.scale[0], n.scale[1], n.scale[2])); + } + } + + return curMatrix; } -void traverseNode ( - std::map & n2m, - const tinygltf::Scene & scene, - const std::string & nodeString, - const glm::mat4 & parentMatrix - ) +void traverseNode( + std::map & n2m, + const tinygltf::Scene & scene, + const std::string & nodeString, + const glm::mat4 & parentMatrix +) { - const tinygltf::Node & n = scene.nodes.at(nodeString); - glm::mat4 M = parentMatrix * getMatrixFromNodeMatrixVector(n); - n2m.insert(std::pair(nodeString, M)); + const tinygltf::Node & n = scene.nodes.at(nodeString); + glm::mat4 M = parentMatrix * getMatrixFromNodeMatrixVector(n); + n2m.insert(std::pair(nodeString, M)); - auto it = n.children.begin(); - auto itEnd = n.children.end(); + auto it = n.children.begin(); + auto itEnd = n.children.end(); - for (; it != itEnd; ++it) { - traverseNode(n2m, scene, *it, M); - } + for (; it != itEnd; ++it) { + traverseNode(n2m, scene, *it, M); + } } void rasterizeSetBuffers(const tinygltf::Scene & scene) { - totalNumPrimitives = 0; + totalNumPrimitives = 0; - std::map bufferViewDevPointers; + std::map bufferViewDevPointers; - // 1. copy all `bufferViews` to device memory - { - std::map::const_iterator it( - scene.bufferViews.begin()); - std::map::const_iterator itEnd( - scene.bufferViews.end()); + // 1. copy all `bufferViews` to device memory + { + std::map::const_iterator it( + scene.bufferViews.begin()); + std::map::const_iterator itEnd( + scene.bufferViews.end()); - for (; it != itEnd; it++) { - const std::string key = it->first; - const tinygltf::BufferView &bufferView = it->second; - if (bufferView.target == 0) { - continue; // Unsupported bufferView. - } + for (; it != itEnd; it++) { + const std::string key = it->first; + const tinygltf::BufferView &bufferView = it->second; + if (bufferView.target == 0) { + continue; // Unsupported bufferView. + } - const tinygltf::Buffer &buffer = scene.buffers.at(bufferView.buffer); + const tinygltf::Buffer &buffer = scene.buffers.at(bufferView.buffer); - BufferByte* dev_bufferView; - cudaMalloc(&dev_bufferView, bufferView.byteLength); - cudaMemcpy(dev_bufferView, &buffer.data.front() + bufferView.byteOffset, bufferView.byteLength, cudaMemcpyHostToDevice); + BufferByte* dev_bufferView; + cudaMalloc(&dev_bufferView, bufferView.byteLength); + cudaMemcpy(dev_bufferView, &buffer.data.front() + bufferView.byteOffset, bufferView.byteLength, cudaMemcpyHostToDevice); - checkCUDAError("Set BufferView Device Mem"); + checkCUDAError("Set BufferView Device Mem"); - bufferViewDevPointers.insert(std::make_pair(key, dev_bufferView)); + bufferViewDevPointers.insert(std::make_pair(key, dev_bufferView)); - } - } + } + } + // 2. for each mesh: + // for each primitive: + // build device buffer of indices, materail, and each attributes + // and store these pointers in a map + { + + std::map nodeString2Matrix; + auto rootNodeNamesList = scene.scenes.at(scene.defaultScene); + + { + auto it = rootNodeNamesList.begin(); + auto itEnd = rootNodeNamesList.end(); + for (; it != itEnd; ++it) { + traverseNode(nodeString2Matrix, scene, *it, glm::mat4(1.0f)); + } + } + + // parse through node to access mesh + + auto itNode = nodeString2Matrix.begin(); + auto itEndNode = nodeString2Matrix.end(); + for (; itNode != itEndNode; ++itNode) { + + const tinygltf::Node & N = scene.nodes.at(itNode->first); + const glm::mat4 & matrix = itNode->second; + const glm::mat3 & matrixNormal = glm::transpose(glm::inverse(glm::mat3(matrix))); + + auto itMeshName = N.meshes.begin(); + auto itEndMeshName = N.meshes.end(); + + for (; itMeshName != itEndMeshName; ++itMeshName) { + + const tinygltf::Mesh & mesh = scene.meshes.at(*itMeshName); + + auto res = mesh2PrimitivesMap.insert(std::pair>(mesh.name, std::vector())); + std::vector & primitiveVector = (res.first)->second; + + // for each primitive + for (size_t i = 0; i < mesh.primitives.size(); i++) { + const tinygltf::Primitive &primitive = mesh.primitives[i]; + + if (primitive.indices.empty()) + return; + + // TODO: add new attributes for your PrimitiveDevBufPointers when you add new attributes + VertexIndex* dev_indices = NULL; + VertexAttributePosition* dev_position = NULL; + VertexAttributeNormal* dev_normal = NULL; + VertexAttributeTexcoord* dev_texcoord0 = NULL; + + // ----------Indices------------- + + const tinygltf::Accessor &indexAccessor = scene.accessors.at(primitive.indices); + const tinygltf::BufferView &bufferView = scene.bufferViews.at(indexAccessor.bufferView); + BufferByte* dev_bufferView = bufferViewDevPointers.at(indexAccessor.bufferView); + + // assume type is SCALAR for indices + int n = 1; + int numIndices = indexAccessor.count; + int componentTypeByteSize = sizeof(VertexIndex); + int byteLength = numIndices * n * componentTypeByteSize; + + dim3 numThreadsPerBlock(128); + dim3 numBlocks((numIndices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); + cudaMalloc(&dev_indices, byteLength); + _deviceBufferCopy << > > ( + numIndices, + (BufferByte*)dev_indices, + dev_bufferView, + n, + indexAccessor.byteStride, + indexAccessor.byteOffset, + componentTypeByteSize); + + + checkCUDAError("Set Index Buffer"); + + + // ---------Primitive Info------- + + // Warning: LINE_STRIP is not supported in tinygltfloader + int numPrimitives; + PrimitiveType primitiveType; + switch (primitive.mode) { + case TINYGLTF_MODE_TRIANGLES: + primitiveType = PrimitiveType::Triangle; + numPrimitives = numIndices / 3; + break; + case TINYGLTF_MODE_TRIANGLE_STRIP: + primitiveType = PrimitiveType::Triangle; + numPrimitives = numIndices - 2; + break; + case TINYGLTF_MODE_TRIANGLE_FAN: + primitiveType = PrimitiveType::Triangle; + numPrimitives = numIndices - 2; + break; + case TINYGLTF_MODE_LINE: + primitiveType = PrimitiveType::Line; + numPrimitives = numIndices / 2; + break; + case TINYGLTF_MODE_LINE_LOOP: + primitiveType = PrimitiveType::Line; + numPrimitives = numIndices + 1; + break; + case TINYGLTF_MODE_POINTS: + primitiveType = PrimitiveType::Point; + numPrimitives = numIndices; + break; + default: + // output error + break; + }; + + + // ----------Attributes------------- + + auto it(primitive.attributes.begin()); + auto itEnd(primitive.attributes.end()); + + int numInstances = NUM_INSTANCES; + int numVertices = 0; + // for each attribute + for (; it != itEnd; it++) { + const tinygltf::Accessor &accessor = scene.accessors.at(it->second); + const tinygltf::BufferView &bufferView = scene.bufferViews.at(accessor.bufferView); + + int n = 1; + if (accessor.type == TINYGLTF_TYPE_SCALAR) { + n = 1; + } + else if (accessor.type == TINYGLTF_TYPE_VEC2) { + n = 2; + } + else if (accessor.type == TINYGLTF_TYPE_VEC3) { + n = 3; + } + else if (accessor.type == TINYGLTF_TYPE_VEC4) { + n = 4; + } + + BufferByte * dev_bufferView = bufferViewDevPointers.at(accessor.bufferView); + BufferByte ** dev_attribute = NULL; + + numVertices = accessor.count; + int componentTypeByteSize; + + // Note: since the type of our attribute array (dev_position) is static (float32) + // We assume the glTF model attribute type are 5126(FLOAT) here + + if (it->first.compare("POSITION") == 0) { + componentTypeByteSize = sizeof(VertexAttributePosition) / n; + dev_attribute = (BufferByte**)&dev_position; + } + else if (it->first.compare("NORMAL") == 0) { + componentTypeByteSize = sizeof(VertexAttributeNormal) / n; + dev_attribute = (BufferByte**)&dev_normal; + } + else if (it->first.compare("TEXCOORD_0") == 0) { + componentTypeByteSize = sizeof(VertexAttributeTexcoord) / n; + dev_attribute = (BufferByte**)&dev_texcoord0; + } + + std::cout << accessor.bufferView << " - " << it->second << " - " << it->first << '\n'; + + dim3 numThreadsPerBlock(128); + dim3 numBlocks((n * numVertices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); + int byteLength = numVertices * n * componentTypeByteSize; + cudaMalloc(dev_attribute, byteLength); + + _deviceBufferCopy << > > ( + n * numVertices, + *dev_attribute, + dev_bufferView, + n, + accessor.byteStride, + accessor.byteOffset, + componentTypeByteSize); + + std::string msg = "Set Attribute Buffer: " + it->first; + checkCUDAError(msg.c_str()); + } + + // malloc for VertexOut + VertexOut* dev_vertexOut; + cudaMalloc(&dev_vertexOut, numVertices * numInstances * sizeof(VertexOut)); + checkCUDAError("Malloc VertexOut Buffer"); + + // ----------Materials------------- + + // You can only worry about this part once you started to + // implement textures for your rasterizer + TextureData* dev_diffuseTex = NULL; + int diffuseTexWidth = 0; + int diffuseTexHeight = 0; + if (!primitive.material.empty()) { + const tinygltf::Material &mat = scene.materials.at(primitive.material); + printf("material.name = %s\n", mat.name.c_str()); + + if (mat.values.find("diffuse") != mat.values.end()) { + std::string diffuseTexName = mat.values.at("diffuse").string_value; + printf("Diffuse Texture name = %s\n", diffuseTexName.c_str()); + if (scene.textures.find(diffuseTexName) != scene.textures.end()) + { + const tinygltf::Texture &tex = scene.textures.at(diffuseTexName); + if (scene.images.find(tex.source) != scene.images.end()) { + const tinygltf::Image &image = scene.images.at(tex.source); + + size_t s = image.image.size() * sizeof(TextureData); + cudaMalloc(&dev_diffuseTex, s); + cudaMemcpy(dev_diffuseTex, &image.image.at(0), s, cudaMemcpyHostToDevice); + + diffuseTexWidth = image.width; + diffuseTexHeight = image.height; + + checkCUDAError("Set Texture Image data"); + } + } + + } + + // TODO: write your code for other materails + // You may have to take a look at tinygltfloader + // You can also use the above code loading diffuse material as a start point + } + + + // ---------Node hierarchy transform-------- + cudaDeviceSynchronize(); + + dim3 numBlocksNodeTransform((numVertices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); + _nodeMatrixTransform << > > ( + numVertices, + dev_position, + dev_normal, + matrix, + matrixNormal); + + checkCUDAError("Node hierarchy transformation"); + + // at the end of the for loop of primitive + // push dev pointers to map + primitiveVector.push_back(PrimitiveDevBufPointers{ + primitive.mode, + primitiveType, + numPrimitives, + numIndices, + numVertices, + numInstances, + + 0, + + dev_indices, + dev_position, + dev_normal, + dev_texcoord0, + + dev_diffuseTex, + diffuseTexWidth, + diffuseTexHeight, + + dev_vertexOut //VertexOut + }); + + totalNumPrimitives += (numPrimitives * numInstances); + + } // for each primitive + + } // for each mesh + } // for each node - // 2. for each mesh: - // for each primitive: - // build device buffer of indices, materail, and each attributes - // and store these pointers in a map - { - - std::map nodeString2Matrix; - auto rootNodeNamesList = scene.scenes.at(scene.defaultScene); - - { - auto it = rootNodeNamesList.begin(); - auto itEnd = rootNodeNamesList.end(); - for (; it != itEnd; ++it) { - traverseNode(nodeString2Matrix, scene, *it, glm::mat4(1.0f)); - } - } + } - // parse through node to access mesh - - auto itNode = nodeString2Matrix.begin(); - auto itEndNode = nodeString2Matrix.end(); - for (; itNode != itEndNode; ++itNode) { - - const tinygltf::Node & N = scene.nodes.at(itNode->first); - const glm::mat4 & matrix = itNode->second; - const glm::mat3 & matrixNormal = glm::transpose(glm::inverse(glm::mat3(matrix))); - - auto itMeshName = N.meshes.begin(); - auto itEndMeshName = N.meshes.end(); - - for (; itMeshName != itEndMeshName; ++itMeshName) { - - const tinygltf::Mesh & mesh = scene.meshes.at(*itMeshName); - - auto res = mesh2PrimitivesMap.insert(std::pair>(mesh.name, std::vector())); - std::vector & primitiveVector = (res.first)->second; - - // for each primitive - for (size_t i = 0; i < mesh.primitives.size(); i++) { - const tinygltf::Primitive &primitive = mesh.primitives[i]; - - if (primitive.indices.empty()) - return; - - // TODO: add new attributes for your PrimitiveDevBufPointers when you add new attributes - VertexIndex* dev_indices = NULL; - VertexAttributePosition* dev_position = NULL; - VertexAttributeNormal* dev_normal = NULL; - VertexAttributeTexcoord* dev_texcoord0 = NULL; - - // ----------Indices------------- - - const tinygltf::Accessor &indexAccessor = scene.accessors.at(primitive.indices); - const tinygltf::BufferView &bufferView = scene.bufferViews.at(indexAccessor.bufferView); - BufferByte* dev_bufferView = bufferViewDevPointers.at(indexAccessor.bufferView); - - // assume type is SCALAR for indices - int n = 1; - int numIndices = indexAccessor.count; - int componentTypeByteSize = sizeof(VertexIndex); - int byteLength = numIndices * n * componentTypeByteSize; - - dim3 numThreadsPerBlock(128); - dim3 numBlocks((numIndices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); - cudaMalloc(&dev_indices, byteLength); - _deviceBufferCopy << > > ( - numIndices, - (BufferByte*)dev_indices, - dev_bufferView, - n, - indexAccessor.byteStride, - indexAccessor.byteOffset, - componentTypeByteSize); - - - checkCUDAError("Set Index Buffer"); - - - // ---------Primitive Info------- - - // Warning: LINE_STRIP is not supported in tinygltfloader - int numPrimitives; - PrimitiveType primitiveType; - switch (primitive.mode) { - case TINYGLTF_MODE_TRIANGLES: - primitiveType = PrimitiveType::Triangle; - numPrimitives = numIndices / 3; - break; - case TINYGLTF_MODE_TRIANGLE_STRIP: - primitiveType = PrimitiveType::Triangle; - numPrimitives = numIndices - 2; - break; - case TINYGLTF_MODE_TRIANGLE_FAN: - primitiveType = PrimitiveType::Triangle; - numPrimitives = numIndices - 2; - break; - case TINYGLTF_MODE_LINE: - primitiveType = PrimitiveType::Line; - numPrimitives = numIndices / 2; - break; - case TINYGLTF_MODE_LINE_LOOP: - primitiveType = PrimitiveType::Line; - numPrimitives = numIndices + 1; - break; - case TINYGLTF_MODE_POINTS: - primitiveType = PrimitiveType::Point; - numPrimitives = numIndices; - break; - default: - // output error - break; - }; - - - // ----------Attributes------------- - - auto it(primitive.attributes.begin()); - auto itEnd(primitive.attributes.end()); - - int numVertices = 0; - // for each attribute - for (; it != itEnd; it++) { - const tinygltf::Accessor &accessor = scene.accessors.at(it->second); - const tinygltf::BufferView &bufferView = scene.bufferViews.at(accessor.bufferView); - - int n = 1; - if (accessor.type == TINYGLTF_TYPE_SCALAR) { - n = 1; - } - else if (accessor.type == TINYGLTF_TYPE_VEC2) { - n = 2; - } - else if (accessor.type == TINYGLTF_TYPE_VEC3) { - n = 3; - } - else if (accessor.type == TINYGLTF_TYPE_VEC4) { - n = 4; - } - - BufferByte * dev_bufferView = bufferViewDevPointers.at(accessor.bufferView); - BufferByte ** dev_attribute = NULL; - - numVertices = accessor.count; - int componentTypeByteSize; - - // Note: since the type of our attribute array (dev_position) is static (float32) - // We assume the glTF model attribute type are 5126(FLOAT) here - - if (it->first.compare("POSITION") == 0) { - componentTypeByteSize = sizeof(VertexAttributePosition) / n; - dev_attribute = (BufferByte**)&dev_position; - } - else if (it->first.compare("NORMAL") == 0) { - componentTypeByteSize = sizeof(VertexAttributeNormal) / n; - dev_attribute = (BufferByte**)&dev_normal; - } - else if (it->first.compare("TEXCOORD_0") == 0) { - componentTypeByteSize = sizeof(VertexAttributeTexcoord) / n; - dev_attribute = (BufferByte**)&dev_texcoord0; - } - - std::cout << accessor.bufferView << " - " << it->second << " - " << it->first << '\n'; - - dim3 numThreadsPerBlock(128); - dim3 numBlocks((n * numVertices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); - int byteLength = numVertices * n * componentTypeByteSize; - cudaMalloc(dev_attribute, byteLength); - - _deviceBufferCopy << > > ( - n * numVertices, - *dev_attribute, - dev_bufferView, - n, - accessor.byteStride, - accessor.byteOffset, - componentTypeByteSize); - - std::string msg = "Set Attribute Buffer: " + it->first; - checkCUDAError(msg.c_str()); - } - - // malloc for VertexOut - VertexOut* dev_vertexOut; - cudaMalloc(&dev_vertexOut, numVertices * sizeof(VertexOut)); - checkCUDAError("Malloc VertexOut Buffer"); - - // ----------Materials------------- - - // You can only worry about this part once you started to - // implement textures for your rasterizer - TextureData* dev_diffuseTex = NULL; - int diffuseTexWidth = 0; - int diffuseTexHeight = 0; - if (!primitive.material.empty()) { - const tinygltf::Material &mat = scene.materials.at(primitive.material); - printf("material.name = %s\n", mat.name.c_str()); - - if (mat.values.find("diffuse") != mat.values.end()) { - std::string diffuseTexName = mat.values.at("diffuse").string_value; - if (scene.textures.find(diffuseTexName) != scene.textures.end()) { - const tinygltf::Texture &tex = scene.textures.at(diffuseTexName); - if (scene.images.find(tex.source) != scene.images.end()) { - const tinygltf::Image &image = scene.images.at(tex.source); - - size_t s = image.image.size() * sizeof(TextureData); - cudaMalloc(&dev_diffuseTex, s); - cudaMemcpy(dev_diffuseTex, &image.image.at(0), s, cudaMemcpyHostToDevice); - - diffuseTexWidth = image.width; - diffuseTexHeight = image.height; - - checkCUDAError("Set Texture Image data"); - } - } - } - - // TODO: write your code for other materails - // You may have to take a look at tinygltfloader - // You can also use the above code loading diffuse material as a start point - } - - - // ---------Node hierarchy transform-------- - cudaDeviceSynchronize(); - - dim3 numBlocksNodeTransform((numVertices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); - _nodeMatrixTransform << > > ( - numVertices, - dev_position, - dev_normal, - matrix, - matrixNormal); - - checkCUDAError("Node hierarchy transformation"); - - // at the end of the for loop of primitive - // push dev pointers to map - primitiveVector.push_back(PrimitiveDevBufPointers{ - primitive.mode, - primitiveType, - numPrimitives, - numIndices, - numVertices, - - dev_indices, - dev_position, - dev_normal, - dev_texcoord0, - - dev_diffuseTex, - diffuseTexWidth, - diffuseTexHeight, - - dev_vertexOut //VertexOut - }); - - totalNumPrimitives += numPrimitives; - - } // for each primitive - - } // for each mesh - - } // for each node - - } - - - // 3. Malloc for dev_primitives - { - cudaMalloc(&dev_primitives, totalNumPrimitives * sizeof(Primitive)); - } - - - // Finally, cudaFree raw dev_bufferViews - { - - std::map::const_iterator it(bufferViewDevPointers.begin()); - std::map::const_iterator itEnd(bufferViewDevPointers.end()); - - //bufferViewDevPointers - - for (; it != itEnd; it++) { - cudaFree(it->second); - } - - checkCUDAError("Free BufferView Device Mem"); - } + // 3. Malloc for dev_primitives + { + cudaMalloc(&dev_primitives, totalNumPrimitives * sizeof(Primitive)); + } + + + // Finally, cudaFree raw dev_bufferViews + { + + std::map::const_iterator it(bufferViewDevPointers.begin()); + std::map::const_iterator itEnd(bufferViewDevPointers.end()); + //bufferViewDevPointers + for (; it != itEnd; it++) { + cudaFree(it->second); + } + + checkCUDAError("Free BufferView Device Mem"); + } } +//------------------------- SHADERS ----------------------------------------------------------------- + +static int curPrimitiveBeginId = 0; +__global__ +void _vertexTransformAndAssembly(PrimitiveDevBufPointers primitive, glm::mat4 MVP, glm::mat4 MV, glm::mat3 MV_normal, glm::mat3 M_InverseTranspose, glm::vec3 eyePos, int width, int height) +{ + // vertex id + const int vid = (blockIdx.x * blockDim.x) + threadIdx.x; + const int numVertices = primitive.numVertices; + const int numInstances = primitive.numInstances; + + if (vid < numVertices) + { + // Attributes + const glm::vec3 inPos = primitive.dev_position[vid]; + const glm::vec3 inNormal = primitive.dev_normal[vid]; + + // This is in NDC Space + glm::vec4 outPos = MVP * glm::vec4(inPos, 1.f); + outPos /= outPos.w; + + // Convert to screen Space + const glm::vec4 screenPos = NDCToScreenSpace(&outPos, width, height); + + // TODO : Change this + const float instanceOffset = 100.f; + + const int gridSize = 3;// glm::sqrt(numInstances); + + for (int instanceId = 0; instanceId < numInstances; ++instanceId) + { + const float offsetX = (instanceId % gridSize) * instanceOffset; + const float offsetY = (instanceId / gridSize) * instanceOffset; + + // Output of vertex shader + primitive.dev_verticesOut[vid * numInstances + instanceId].pos = screenPos + glm::vec4(offsetX, offsetY, 0.f, 0.f);// instanceOffset * float(instanceId); + primitive.dev_verticesOut[vid * numInstances + instanceId].eyePos = glm::vec3(MV * glm::vec4(inPos, 1.f)); + primitive.dev_verticesOut[vid * numInstances + instanceId].eyeNor = MV_normal * inNormal; + + if (primitive.dev_diffuseTex != NULL) + { + primitive.dev_verticesOut[vid * numInstances + instanceId].texcoord0 = primitive.dev_texcoord0[vid]; + } + else + { + primitive.dev_verticesOut[vid * numInstances + instanceId].color = glm::vec3(1.0f);//glm::clamp(glm::vec3(outPos), glm::vec3(0.f), glm::vec3(1.f)); + } + } + } +} -__global__ -void _vertexTransformAndAssembly( - int numVertices, - PrimitiveDevBufPointers primitive, - glm::mat4 MVP, glm::mat4 MV, glm::mat3 MV_normal, - int width, int height) { +__global__ +void _primitiveAssembly(int numIndices, int curPrimitiveBeginId, Primitive* dev_primitives, PrimitiveDevBufPointers primitive) { - // vertex id - int vid = (blockIdx.x * blockDim.x) + threadIdx.x; - if (vid < numVertices) { + // index id + const int iid = (blockIdx.x * blockDim.x) + threadIdx.x; + const int numInstances = primitive.numInstances; + + if (iid < numIndices) { + + // This is primitive assembly for triangles + + for (int instanceId = 0; instanceId < numInstances; ++instanceId) + { + int pid = iid / (int)primitive.primitiveType;; // id for cur primitives vector + if (primitive.primitiveMode == TINYGLTF_MODE_TRIANGLES) + { + const int devBufferIndexPos = primitive.dev_indices[iid] * numInstances + instanceId; + const int primitiveIndexPos = (pid + curPrimitiveBeginId) * numInstances + instanceId; + + dev_primitives[primitiveIndexPos].v[iid % (int)primitive.primitiveType] = primitive.dev_verticesOut[devBufferIndexPos]; + dev_primitives[primitiveIndexPos].primitiveType = primitive.primitiveType; + dev_primitives[primitiveIndexPos].instanceId = instanceId; + + if (primitive.dev_diffuseTex != NULL) + { + dev_primitives[primitiveIndexPos].dev_diffuseTex = primitive.dev_diffuseTex; + dev_primitives[primitiveIndexPos].diffuseTexWidth = primitive.diffuseTexWidth; + dev_primitives[primitiveIndexPos].diffuseTexHeight = primitive.diffuseTexHeight; + } + else + { + dev_primitives[primitiveIndexPos].dev_diffuseTex = NULL; + dev_primitives[primitiveIndexPos].diffuseTexWidth = 0; + dev_primitives[primitiveIndexPos].diffuseTexHeight = 0; + } + } + } + } +} - // TODO: Apply vertex transformation here - // Multiply the MVP matrix for each vertex position, this will transform everything into clipping space - // Then divide the pos by its w element to transform into NDC space - // Finally transform x and y to viewport space +__global__ +void _rasterizePrimitive(int width, int height, int totalNumPrimitives, Primitive* dev_primitives, Fragment* dev_fragmentBuffer, float* dev_depth, unsigned int* mutexLock) +{ + // primitive id + const int primitiveId = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (primitiveId < totalNumPrimitives) + { + Primitive primitive = dev_primitives[primitiveId]; + + PrimitiveType test = Triangle; + + if (test == Triangle) + { + // Vertices in screen Space + VertexOut v1 = primitive.v[0]; + VertexOut v2 = primitive.v[1]; + VertexOut v3 = primitive.v[2]; + + glm::vec3 triangle[3]; + triangle[0] = glm::vec3(v1.pos); + triangle[1] = glm::vec3(v2.pos); + triangle[2] = glm::vec3(v3.pos); + + TextureData* dev_diffuseTex = primitive.dev_diffuseTex; + int textureWidth = primitive.diffuseTexWidth; + int textureHeight = primitive.diffuseTexHeight; + + const AABB bounds = getAABBForTriangle(triangle); + + // Clamp To Screen + float minX = glm::clamp(bounds.min[0], 0.f, width - 1.f); + float maxX = glm::clamp(bounds.max[0], 0.f, width - 1.f); + float minY = glm::clamp(bounds.min[1], 0.f, height - 1.f); + float maxY = glm::clamp(bounds.max[1], 0.f, height - 1.f); + + for (int row = minY; row <= maxY; ++row) + { + for (int col = minX; col <= maxX; ++col) + { + const int pixelIndex = col + row * width; + const glm::vec2 currPos(col, row); + + // Calculate BaryCentric coordinates + const glm::vec3 baryCoord = calculateBarycentricCoordinate(triangle, currPos); + + // Check if point is inside triangle + const bool isInside = isBarycentricCoordInBounds(baryCoord); + + if (isInside) + { + // Get the interop depth + const float currDepth = -getZAtCoordinate(baryCoord, triangle); + + bool isSet; + do { + isSet = (atomicCAS(&mutexLock[pixelIndex], 0, 1) == 0); + if (isSet) + { + if (currDepth < dev_depth[pixelIndex]) + { + dev_fragmentBuffer[pixelIndex].eyeNor = (baryCoord.x * v1.eyeNor) + (baryCoord.y * v2.eyeNor) + (baryCoord.z * v3.eyeNor); + dev_fragmentBuffer[pixelIndex].eyePos = (baryCoord.x * v1.eyePos) + (baryCoord.y * v2.eyePos) + (baryCoord.z * v3.eyePos); + + if (dev_diffuseTex != NULL) + { + glm::vec2 textureCoord = baryCoord[0] * v1.texcoord0 + baryCoord[1] * v2.texcoord0 + baryCoord[2] * v3.texcoord0; + textureCoord = glm::vec2(textureCoord.x * textureWidth, textureCoord.y * textureHeight); + + const int u = int(textureCoord.x); + const int v = int(textureCoord.y); + +#ifdef BI_LINEAR_INTERPOLATION + dev_fragmentBuffer[pixelIndex].color = sampleTextureBiLinear(dev_diffuseTex, u, v, glm::fract(textureCoord), textureWidth, BYTES_PER_PIXEL); +#else + dev_fragmentBuffer[pixelIndex].color = sampleTextureSimple(dev_diffuseTex, u, v, textureWidth, BYTES_PER_PIXEL); +#endif + } + else + { + dev_fragmentBuffer[pixelIndex].color = (baryCoord.x * v1.color) + (baryCoord.y * v2.color) + (baryCoord.z * v3.color); + } + dev_depth[pixelIndex] = currDepth; + } + } + if (isSet) + { + mutexLock[pixelIndex] = 0; + } + } while (!isSet); + + } + } + } + } + else if (test == Line) + { + // Vertices in screen Space + VertexOut v1 = primitive.v[0]; + VertexOut v2 = primitive.v[1]; + VertexOut v3 = primitive.v[2]; + + glm::vec3 triangle[3]; + triangle[0] = glm::vec3(v1.pos); + triangle[1] = glm::vec3(v2.pos); + triangle[2] = glm::vec3(v3.pos); + + TextureData* dev_diffuseTex = primitive.dev_diffuseTex; + int textureWidth = primitive.diffuseTexWidth; + int textureHeight = primitive.diffuseTexHeight; + + const AABB bounds = getAABBForTriangle(triangle); + + // Clamp To Screen + float minX = glm::clamp(bounds.min[0], 0.f, width - 1.f); + float maxX = glm::clamp(bounds.max[0], 0.f, width - 1.f); + float minY = glm::clamp(bounds.min[1], 0.f, height - 1.f); + float maxY = glm::clamp(bounds.max[1], 0.f, height - 1.f); + + for (int row = minY; row <= maxY; ++row) + { + for (int col = minX; col <= maxX; ++col) + { + const int pixelIndex = col + row * width; + const glm::vec2 currPos(col, row); + + // Calculate BaryCentric coordinates + const glm::vec3 baryCoord = calculateBarycentricCoordinate(triangle, currPos); + + // Check if point is inside triangle + const bool isInside = isBarycentricCoordOnBoundary(baryCoord); + + if (isInside) + { + // Get the interop depth + const float currDepth = -getZAtCoordinate(baryCoord, triangle); + + bool isSet; + do { + isSet = (atomicCAS(&mutexLock[pixelIndex], 0, 1) == 0); + if (isSet) + { + if (currDepth < dev_depth[pixelIndex]) + { + dev_fragmentBuffer[pixelIndex].eyeNor = (baryCoord.x * v1.eyeNor) + (baryCoord.y * v2.eyeNor) + (baryCoord.z * v3.eyeNor); + dev_fragmentBuffer[pixelIndex].eyePos = (baryCoord.x * v1.eyePos) + (baryCoord.y * v2.eyePos) + (baryCoord.z * v3.eyePos); + + if (dev_diffuseTex != NULL) + { + glm::vec2 textureCoord = baryCoord[0] * v1.texcoord0 + baryCoord[1] * v2.texcoord0 + baryCoord[2] * v3.texcoord0; + textureCoord = glm::vec2(textureCoord.x * textureWidth, textureCoord.y * textureHeight); + + const int u = int(textureCoord.x); + const int v = int(textureCoord.y); + +#ifdef BI_LINEAR_INTERPOLATION + dev_fragmentBuffer[pixelIndex].color = sampleTextureBiLinear(dev_diffuseTex, u, v, glm::fract(textureCoord), textureWidth, BYTES_PER_PIXEL); +#else + dev_fragmentBuffer[pixelIndex].color = sampleTextureSimple(dev_diffuseTex, u, v, textureWidth, BYTES_PER_PIXEL); +#endif + } + else + { + dev_fragmentBuffer[pixelIndex].color = (baryCoord.x * v1.color) + (baryCoord.y * v2.color) + (baryCoord.z * v3.color); + } + dev_depth[pixelIndex] = currDepth; + } + } + if (isSet) + { + mutexLock[pixelIndex] = 0; + } + } while (!isSet); + + } + } + } + + } + else if (test == Point) + { + // Vertices in screen Space + VertexOut v1 = primitive.v[0]; + VertexOut v2 = primitive.v[1]; + VertexOut v3 = primitive.v[2]; + + glm::vec3 triangle[3]; + triangle[0] = glm::vec3(v1.pos); + triangle[1] = glm::vec3(v2.pos); + triangle[2] = glm::vec3(v3.pos); + + TextureData* dev_diffuseTex = primitive.dev_diffuseTex; + int textureWidth = primitive.diffuseTexWidth; + int textureHeight = primitive.diffuseTexHeight; + + const AABB bounds = getAABBForTriangle(triangle); + + // Clamp To Screen + float minX = glm::clamp(bounds.min[0], 0.f, width - 1.f); + float maxX = glm::clamp(bounds.max[0], 0.f, width - 1.f); + float minY = glm::clamp(bounds.min[1], 0.f, height - 1.f); + float maxY = glm::clamp(bounds.max[1], 0.f, height - 1.f); + + for (int row = minY; row <= maxY; ++row) + { + for (int col = minX; col <= maxX; ++col) + { + const int pixelIndex = col + row * width; + const glm::vec2 currPos(col, row); + + // Calculate BaryCentric coordinates + const glm::vec3 baryCoord = calculateBarycentricCoordinate(triangle, currPos); + + // Check if point is on the vertices of the triangle + const bool isInside = isBarycentricCoordOnVertices(triangle, currPos); + + if (isInside) + { + // Get the interop depth + const float currDepth = -getZAtCoordinate(baryCoord, triangle); + + bool isSet; + do { + isSet = (atomicCAS(&mutexLock[pixelIndex], 0, 1) == 0); + if (isSet) + { + if (currDepth < dev_depth[pixelIndex]) + { + dev_fragmentBuffer[pixelIndex].eyeNor = (baryCoord.x * v1.eyeNor) + (baryCoord.y * v2.eyeNor) + (baryCoord.z * v3.eyeNor); + dev_fragmentBuffer[pixelIndex].eyePos = (baryCoord.x * v1.eyePos) + (baryCoord.y * v2.eyePos) + (baryCoord.z * v3.eyePos); + + if (dev_diffuseTex != NULL) + { + glm::vec2 textureCoord = baryCoord[0] * v1.texcoord0 + baryCoord[1] * v2.texcoord0 + baryCoord[2] * v3.texcoord0; + textureCoord = glm::vec2(textureCoord.x * textureWidth, textureCoord.y * textureHeight); + + const int u = int(textureCoord.x); + const int v = int(textureCoord.y); + +#ifdef BI_LINEAR_INTERPOLATION + dev_fragmentBuffer[pixelIndex].color = sampleTextureBiLinear(dev_diffuseTex, u, v, glm::fract(textureCoord), textureWidth, BYTES_PER_PIXEL); +#else + dev_fragmentBuffer[pixelIndex].color = sampleTextureSimple(dev_diffuseTex, u, v, textureWidth, BYTES_PER_PIXEL); +#endif + } + else + { + dev_fragmentBuffer[pixelIndex].color = (baryCoord.x * v1.color) + (baryCoord.y * v2.color) + (baryCoord.z * v3.color); + } + dev_depth[pixelIndex] = currDepth; + } + } + if (isSet) + { + mutexLock[pixelIndex] = 0; + } + } while (!isSet); + + } + } + } + + } + } +} - // TODO: Apply vertex assembly here - // Assemble all attribute arraies into the primitive array - - } +__global__ +void _rasterizePrimitiveWithMSAA(int width, int height, int totalNumPrimitives, Primitive* dev_primitives, Fragment* dev_fragmentBuffer, float* dev_depth, unsigned int* mutexLock) +{ + // primitive id + /*const int primitiveId = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (primitiveId < totalNumPrimitives) + { + Primitive primitive = dev_primitives[primitiveId]; + + if (primitive.primitiveType == Triangle) + { + // Vertices in screen Space + VertexOut v1 = primitive.v[0]; + VertexOut v2 = primitive.v[1]; + VertexOut v3 = primitive.v[2]; + + glm::vec3 triangle[3]; + triangle[0] = glm::vec3(v1.pos); + triangle[1] = glm::vec3(v2.pos); + triangle[2] = glm::vec3(v3.pos); + + TextureData* dev_diffuseTex = primitive.dev_diffuseTex; + int textureWidth = primitive.diffuseTexWidth; + int textureHeight = primitive.diffuseTexHeight; + + const AABB bounds = getAABBForTriangle(triangle); + + // Clamp To Screen + float minX = glm::clamp(bounds.min[0], 0.f, width - 1.f); + float maxX = glm::clamp(bounds.max[0], 0.f, width - 1.f); + float minY = glm::clamp(bounds.min[1], 0.f, height - 1.f); + float maxY = glm::clamp(bounds.max[1], 0.f, height - 1.f); + + for (int row = minY; row <= maxY; row += MSAA_LEVEL) + { + for (int col = minX; col <= maxX; col += MSAA_LEVEL) + { + + + bool isInside = false; + + const int centerPixelIndex = (col + MSAA_LEVEL / 2) + (row * MSAA_LEVEL / 2) * width; + const glm::vec2 centerPixelPos((col + MSAA_LEVEL / 2), (row * MSAA_LEVEL / 2)); + const glm::vec3 centerBaryCoord = calculateBarycentricCoordinate(triangle, centerPixelPos); + + for (int subSampleX = 0; subSampleX < MSAA_LEVEL; ++subSampleX) + { + for (int subSampleY = 0; subSampleY < MSAA_LEVEL; ++subSampleY) + { + + const glm::vec2 currPosSubSample(col + subSampleX + subSampleY, row); + // Calculate BaryCentric coordinates + subSampleBarCoord[subSampleX][subSampleY] = calculateBarycentricCoordinate(triangle, currPosSubSample); + + // Check if point is inside triangle + const bool subSampleHit = isBarycentricCoordInBounds(subSampleBarCoord[subSampleX][subSampleY]); + subSubSampleHit[subSampleX][subSampleY] = subSampleHit; + if (subSampleHit) + { + isInside = true; + } + } + } + + if (isInside) + { + + // 1. Sample the center pixel + glm::vec3 centerPixelColor(0.f); + + bool isSet; + do { + isSet = (atomicCAS(&mutexLock[centerPixelIndex], 0, 1) == 0); + if (isSet) + { + if (currDepth < dev_depth[centerPixelIndex]) + { + dev_fragmentBuffer[centerPixelIndex].eyeNor = (centerBaryCoord.x * v1.eyeNor) + (centerBaryCoord.y * v2.eyeNor) + (centerBaryCoord.z * v3.eyeNor); + dev_fragmentBuffer[centerPixelIndex].eyePos = (centerBaryCoord.x * v1.eyePos) + (centerBaryCoord.y * v2.eyePos) + (centerBaryCoord.z * v3.eyePos); + + if (dev_diffuseTex != NULL) + { + glm::vec2 textureCoord = (centerBaryCoord.x * v1.texcoord0) + (centerBaryCoord.y * v2.texcoord0) + (centerBaryCoord.z * v3.texcoord0); + textureCoord = glm::vec2(textureCoord.x * textureWidth, textureCoord.y * textureHeight); + + textureCoord = glm::clamp(textureCoord, glm::vec2(0.f), glm::vec2(textureWidth - 1, textureHeight - 1)); + + // Apparently there are 3 bytes per pixel based on the texture array size and texture size. + const int startPixelIndex = int(textureCoord.x + textureCoord.y * textureWidth) * 3; + float r = dev_diffuseTex[startPixelIndex]; + float g = dev_diffuseTex[startPixelIndex + 1]; + float b = dev_diffuseTex[startPixelIndex + 2]; + centerPixelColor = glm::vec3(r, g, b) / 255.f; + dev_fragmentBuffer[pixelIndex].color = centerPixelColor; + } + else + { + centerPixelColor = glm::vec3(1.0f);// (baryCoord.x * v1.color) + (baryCoord.y * v2.color) + (baryCoord.z * v3.color); + dev_fragmentBuffer[centerPixelIndex].color = centerPixelColor; + } + dev_depth[centerPixelIndex] = currDepth; + } + } + if (isSet) + { + mutexLock[centerPixelIndex] = 0; + } + } while (!isSet); + + // 2. Do a depth test for all the remaining subSamples and update framebuffer + + + for (int subSampleX = 0; subSampleX < MSAA_LEVEL; ++subSampleX) + { + for (int subSampleY = 0; subSampleY < MSAA_LEVEL; ++subSampleY) + { + const glm::vec2 currPosSubSample(col + subSampleX + subSampleY, row); + + } + } + } + } + } + } + }*/ } +//#define MAT_PLAIN +#define MAT_LAMBERT +//#define MAT_BLINN_PHONG -static int curPrimitiveBeginId = 0; +#define BLINN_PHONEXP 64.f +#define AMBIENT_LIGHT 0.2f -__global__ -void _primitiveAssembly(int numIndices, int curPrimitiveBeginId, Primitive* dev_primitives, PrimitiveDevBufPointers primitive) { +/** +* Writes fragment colors to the framebuffer +*/ +__global__ +void render(int w, int h, Fragment *fragmentBuffer, glm::vec3 *framebuffer) +{ + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * w); + + if (x < w && y < h) + { + + // No Shading +#ifdef MAT_PLAIN + framebuffer[index] = fragmentBuffer[index].color; +#endif - // index id - int iid = (blockIdx.x * blockDim.x) + threadIdx.x; + // Lamberts Material +#ifdef MAT_LAMBERT - if (iid < numIndices) { + float lambertsTerm = glm::abs(glm::dot(fragmentBuffer[index].eyeNor, glm::vec3(0, 0, 1))) + AMBIENT_LIGHT; + lambertsTerm = glm::clamp(lambertsTerm, 0.f, 1.f); + framebuffer[index] = fragmentBuffer[index].color * lambertsTerm; +#endif - // TODO: uncomment the following code for a start - // This is primitive assembly for triangles +#ifdef MAT_BLINN_PHONG - //int pid; // id for cur primitives vector - //if (primitive.primitiveMode == TINYGLTF_MODE_TRIANGLES) { - // pid = iid / (int)primitive.primitiveType; - // dev_primitives[pid + curPrimitiveBeginId].v[iid % (int)primitive.primitiveType] - // = primitive.dev_verticesOut[primitive.dev_indices[iid]]; - //} + const glm::vec3 fsNorm = fragmentBuffer[index].eyeNor; + const glm::vec3 fsCamera = glm::vec3(0, 0, 1.f); + const glm::vec3 hVec = (fsNorm + fsCamera) / 2.f; + const float specular = glm::max(glm::pow(glm::dot(glm::normalize(hVec), glm::normalize(fsNorm)), BLINN_PHONEXP), 0.f); + const float lambertsTerm = glm::clamp(glm::abs(glm::dot(fsNorm, fsCamera)) + AMBIENT_LIGHT, 0.f, 1.f); - // TODO: other primitive types (point, line) - } - + framebuffer[index] = fragmentBuffer[index].color * (lambertsTerm + specular); +#endif + + } } +#define SSAA_UNIFORM_GRID + +/** +* Performs Anti-Aliasing +*/ +__global__ +void _SSAA(int downWidth, int downHeight, int originalWidth, int originalHeight, glm::vec3 *inputFrameBuffer, glm::vec3 *aaFrameBuffer) +{ + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * downWidth); + + if (x < downWidth && y < downHeight) + { +#ifdef SSAA_UNIFORM_GRID + + const int upscaledIndexX = x * SSAA_LEVEL; + const int upscaledIndexY = y * SSAA_LEVEL; + + glm::vec3 averageColor(0.f); + int numColors = 0; + + for (int i = 0; i < SSAA_LEVEL; ++i) + { + for (int j = 0; j < SSAA_LEVEL; ++j) + { + const int newIndex = (upscaledIndexX + i) + ((upscaledIndexY + i) * originalWidth); + averageColor += inputFrameBuffer[newIndex]; + numColors++; + } + } + aaFrameBuffer[index] = (averageColor / float(numColors)); + +#endif // SSAA_UNIFORM_GRID + + } +} /** * Perform rasterization. */ -void rasterize(uchar4 *pbo, const glm::mat4 & MVP, const glm::mat4 & MV, const glm::mat3 MV_normal) { +void rasterize(uchar4 *pbo, const glm::mat4 & MVP, const glm::mat4 & MV, const glm::mat3& MV_normal, const glm::mat3& M_inverseTranspose, const glm::vec3& eyePos) { + int sideLength2d = 8; dim3 blockSize2d(sideLength2d, sideLength2d); - dim3 blockCount2d((width - 1) / blockSize2d.x + 1, - (height - 1) / blockSize2d.y + 1); - - // Execute your rasterization pipeline here - // (See README for rasterization pipeline outline.) - - // Vertex Process & primitive assembly - { - curPrimitiveBeginId = 0; - dim3 numThreadsPerBlock(128); - - auto it = mesh2PrimitivesMap.begin(); - auto itEnd = mesh2PrimitivesMap.end(); - - for (; it != itEnd; ++it) { - auto p = (it->second).begin(); // each primitive - auto pEnd = (it->second).end(); - for (; p != pEnd; ++p) { - dim3 numBlocksForVertices((p->numVertices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); - dim3 numBlocksForIndices((p->numIndices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); - - _vertexTransformAndAssembly << < numBlocksForVertices, numThreadsPerBlock >> >(p->numVertices, *p, MVP, MV, MV_normal, width, height); - checkCUDAError("Vertex Processing"); - cudaDeviceSynchronize(); - _primitiveAssembly << < numBlocksForIndices, numThreadsPerBlock >> > - (p->numIndices, - curPrimitiveBeginId, - dev_primitives, - *p); - checkCUDAError("Primitive Assembly"); - - curPrimitiveBeginId += p->numPrimitives; - } - } - - checkCUDAError("Vertex Processing and Primitive Assembly"); - } - - cudaMemset(dev_fragmentBuffer, 0, width * height * sizeof(Fragment)); - initDepth << > >(width, height, dev_depth); - - // TODO: rasterize - - - - // Copy depthbuffer colors into framebuffer - render << > >(width, height, dev_fragmentBuffer, dev_framebuffer); - checkCUDAError("fragment shader"); + + dim3 blockCount2d((width - 1) / blockSize2d.x + 1, + (height - 1) / blockSize2d.y + 1); + + dim3 blockAACount2d((originalWidth - 1) / blockSize2d.x + 1, + (originalHeight - 1) / blockSize2d.y + 1); + + dim3 numThreadsPerBlock(128); + dim3 blocksPrimitives((totalNumPrimitives + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); + + // Execute your rasterization pipeline here + // (See README for rasterization pipeline outline.) + + // Vertex Process & primitive assembly + { + curPrimitiveBeginId = 0; + + auto it = mesh2PrimitivesMap.begin(); + auto itEnd = mesh2PrimitivesMap.end(); + + for (; it != itEnd; ++it) + { + auto p = (it->second).begin(); + auto pEnd = (it->second).end(); + + for (; p != pEnd; ++p) + { + dim3 numBlocksForVertices((p->numVertices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); + dim3 numBlocksForIndices((p->numIndices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x); + + // 1. Vertex Assembly and Shader + _vertexTransformAndAssembly << < numBlocksForVertices, numThreadsPerBlock >> > (*p, MVP, MV, MV_normal, M_inverseTranspose, eyePos, width, height); + checkCUDAError("Vertex Processing"); + + cudaDeviceSynchronize(); + + // 2. Primitive Assembly + _primitiveAssembly << < numBlocksForIndices, numThreadsPerBlock >> > + (p->numIndices, + curPrimitiveBeginId, + dev_primitives, + *p); + checkCUDAError("Primitive Assembly"); + + curPrimitiveBeginId += (p->numPrimitives * p->numInstances); + } + } + + checkCUDAError("Vertex Processing and Primitive Assembly"); + } + + cudaMemset(dev_fragmentBuffer, 0, width * height * sizeof(Fragment)); + + // 3. Depth Check + initDepth << > > (width, height, dev_depth); + + // 4. Rasterize - Call per primitive +#ifdef AA_MULTI_SAMPLE + _rasterizePrimitive << > > (width, height, totalNumPrimitives, dev_primitives, dev_fragmentBuffer, dev_depth, dev_mutex); +#else + _rasterizePrimitive << > > (width, height, totalNumPrimitives, dev_primitives, dev_fragmentBuffer, dev_depth, dev_mutex); +#endif + checkCUDAError("Rasterizer"); + + // Copy fragmentBuffer colors into framebuffer + render << > > (width, height, dev_fragmentBuffer, dev_framebuffer); + checkCUDAError("fragment shader"); + + // perform SSAA + _SSAA << < blockAACount2d, blockSize2d >> > (originalWidth, originalHeight, width, height, dev_framebuffer, dev_AAFrameBuffer); + checkCUDAError("SSAA"); + // Copy framebuffer into OpenGL buffer for OpenGL previewing - sendImageToPBO<<>>(pbo, width, height, dev_framebuffer); + sendImageToPBO << > > (pbo, originalWidth, originalHeight, dev_AAFrameBuffer); checkCUDAError("copy render result to pbo"); } @@ -741,36 +1329,43 @@ void rasterizeFree() { // deconstruct primitives attribute/indices device buffer - auto it(mesh2PrimitivesMap.begin()); - auto itEnd(mesh2PrimitivesMap.end()); - for (; it != itEnd; ++it) { - for (auto p = it->second.begin(); p != it->second.end(); ++p) { - cudaFree(p->dev_indices); - cudaFree(p->dev_position); - cudaFree(p->dev_normal); - cudaFree(p->dev_texcoord0); - cudaFree(p->dev_diffuseTex); + auto it(mesh2PrimitivesMap.begin()); + auto itEnd(mesh2PrimitivesMap.end()); + for (; it != itEnd; ++it) { + for (auto p = it->second.begin(); p != it->second.end(); ++p) { + cudaFree(p->dev_indices); + cudaFree(p->dev_position); + cudaFree(p->dev_normal); + cudaFree(p->dev_texcoord0); + cudaFree(p->dev_diffuseTex); + + cudaFree(p->dev_verticesOut); - cudaFree(p->dev_verticesOut); - - //TODO: release other attributes and materials - } - } + //TODO: release other attributes and materials + } + } - //////////// + //////////// cudaFree(dev_primitives); dev_primitives = NULL; - cudaFree(dev_fragmentBuffer); - dev_fragmentBuffer = NULL; + cudaFree(dev_fragmentBuffer); + dev_fragmentBuffer = NULL; cudaFree(dev_framebuffer); dev_framebuffer = NULL; - cudaFree(dev_depth); - dev_depth = NULL; + cudaFree(dev_AAFrameBuffer); + dev_AAFrameBuffer = NULL; + + cudaFree(dev_depth); + dev_depth = NULL; + + cudaFree(dev_mutex); + dev_mutex = NULL; + checkCUDAError("rasterize Free"); } diff --git a/src/rasterize.h b/src/rasterize.h index 560aae9..a82bbfb 100644 --- a/src/rasterize.h +++ b/src/rasterize.h @@ -20,5 +20,5 @@ namespace tinygltf{ void rasterizeInit(int width, int height); void rasterizeSetBuffers(const tinygltf::Scene & scene); -void rasterize(uchar4 *pbo, const glm::mat4 & MVP, const glm::mat4 & MV, const glm::mat3 MV_normal); +void rasterize(uchar4 *pbo, const glm::mat4 & MVP, const glm::mat4 & MV, const glm::mat3& MV_normal, const glm::mat3& M_inverseTranspose, const glm::vec3& eyePos); void rasterizeFree(); diff --git a/src/rasterizeTools.h b/src/rasterizeTools.h index 46c701e..b86a13c 100644 --- a/src/rasterizeTools.h +++ b/src/rasterizeTools.h @@ -11,6 +11,7 @@ #include #include #include +#include struct AABB { glm::vec3 min; @@ -25,6 +26,16 @@ glm::vec3 multiplyMV(glm::mat4 m, glm::vec4 v) { return glm::vec3(m * v); } +__host__ __device__ static +glm::vec4 NDCToScreenSpace(const glm::vec4* v, int width, int height) +{ + glm::vec4 screenCoords(0.f, 0.f, 0.f, 1.f); + screenCoords[0] = ((*v)[0] + 1.f) * width / 2.f; + screenCoords[1] = (1.f - (*v)[1]) * height / 2.f; + screenCoords[2] = ((*v)[2] + 1.f) / 2.f;; + return screenCoords; +} + // CHECKITOUT /** * Finds the axis aligned bounding box for a given triangle. @@ -33,13 +44,13 @@ __host__ __device__ static AABB getAABBForTriangle(const glm::vec3 tri[3]) { AABB aabb; aabb.min = glm::vec3( - min(min(tri[0].x, tri[1].x), tri[2].x), - min(min(tri[0].y, tri[1].y), tri[2].y), - min(min(tri[0].z, tri[1].z), tri[2].z)); + glm::min(glm::min(tri[0].x, tri[1].x), tri[2].x), + glm::min(glm::min(tri[0].y, tri[1].y), tri[2].y), + glm::min(glm::min(tri[0].z, tri[1].z), tri[2].z)); aabb.max = glm::vec3( - max(max(tri[0].x, tri[1].x), tri[2].x), - max(max(tri[0].y, tri[1].y), tri[2].y), - max(max(tri[0].z, tri[1].z), tri[2].z)); + glm::max(glm::max(tri[0].x, tri[1].x), tri[2].x), + glm::max(glm::max(tri[0].y, tri[1].y), tri[2].y), + glm::max(glm::max(tri[0].z, tri[1].z), tri[2].z)); return aabb; } @@ -88,6 +99,32 @@ bool isBarycentricCoordInBounds(const glm::vec3 barycentricCoord) { barycentricCoord.z >= 0.0 && barycentricCoord.z <= 1.0; } +#define POINT_TOLERANCE 2.5f +#define LINE_TOLERANCE 0.02f + +#define IS_SAME(a,b) abs(a - b) < POINT_TOLERANCE + +/** +* Check if a barycentric coordinate is on the boundary of a triangle. +*/ +__host__ __device__ static +bool isBarycentricCoordOnBoundary(const glm::vec3 barycentricCoord) +{ + return isBarycentricCoordInBounds(barycentricCoord) && (barycentricCoord.x <= LINE_TOLERANCE || barycentricCoord.y <= LINE_TOLERANCE || + barycentricCoord.z <= LINE_TOLERANCE); +} + +/** +* Check if a barycentric coordinate is points of a triangle. +*/ +__host__ __device__ static +bool isBarycentricCoordOnVertices(const glm::vec3 tri[3], glm::vec2 point) +{ + return (IS_SAME(point.x, tri[0].x) && IS_SAME(point.y, tri[0].y)) || + (IS_SAME(point.x, tri[1].x) && IS_SAME(point.y, tri[1].y)) || + (IS_SAME(point.x, tri[2].x) && IS_SAME(point.y, tri[2].y)); +} + // CHECKITOUT /** * For a given barycentric coordinate, compute the corresponding z position @@ -99,3 +136,131 @@ float getZAtCoordinate(const glm::vec3 barycentricCoord, const glm::vec3 tri[3]) + barycentricCoord.y * tri[1].z + barycentricCoord.z * tri[2].z); } + + +/** +* Sample a given texture at the given UV using no filtering +*/ +__host__ __device__ static +glm::vec3 sampleTextureSimple(const unsigned char* textureData, int u, int v, int textureWidth, int bytesPerPixel) +{ + const int startPixelIndex = int(u + v * textureWidth) * bytesPerPixel; + return glm::vec3(textureData[startPixelIndex], textureData[startPixelIndex + 1], textureData[startPixelIndex + 2]); +} + +/** +* Sample a given texture at the given UV using bilinear filtering +*/ +__host__ __device__ static +glm::vec3 sampleTextureBiLinear(const unsigned char* textureData, int u, int v, const glm::vec2& mixRatio, int textureWidth, int bytesPerPixel) +{ + int sampleIndex = (u + v * textureWidth) * bytesPerPixel; + const glm::vec3 sample1 = glm::vec3(textureData[sampleIndex], textureData[sampleIndex + 1], textureData[sampleIndex + 2]); + + sampleIndex = (u + 1 + v * textureWidth) * bytesPerPixel; + const glm::vec3 sample2 = glm::vec3(textureData[sampleIndex], textureData[sampleIndex + 1], textureData[sampleIndex + 2]); + + sampleIndex = (u + (v + 1) * textureWidth) * bytesPerPixel; + const glm::vec3 sample3 = glm::vec3(textureData[sampleIndex], textureData[sampleIndex + 1], textureData[sampleIndex + 2]); + + sampleIndex = (u + 1 + (v + 1) * textureWidth) * bytesPerPixel; + const glm::vec3 sample4 = glm::vec3(textureData[sampleIndex], textureData[sampleIndex + 1], textureData[sampleIndex + 2]); + + const glm::vec3 mixInX = glm::mix(sample2, sample4, mixRatio.x); + const glm::vec3 mixInY = glm::mix(sample1, sample3, mixRatio.x); + + return glm::mix(mixInX, mixInY, mixRatio.y) / 255.f; +} + + +/** +* This class is used for timing the performance +* Uncopyable and unmovable +* +* Adapted from WindyDarian(https://github.com/WindyDarian) +*/ +class PerformanceTimer +{ +public: + PerformanceTimer() + { + cudaEventCreate(&event_start); + cudaEventCreate(&event_end); + } + + ~PerformanceTimer() + { + cudaEventDestroy(event_start); + cudaEventDestroy(event_end); + } + + void startCpuTimer() + { + if (cpu_timer_started) { throw std::runtime_error("CPU timer already started"); } + cpu_timer_started = true; + + time_start_cpu = std::chrono::high_resolution_clock::now(); + } + + void endCpuTimer() + { + time_end_cpu = std::chrono::high_resolution_clock::now(); + + if (!cpu_timer_started) { throw std::runtime_error("CPU timer not started"); } + + std::chrono::duration duro = time_end_cpu - time_start_cpu; + prev_elapsed_time_cpu_milliseconds = + static_cast(duro.count()); + + cpu_timer_started = false; + } + + void startGpuTimer() + { + if (gpu_timer_started) { throw std::runtime_error("GPU timer already started"); } + gpu_timer_started = true; + + cudaEventRecord(event_start); + } + + void endGpuTimer() + { + cudaEventRecord(event_end); + cudaEventSynchronize(event_end); + + if (!gpu_timer_started) { throw std::runtime_error("GPU timer not started"); } + + cudaEventElapsedTime(&prev_elapsed_time_gpu_milliseconds, event_start, event_end); + gpu_timer_started = false; + } + + float getCpuElapsedTimeForPreviousOperation() //noexcept //(damn I need VS 2015 + { + return prev_elapsed_time_cpu_milliseconds; + } + + float getGpuElapsedTimeForPreviousOperation() //noexcept + { + return prev_elapsed_time_gpu_milliseconds; + } + + // remove copy and move functions + PerformanceTimer(const PerformanceTimer&) = delete; + PerformanceTimer(PerformanceTimer&&) = delete; + PerformanceTimer& operator=(const PerformanceTimer&) = delete; + PerformanceTimer& operator=(PerformanceTimer&&) = delete; + +private: + cudaEvent_t event_start = nullptr; + cudaEvent_t event_end = nullptr; + + using time_point_t = std::chrono::high_resolution_clock::time_point; + time_point_t time_start_cpu; + time_point_t time_end_cpu; + + bool cpu_timer_started = false; + bool gpu_timer_started = false; + + float prev_elapsed_time_cpu_milliseconds = 0.f; + float prev_elapsed_time_gpu_milliseconds = 0.f; +};