关于核函数中的循环和判断

各位[名词2]好，最近在研究一个GPU加速的例子，本身是一个查找表，要对这个查找表进行并行，大概的结构是一个（N*M）的矩阵，N行都各自独立，且分配一个线程，核函数中是对每一行的数据进行处理。但是现在有个问题是，运行一次核函数，里面需要循环，每一行的结构大概是这样（1+4*n），1代表是要循环多少次（times），然后每次一次的计算是涉及到4个数据，当然这里的n是固定的，并且这一行不一定填满，主要还是看这一行最开头的数据是多少。另外每进行一次循环利用这四个数据的时候，都会涉及到if判断。最终我发现，这样子的核函数运行似乎很慢，我的理解是因为循环和判断过多。（当然我的查找表是直接存放在全局内存，不涉及内存加速访问的效果），想请教一下，我该如何才能加速这个例子。下面是我的代码框架，其中涉及的一些参数我没有给出，大致的意思就是核函数包含一个循环以及一些判断，是否能够通过改进内存的调用或者改进核函数的写法来提高速度？？？    拜谢！！！
__global__ void kernel()
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int i = tid / mappingWidth;
    int j = tid - i * mappingWidth;
    if (i > 0 && i < LStitching_height - 1 && j>0 && j < LStitching_width - 1)
   {
      int numpic=*(lastpano_gpu+(i*mappingWidth+j)*(1+4*+0);

      for(int k=0;k<numpic;k++)
      {
         int nowpic=*(lastpano_gpu+(i*mappingWidth+j)*(1+4*+4*k+1);
         
         
         srcLocationX = *(lastpano_gpu+(i*mappingWidth+j)*(1+4*+4*k+2);
         srcLocationY = *(lastpano_gpu+(i*mappingWidth+j)*(1+4*+4*k+3);
         alpha=*(lastpano_gpu+(i*mappingWidth+j)*(1+4*+4*k+4);
                        if(nowpic==0)
                        {
                            .....
                        }
                        if(nowpic==1)
                        {
                            .....
                        }
                        if(nowpic==2)
                        {
                            .....
                        }
                        if(nowpic==3)
                        {
                            .....
                        }
                        if(nowpic==4)
                        {
                            .....
                        }
                        if(nowpic==5)
                        {
                            .....
                        }
                        if(nowpic==6)
                        {
                            .....
                        }
                        if(nowpic==7)
                        {
                            .....
                        }
}

各位[名词2]好，最近在研究一个GPU加速的例子，本身是一个查找表，要对这个查找表进行并行，大概的结构是一个（N*M）的矩阵，N行都各自独立，且分配一个线程，核函数中是对每一行的数据进行处理。但是现在有个问题是，运行一次核函数，里面需要循环，每一行的结构大概是这样（1+4*n），1代表是要循环多少次（times），然后每次一次的计算是涉及到4个数据，当然这里的n是固定的，并且这一行不一定填满，主要还是看这一行最开头的数据是多少。另外每进行一次循环利用这四个数据的时候，都会涉及到if判断。最终我发现，这样子的核函数运行似乎很慢，我的理解是因为循环和判断过多。（当然我的查找表是直接存放在全局内存，不涉及内存加速访问的效果），想请教一下，我该如何才能加速这个例子。下面是我的代码框架，其中涉及的一些参数我没有给出，大致的意思就是核函数包含一个循环以及一些判断，是否能够通过改进内存的调用或者改进核函数的写法来提高速度？？？    拜谢！！！
__global__ void kernel()
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int i = tid / mappingWidth;
    int j = tid - i * mappingWidth;
    if (i > 0 && i < LStitching_height - 1 && j>0 && j < LStitching_width - 1)
   {
      int numpic=*(lastpano_gpu+(i*mappingWidth+j)*(1+4*+0);

      for(int k=0;k<numpic;k++)
      {
         int nowpic=*(lastpano_gpu+(i*mappingWidth+j)*(1+4*+4*k+1);
         
         
         srcLocationX = *(lastpano_gpu+(i*mappingWidth+j)*(1+4*+4*k+2);
         srcLocationY = *(lastpano_gpu+(i*mappingWidth+j)*(1+4*+4*k+3);
         alpha=*(lastpano_gpu+(i*mappingWidth+j)*(1+4*+4*k+4);
                        if(nowpic==0)
                        {
                            .....
                        }
                        if(nowpic==1)
                        {
                            .....
                        }
                        if(nowpic==2)
                        {
                            .....
                        }
                        if(nowpic==3)
                        {
                            .....
                        }
                        if(nowpic==4)
                        {
                            .....
                        }
                        if(nowpic==5)
                        {
                            .....
                        }
                        if(nowpic==6)
                        {
                            .....
                        }
                        if(nowpic==7)
                        {
                            .....
                        }
}

因为楼主你的代码只给出了一个示范的骨架，我只能针对当前的骨架进行大体上的估计（以及，"kernel运行较慢"是一个主观上的概念）：

(1)首先楼主你做了复杂的1D到2D的变换（用来判断X和Y方向是否超出边界），最终访存的时候，依然使用了再从2D到1D回来的，有点绕。 这没有必要。CUDA里可以直接写成2D的block和grid的形状启动即可。无需反复变换过去变换过来的，清晰又易懂。

(2)根据你的if判断边界是否越界（如果你是这个意思）后，里面的代码看，每行的每个4x32-bit的元素，似乎可以在单独并行（如果你简化掉了重要的前后依赖关系，则当我没说）。但是因为开头你存放了1个4B元素，导致了所有的后续的4x4B都被错开了对齐的位置，导致实际上每次只能硬生生的读取4次，这个不太好。建议直接将行首的指示性的那个4B，单独存放一个数组，或者想法让开头的这个4B存放在上一行的行尾（你得添加一些padding类的东西），从而能让每行行首对齐到16B的边界，这样可以直接用int4类型一次性读完，而且多个线程间合并.

(3)以上都是小问题，大问题在于被你节选掉的部分：线程线程都要判断该4元组的第一个分量，然后走8种可能的不同的代码分支。这里可能是最影响性能的所在（只是可能，被你节选掉了）。

我如果是你，我会考虑只解决问题3. 而问题3如何解决，有很多种做法。常见的一种做法是在shared memory内部开辟8个缓冲区/数组，每行的相互有关的线程们，将数据从global中读取一行，直接分类到这8个数组中（同时维持对应的原始位置的索引，方便你回写结果缓冲区---当然，这里你也节选掉了）。然后在shared中分类完成中，直接每次集体的做每种。这样可以尽量的减少warp内部的分支，提高效率（这需要一定长度的行才能比较有效率）。

所以问题最终变成了，我如何能在shared中理顺他们。这个不妨自行搜索一下，如果每行的多个处理目标之间完全无关的话，简单的原子操作追加位置，再存放到每个类型的缓冲区中是最快的。类似一种直方图统计的变种，将需要处理的8种类型的，直接存放在8种bucket中，然后再集体的处理每个bucket里面的任务即可。

因为楼主你的代码只给出了一个示范的骨架，我只能针对当前的骨架进行大体上的估计（以及，"kernel运行较慢"是一个主观上的概念）：

(1)首先楼主你做了复杂的1D到2D的变换（用来判断X和Y方向是否超出边界），最终访存的时候，依然使用了再从2D到1D回来的，有点绕。 这没有必要。CUDA里可以直接写成2D的block和grid的形状启动即可。无需反复变换过去变换过来的，清晰又易懂。

(2)根据你的if判断边界是否越界（如果你是这个意思）后，里面的代码看，每行的每个4x32-bit的元素，似乎可以在单独并行（如果你简化掉了重要的前后依赖关系，则当我没说）。但是因为开头你存放了1个4B元素，导致了所有的后续的4x4B都被错开了对齐的位置，导致实际上每次只能硬生生的读取4次，这个不太好。建议直接将行首的指示性的那个4B，单独存放一个数组，或者想法让开头的这个4B存放在上一行的行尾（你得添加一些padding类的东西），从而能让每行行首对齐到16B的边界，这样可以直接用int4类型一次性读完，而且多个线程间合并.

(3)以上都是小问题，大问题在于被你节选掉的部分：线程线程都要判断该4元组的第一个分量，然后走8种可能的不同的代码分支。这里可能是最影响性能的所在（只是可能，被你节选掉了）。

我如果是你，我会考虑只解决问题3. 而问题3如何解决，有很多种做法。常见的一种做法是在shared memory内部开辟8个缓冲区/数组，每行的相互有关的线程们，将数据从global中读取一行，直接分类到这8个数组中（同时维持对应的原始位置的索引，方便你回写结果缓冲区---当然，这里你也节选掉了）。然后在shared中分类完成中，直接每次集体的做每种。这样可以尽量的减少warp内部的分支，提高效率（这需要一定长度的行才能比较有效率）。

所以问题最终变成了，我如何能在shared中理顺他们。这个不妨自行搜索一下，如果每行的多个处理目标之间完全无关的话，简单的原子操作追加位置，再存放到每个类型的缓冲区中是最快的。类似一种直方图统计的变种，将需要处理的8种类型的，直接存放在8种bucket中，然后再集体的处理每个bucket里面的任务即可。

此外，如同我上面发的，没有必要1个线程逐步处理一行（这样不仅仅4x4B不能一次性读取，warp的每个线程之间还在大跨步的跳跃式读取，没有必要）。如果每个分支超级小，则这里也应当考虑。

当然，如果你直接上问题3， 多个线程集体将行上的多个4元组给分类了，每个线程只逐步处理一行的问题也自动不存在了（因为你将集体的处理每个bucket中的问题了）。

此外，如同我上面发的，没有必要1个线程逐步处理一行（这样不仅仅4x4B不能一次性读取，warp的每个线程之间还在大跨步的跳跃式读取，没有必要）。如果每个分支超级小，则这里也应当考虑。

当然，如果你直接上问题3， 多个线程集体将行上的多个4元组给分类了，每个线程只逐步处理一行的问题也自动不存在了（因为你将集体的处理每个bucket中的问题了）。

感谢您的回复，我是小白，大致理解了一下您的意思，您看一下对不对。您的意思是处理一行的数据时，可以直接进行分类，我这是8个分类，分成8个类别后，在核函数里面对这8个类别直接都进行计算，然后返回值就行。这样能够避免循环和多次判断的问题，应该是这样理解的吧？
这里我需要解释一下，因为这8个类别不是每一行我都需要进行处理，所以我才进行循环numpic次，然后对每次循环的nowpic进行判断来自哪个类别然后进行计算。所以问题回归到您说的直接就计算8个类别，我不是很清楚这样子是否能够更快？
此外，我还有一点优化方向就是，其实我实际工作中不可能有8个类别，最多可能也就4个类别，然后一行数据中大部分都只需要计算1次，少部分计算2次，最少的部分计算3次或者4次。面对这种情况，不知道您是否有更好的建议，谢谢！

此外，如同我上面发的，没有必要1个线程逐步处理一行（这样不仅仅4x4B不能一次性读取，warp的每个线程之间还在大跨步的跳跃式读取，没有必要）。如果每个分支超级小，则这里也应当考虑。

当然，如果你直接上问题3， 多个线程集体将行上的多个4元组给分类了，每个线程只逐步处理一行的问题也自动不存在了（因为你将集体的处理每个bucket中的问题了）。

还有一个我需要解释给您的是，我每一行这8个分类排序其实都不是固定的，如果按照您的意思创建共享内存对每一行都预先分类和排序，那么我其实还需要创建一个数组专门用来预先保存好这每一行的中类别各自的排序是吗？然后我对8个类别进行计算的时候，就是去访问共享内存来知道我当前计算的类别排第几？我又仔细的读了一下您的回复，您的意思是我可以在共享内存中直接集体处理某一类的数据，是的，我的每个类型之间是无关的，至于您说的追加原子操作我尚未了解，我需要再进一步学习。关于共享内存的处理我还没能很好的学习，我再看看，再次感谢！

感谢您的回复，我是小白，大致理解了一下您的意思，您看一下对不对。您的意思是处理一行的数据时，可以直接进行分类，我这是8个分类，分成8个类别后，在核函数里面对这8个类别直接都进行计算，然后返回值就行。这样能够避免循环和多次判断的问题，应该是这样理解的吧？
这里我需要解释一下，因为这8个类别不是每一行我都需要进行处理，所以我才进行循环numpic次，然后对每次循环的nowpic进行判断来自哪个类别然后进行计算。所以问题回归到您说的直接就计算8个类别，我不是很清楚这样子是否能够更快？
此外，我还有一点优化方向就是，其实我实际工作中不可能有8个类别，最多可能也就4个类别，然后一行数据中大部分都只需要计算1次，少部分计算2次，最少的部分计算3次或者4次。面对这种情况，不知道您是否有更好的建议，谢谢！

我的意思主要有两点：
（1）代码只能看到一个骨架，最多只能看出多种计算之间是独立的，而不能知道每种计算究竟是否很简单/很复杂。如果计算很复杂的话，则可以在一定范围内分类（不一定是排序，分类即可），这样每种分类可以warp集中的计算，减少warp内分支，从而规避你担忧的因为divergent branch而导致的性能下降。

（2）你可以可以不分类整理到shared中的数组中。可以直接快速每次每个warp对一定范围内的数据（例如一个动态的，根据numpic，和行数之类的，所决定的范围），重复扫描多次，每次只处理一种。例如一个256个目标区域，第一次扫描需要处理的类型1的，warp能攒到～32个目标就集体计算一次。然后再扫描类型2，类似3的。这种不需要分配额外的分配的bucket的空间，但可能需要重复的读取原始数据多次。

（3）如果实际的问题和示范的代码不同，实际上只有有限很少量的几种，甚至一定范围内的，能落入分类1的就肯定不走范围2之类的，也可以考虑直接上，说不定更快。

提供更详细的代码有助于更进一步的分析。

我的意思主要有两点：
（1）代码只能看到一个骨架，最多只能看出多种计算之间是独立的，而不能知道每种计算究竟是否很简单/很复杂。如果计算很复杂的话，则可以在一定范围内分类（不一定是排序，分类即可），这样每种分类可以warp集中的计算，减少warp内分支，从而规避你担忧的因为divergent branch而导致的性能下降。

（2）你可以可以不分类整理到shared中的数组中。可以直接快速每次每个warp对一定范围内的数据（例如一个动态的，根据numpic，和行数之类的，所决定的范围），重复扫描多次，每次只处理一种。例如一个256个目标区域，第一次扫描需要处理的类型1的，warp能攒到～32个目标就集体计算一次。然后再扫描类型2，类似3的。这种不需要分配额外的分配的bucket的空间，但可能需要重复的读取原始数据多次。

（3）如果实际的问题和示范的代码不同，实际上只有有限很少量的几种，甚至一定范围内的，能落入分类1的就肯定不走范围2之类的，也可以考虑直接上，说不定更快。

提供更详细的代码有助于更进一步的分析。

此外，根据你的中文（而不是你的代码，代码并不能看出这点），你的问题实际上是一个行宽不定的数组，例如：
**
*
***
**
***
*
*
**
...

其中每个*代表一个处理任务，则等于横向有1-4个不定长度，纵向也是动态的这样的一个任务数组。

而这个动态的2D表格中的任务有多种类型。所以直接每个线程黑上一个*, 必然导致warp中的32个线程在多个类别中分支，降低效率。此时不如将一个2D的区域整体考虑，无论你是扫描多次，还是用shared中分类，这样使得每个warp能集体的处理1种类型的任务，从而尽量消除了warp内的分支，则可能有助于提升效率（得看每个分类的计算复杂度，上面已经说过了，别每种计算就一行代码，还不如你分类或者扫描的代价高，就不如直接计算了）。