工作的处理变量固定尺寸组在CUDA

我发现自己想要写类似于CUDA代码：工作的处理变量固定尺寸组在CUDA

__constant__ int cumulativeGroupThreadCount[GroupCount]; 

__device__ int SearchForGroupIndex() 
{ 
    int gridThreadID = blockIdx.x*blockDim.x + threadIdx.x; 
    int groupIndex = 0; 
    int sum = 0; 


    while(groupIndex < GroupCount && gridThreadID <= sum) 
    { 
     sum += cumulativeGroupThreadCount[groupIndex]; 
     ++groupIndex; 
    } 

    return groupIndex; 
} 

__device__ KernelMain() 
{ 
    GroupData data = groupData[SearchForGroupIndex()]; 

    ... 
}

随着意图来处理数据，其中每个基团可以是具有不同的尺寸的基团，但我希望使用一个CUDA线程处理每个组项目。每个组都有一组特定的关联组数据。

如果组数不是那么大，那么我认为它可能不是最差的方法？

随着组数的增加，我可以开始考虑更复杂的warp/block wide二进制搜索类型的行为，但我觉得在那一点上它不是一个好主意。

另一种方法可能是按大小对每个组进行排序/存储，然后使用相同的大小分别处理每个组。或者是在每个组中最多的项目被采用的变体，导致可能多个无用的线程需要被屏蔽掉。可能分成多批以减少浪费。

更简单的方法当然是将索引存储到每个线程的组数据，但这可能会浪费更多的空间和内存带宽，然后需要？

那么有没有更好的方法来处理这种类型的问题，一般在CUDA中？

来源

2016-12-01 iam

二分查找应该可以正常工作。它将具有良好的内存局部性，因为相邻的线程将位于相同或相邻的组中，并且如果平均组大小相对于warp大小较大，则分支散度将最小。

事情是这样的：

template<typename T, typename I> 
__device__ 
I upper_bound_index(T const* data, 
        I  count, 
        T const& value) { 
    I start = 0; 
    while(count > 0) { 
     I step = count/2; 
     if(!(value < data[start + step])) { 
      start += step + 1; 
      count -= step + 1; 
     } else { 
      count = step; 
     } 
    } 
    return start; 
} 

__global__ 
void group_kernel(int       numGroups, 
        int  const* __restrict__ cumulativeGroupThreadCount, 
        GroupData const* __restrict__ groupData) { 
    int gridThreadID = blockIdx.x*blockDim.x + threadIdx.x; 
    int groupID = upper_bound_index(cumulativeGroupThreadCount, 
            numGroups, 
            gridThreadID); 
    if(groupID == numGroups) { 
     // Excess threads 
     return; 
    } 
    int itemID = gridThreadID - (groupID > 0 ? 
           cumulativeGroupThreadCount[groupID-1] : 
           0); 
    GroupData data = groupData[groupID]; 
    // ... 
}

来源

2016-12-04 21:02:03

工作的处理变量固定尺寸组在CUDA

回答

相关问题