作者：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 Blogdevblogs.nvidia.com
其中在第II部分，Minimizing Divergence 部分中的traverseIterative函数中。我创建的是
int stack[64];
并且这部分在每个像素的最初始被创建，每次查找时只是重设为0，最后记得free掉这个数组。
目前调试BVH依然有问题，渲染个茶壶是没有问题的，
但是换成Bunny就会抛出Error 77。目前还没有解决。
代码目前开源在：
由于之前没怎么写过C++项目，代码可能有些乱，深表歉意，明天开学，等过一阵子可能才开
始修BUG和整理代码。
关于为什么我为什么全都写在头文件里，因为CUDA的编译器如果想要代码分离的话需要开启
【generate relocatable device code】但是这样会导致编译器无法进行代码优化。似乎另一种
解决方式是使用CUDA的*.cuh和*.cu文件进行代码分离，但是我目前还没有测试成功。如果这
样可以的话之后整理代码的时候会进行代码分离。