University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 3
- Trung Le
- Windows 10 Home, i7-4790 CPU @ 3.60GHz 12GB, GTX 980 Ti (Person desktop)
- [Perspective correct UV texture mapping with bilinear filtering](## UV Texture Mapping)
- [Order independent transparency using k-buffer](## Order independent transparency using k-buffer)
The following header flags can be found in
#define RENDER_DEPTH_ONLY: Uncomment to render with z-buffer only
#define RENDER_NORMAL_ONLY: Uncomment to render with eye space normals only
#define BILINEAR_FILTERING: Uncomment to enable bilinear filtering
#define NOT_USE_TEXTURE: Uncomment if the model doesn't have a texture file (cow, centaur, 2 cylinder engine, wolf, octocat, flower)
#define USE_K_BUFFER: Uncomment if using independent order transparency with k-buffer
#define SHARED_MEMORY_MATERIALS: Uncomment to use shared memory for materials
(not working, please ignore)
The following header flags can be found in
#define USE_CENTAUR_MODEL: Uncomment if using the
#define USE_HEAD_MODEL: Uncomment if using the
#define USE_ENGINE_MODEL: Uncomment if using the
#define USE_TURNTABLE: Uncomment to rotate model on a turntable
UV Texture Mapping
1. Perspective correct
An attribute of a fragment inside a triangle can be computed by using barycentric interpolation of the triangle vertices. A naive approach for this interpolation is directly interpolate the attribute in raster space after the vertices have been projected onto the screen. However, since the projected vertices don't preserve the correct perspective, or depth, of the fragment computed, we have to perform perspective correction. I implemented perspective correct UV texture mapping and depth interpolation using the explanation at scratchapixel.
2. Bilinear filtering
Bilinear filtering is an antialiasing technique that averages the nearest four texels of the lookup texel to create smooth effect when the texture is zoomed in or zoomed out.
Below shows the comparison with bilinear filtering on (left) and off (right). Bilinear filtering makes the wavy curve on the Cesium logo appear to be smoother and not pixelated.
|Bilinear ON||Bilinear OFF|
At first, I expected bilinear filtering might have a performace hit on the
_rasterize function because of the nearest texels lookup. However, after profiling, it appears to only affect rasterization minimally. See the [performance analysis](# Performance analysis) section.
Order independent transparency using k-buffer
k-buffer is a generalized version of the traditional z-buffer for depth. Instead of storing a single depth as in the z-buffer, k-buffer stores a list of depths for overlapping fragments. This information allows us to compute the alpha blending color of these multiple overlapping fragments in screenspace.
|k-buffer ON||k-buffer OFF|
To make sure that colors from overlapping fragments are accumulated correctly,
atomicAdd is used to synchronize across all threads that attempt to accumulate to the
All profiling was done using NSight Performance Analysis tool with the following parameters
Kernel execution time
| Scene used
The bottleneck happens in the
_rasterize kernel because we have to loop through each pixel in every triangle's bounding box. Therefore, each
_rasterize kernel is bounded by O(n2), where n is the size of the triangle's bounding box in screen space. This means that a large triangle with a large bounding box will have a performance hit. To compare, I profiled a scene with the head model where the camera is located at the origin, versus a scene where the camera is zoomed in.
A close up at
|Camera at origin||Camera zoomed in|
As we can see, the
_rasterize kernel increases significantly in kernel time because each triangle that it has to rasterize now has a larger area with more number of fragments.
Kernel execution time vs. features
Profile is done by enabling each feature one by one
As expected, bilinear filtering and k-buffer occupy more device time. However, the performace decrease isn't significant enough. For the k-buffer, instead of using a linked list of depth buffers, I only created an additional buffer of accumulated alpha colors of overlapping fragments. This optimized for having to look several depth buffer, which could make memory read and write from global buffer slower.
Registers used vs. kernel
_rasterize and (interestingly)
_vertexTransformAndAssembly take up a large number of registers, 112 and 95, respectively. This significantly severes the ability for GPU scheduler to optimize the number of active blocks per kernel launch.
Shared memory for materials
To optimize for
render kernel global memory read from materials list, this data is copied over to shared memory. In the beginning of the
render kernel call, threads within the same block will copy over the materials data from globabal memory to shared memory. A CUDA
__syncthreads() is used to make sure this shared data is initialized properly before used.
The following uses three different scene with varying number of materials and ran for 20,000ms each:
|1 material||5 materials||115 materials|
Let's take a look at the kernel analysis for
Looks quite bad! It seems that having sharing memory isn't that great at all. The occupancy for active warps are quite low. This is due to the fact that an additional step is required to transfer the materials data over from global so shared memory with a thread sync.
Backface culling is implemented, but without stream compaction. I only include here the reference renders for this feature but it is incomlete as to gain proper performace increase.
|Flower with backface culling||Flower with backface culling|
|4212 tris||18342 tris||15708 tris|
|Centaur||Cesium Milk Truck||flower|
|34670 tris||3624 tris||640 tris|
|5804 tris||17684 tris||121496 tris|
Renders of normal and depth attributes:
Bonus artistic renders!
int conversion goes wrong...
When your ghost friend won't stop staring at you :-)
General information for CUDA device
- Device name: GeForce GTX 980 Ti
- Compute capability: 5.2
- Compute mode: Default
- Clock rate: 1076000
- Integrated: 0
- Device copy overlap: Enabled
- Kernel execution timeout: Enabled
Memory information for CUDA device
- Total global memory: 6442450944
- Total constant memory: 65536
- Multiprocessor count: 22
- Shared memory per multiprocessor: 98304
- Registers per multiprocessor: 65536
- Max threads per multiprocessor: 2048
- Max grid dimensions: [2147483647, 65535, 65535]
- Max threads per block: 1024
- Max registers per block: 65536
- Max thread dimensions: [1024, 1024, 64]
- Threads per block: 512
- tinygltfloader by @soyoyo
- obj2gltf by AnalyticalGraphicsInc
- Implementing Weighted, Blended Order-Independent Transparency by Morgan McGuire
- k-buffer by Morgan McGuire
- glTF Sample Models
- octocat by Sally Kong. Model is converted using obj2gltf.
- wolf by Rachel Hwang. Model is converted using obj2gltf.
- centaur model by BAQStudio, Model is converted using obj2gltf.
- Infinite, 3D Head Scan by Lee Perry-Smith is licensed under a Creative Commons Attribution 3.0 Unported License. Based on a work at www.triplegangers.com. This distribution was created by Morgan McGuire and Guedis Cardenas http://graphics.cs.williams.edu/data/. See LICENSE. Model is converted using obj2gltf.