CUDA是GPU通用计算的一种，其中现在大热的深度学习底层GPU计算差不多都选择的CUDA，在这我们先简单了解下其中的一些概念，为了好理解，我们先用DX11里的Compute shader来和CUDA比较下，这二者都可用于GPU通用计算。
先上一张微软MSDN上的图.
Compute shader:
线程块: Dispatch(x,y,z), 索引SV_GroupID
线程组: [numthreads(SIZE_X, SIZE_Y, 1)], 索引SV_GroupThreadID.
组内索引: CS以组为单位，shader共享在一个组内，groupshared / GroupMemoryBarrierWithGroupSync,其SV_GroupIndex为组内索引，组内共享块一般用此SV_GroupIndex做为索引，或是这个的倍数，SV_GroupIndex = SV_GroupThreadID.x + SV_GroupThreadID.y*SIZE_X(在这假定二维)
所有线程唯一索引:在整个空间的索引三维索引为SV_DispatchThreadID，SV_DispatchThreadID = SV_GroupThreadID+SV_GroupID*numthreads；
如果提供一个width,height的数据，有关系width=x*SIZE_X，height=y*SIZE_Y.(所以一般我们得到数据的长宽，然后设定线程组后，调度就直接求出来了，但是可能不是整除，所以可以把真实的传入进去).而SV_DispatchThreadID表示在整个width,height中的索引，一般来说，我们直接用SV_DispatchThreadID就够了，但是如果使用了groupshared/GroupMemoryBarrierWithGroupSync，就会用SV_GroupIndex来整个当个线程组计算。
同样的概念CUDA中：increment_kernel<<<gridDim, blockDim, 0, 0>>>
线程块: gridDim, 索引 blockIdx
线程组: blockDim 索引 threadIdx
组内索引：给组内共享块索引用.__shared__/__syncthreads，那针对的对象应该用groupIndex来当索引。
int groupIndex = threadIdx.x；(假定一维）
int groupIndex = threadIdx.x + threadIdx.y*blockDim.x;(假定二维)
所有线程唯一索引: 在线程组里的索引 threadIdx ，和dx cs不同，这里是三维的。
如上找到在整个width,height中的位置和上面的SV_DispatchThreadID一样。
const int idx = threadIdx.x + blockIdx.x * blockDim.x; const int idy = threadIdx.y + blockIdx.y * blockDim.y;
同理dx11里常用内存显存交换API如map/unmap对应cudaMemcpyAsync cudaMemcpyDeviceToHost/cudaMemcpyHostToDevice这些。
回到正题上，导向滤波算法 是何凯明提出的一种保边滤波算法，这种算法在这不细说，opencv已经自带这种算法，这篇文章只是针对这种算法实现一个cuda版本的，毕竟对于图像处理来说，很多点都是并行计算的，这个链接里有提到matlab的实现，我们看下多通道下的快速实现来看下效果。
效果非常牛，其中我把相应matlab实现贴出来，先简单理解下，其中matlab的这种写法很像一些深度学习库python接口里提供的数据操作类的用法，看起来还是很好理解的。