（轉）基於CUDA的GPU光線追蹤
作者：Asixa
鏈接：https://zhuanlan.zhihu.com/p/55855479
來源：知乎
著作權歸作者所有。商業轉載請聯系作者獲得授權，非商業轉載請註明出處。
? ?
替STL。
4. 顯存層級
GPU上的顯存共分為三個層級，Global Memory, Shared Memory, Local Memory.
讀取耗時方面，L存 < S存 <<G存 << 內存。
local memory是最快的，但是需要註意一個問題，每個Kernel的local memory大
小是650000字節，如果使用量超過了這個量，就會崩潰。
5.在Device慎用 new，malloc
這兩個操作是在Device端創建一個Global Memory，這個弊端是很慢。我相信每個
使用GPU加速的程序都是對效率敏感的。
6. 異常處理
可以在每次調用完Kernel寫
auto error = cudaGetLastError();
if(error!=0)printf("error %d\n", error);
來檢測有沒有Error拋出。（我每次使用Nsight調試整個電腦就會崩，可能是我自己
的問題）
而在我寫渲染器的時候出現最多的是Error 77"內存越界"，一般的內存越界很容易避
免，但是我依然遇到很多很迷的崩潰然後拋出Error77，據我猜測應該包含但不限於以下
兩種情況：
- kernel棧溢出
前面說過，Kernel的棧深度並不夠用，第一種解決辦法是消除遞歸，減少函數相互調用
等。第二種是 將項目從Debug模式改成Release模式，這樣編譯器的優化就會發揮作用。
- Local Memory超過了極限,
將不需要的對象及時的free掉，或者使用cudaDeviceSetLimitAPI設置最低Local Memory
大小。
7. 隨機數
在之前的項目中我的隨機數使用的是drand48()，但是CUDA提供了一個更高效的隨機數
生成器curand。
curand提供多種隨機數序列。我用的最多的是最普通的curand_uniform，在我的光線追
蹤采樣中，我確保每個像素的采樣序列都不一樣，不然就會出現很多奇怪的效果
我為每個像素都創建了一個currandState
//Host
#include <curand_kernel.h>
//...
curandState *d_rng_states = nullptr;
cudaMalloc(reinterpret_cast<void **>(&d_rng_states), height * width * sizeof(curandState));
而種子方面，使用像素的唯一id。
//Device
const auto tidx = blockIdx.x * blockDim.x + threadIdx.x;
const auto tidy = blockIdx.y * blockDim.y + threadIdx.y;
curand_init(seed + tidx + tidy * d_width, 0, 0, &rngStates[tidx]);
這樣在每次調用
curand_uniform(&rngStates[tid]) //tid = tidx + tidy * width
就可以生成一個0~1的隨機浮點數了。
7. 紋理
在CPU渲染器中我使用byte[] 儲存的紋理信息，如果在Cuda中也使用 unsigned char* 的話，
會消耗很多的Global Memory，並且最重要的是，Global Memory很慢。
幸運的是Cuda提供了一個Texture解決方案，這個Texture儲存在一個特定的顯存區域可以
極大地提高讀取速度。
在Cuda的示例 0_Simple/simpleTexture項目中，項目實現了一個簡單Texture，這個Texture
通過綁定到了一部分顯存提供更快的讀取。甚至不需要傳遞指針到kernal即可當全局變量使用。
但是有兩個問題：
第一個問題，這個Texture不能是數組或者指針數組。也就是說Texture的數量在編譯的時候
就是寫死的。
解決方案：1. 將所有的紋理都合並到一張Atlas，這理論上是最快的，效果大概是這樣:
圖自Unity Form by gary_bbgames
第二個方案是使用Texture的BindlessTexture功能，這個在CUDA的示例 2_Graphics/bindle
ssTexture項目中有實現。而我采用的就是這種方法。
CudaTexture第二個問題是如何綁定RGB三通道，示例項目中的顏色通道只有一個，並且值類型
是float，我嘗試使用uchar3類型來儲存三個RGB值但是沒有成功。我最後使用的是LayeredTe
xture來創建三個層,代碼在Cuda示例 0_Simple/simpleLayeredTexture項目。我不確定這是否
是創建三通道紋理的最優方法，如果有其他寫法，請讓我知道謝謝。
三通道紋理的緩沖有點奇怪，是這樣的，在創建之前需要修改一下。
//類型float
RRRRRRRRRRGGGGGGGGGGBBBBBBBBBB
下面附Texture相關代碼
//Host
inline void InitTextureList()
????????{
????????????????for (auto i = 0; i < TEXTURE_COUNT; i++) {
????????????????????????//讀取紋理，使用了stb_image庫
????????????????????????int width, height, depth;
????????????????????????const auto tex_data = stbi_load(imageFilenames[i],&width, &height, &depth, 0);
????????????????????????const auto size = width * height * depth;
????????????????????????float* h_data = new float[size];
????????????????????????printf("LoadTexture %d,%d,%d\n", width, height, depth);
????????????????????????for (unsigned int layer = 0; layer < 3; layer++)
????????????????????????????????for (auto i = 0; i < static_cast<int>(width * height); i++)h_data[layer*width*height + i] = tex_data[i * 3 + layer] / 255.0;
//cudaArray Descriptor
????????????????????????cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
????????????????????????//cuda Array
????????????????????????cudaArray *d_cuArr;
????????????????????????cudaMalloc3DArray(&d_cuArr, &channelDesc, make_cudaExtent(width, height, 3), cudaArrayLayered);
cudaMemcpy3DParms myparms = { 0 };
????????????????????????myparms.srcPos = make_cudaPos(0, 0, 0);
????????????????????????myparms.dstPos = make_cudaPos(0, 0, 0);
????????????????????????myparms.srcPtr = make_cudaPitchedPtr(h_data, width * sizeof(float), width, height);
????????????????????????myparms.dstArray = d_cuArr;
????????????????????????myparms.extent = make_cudaExtent(width, height, 3);
????????????????????????myparms.kind = cudaMemcpyHostToDevice;
????????????????????????cudaMemcpy3D(&myparms);
????????????????????????
cudaResourceDesc texRes;
????????????????????????memset(&texRes, 0, sizeof(cudaResourceDesc));
????????????????????????texRes.resType = cudaResourceTypeArray;
????????????????????????texRes.res.array.array = d_cuArr;
????????????????????????cudaTextureDesc texDescr;
????????????????????????memset(&texDescr, 0, sizeof(cudaTextureDesc));
????????????????????????texDescr.filterMode = cudaFilterModeLinear;
????????????????????????texDescr.addressMode[0] = cudaAddressModeWrap; // clamp
????????????????????????texDescr.addressMode[1] = cudaAddressModeWrap;
????????????????????????texDescr.addressMode[2] = cudaAddressModeWrap;
????????????????????????texDescr.readMode = cudaReadModeElementType;
????????????????????????texDescr.normalizedCoords = true;
????????????????????????cudaCreateTextureObject(&textlist[i], &texRes, &texDescr, NULL);
????????????????}
????????}
? ?
//Device
const auto albedo =Vec3(
????????????????tex2DLayered<float>(texs[texid], rec.u, 1-rec.v, 0), //R
????????????????tex2DLayered<float>(texs[texid], rec.u, 1 - rec.v, 1),//G
????????????????tex2DLayered<float>(texs[texid], rec.u, 1 - rec.v, 2));//B
? ?
8. BVH層次包圍盒
在Kernel寫BVH真的是刺激....
首先正如前面所說，BVH必須在CPU創建，所以從Host向Device復制數據時候，需要復制
一棵二叉樹，二叉樹的子節點還是個派生類的指針.....
由於我之前沒單獨學過C語言的內存管理，所以這部分消耗了我整整兩天一夜的精力。
我最後的解決方案是將所有對象包括BVH節點放在一個父類指針數組(Hitable**)中先傳到
Device。每個對象都被賦予一個id，也就是在數組中的位置。而BVH樹的左右節點只是個int
對象。
二分查找部分，由於這部分原始代碼高度依賴於遞歸，需要改成循環。這部分我參考了
https://devblogs.nvidia.com/thinking-parallel-part-i-collision-detection-gpu/?devblogs.nvidia.com
Thinking Parallel, Part II: Tree Traversal on the GPU | NVIDIA Developer Blog?devblogs.nvidia.com
Thinking Parallel, Part III: Tree Construction on the GPU | NVIDIA Developer Blog?devblogs.nvidia.com
其中在第II部分，Minimizing Divergence 部分中的traverseIterative函數中。我創建的是
int stack[64];
並且這部分在每個像素的最初始被創建，每次查找時只是重設為0，最後記得free掉這個數組。
目前調試BVH依然有問題，渲染個茶壺是沒有問題的，
但是換成Bunny就會拋出Error 77。目前還沒有解決。
? ?
代碼目前開源在：
Asixa/ALightGPU?github.com
由於之前沒怎麽寫過C++項目，代碼可能有些亂，深表歉意，明天開學，等過一陣子可能才開
始修BUG和整理代碼。
關於為什麽我為什麽全都寫在頭文件裏，因為CUDA的編譯器如果想要代碼分離的話需要開啟
【generate relocatable device code】但是這樣會導致編譯器無法進行代碼優化。似乎另一種
解決方式是使用CUDA的*.cuh和*.cu文件進行代碼分離，但是我目前還沒有測試成功。如果這
樣可以的話之後整理代碼的時候會進行代碼分離。
（轉）基於CUDA的GPU光線追蹤