0

我正在处理与图遍历（维特比算法）相关的一些任务每个时间步骤我都有一组紧凑的活动状态，每个状态都完成了一些工作，然后结果通过传出弧传播到每个弧的目标状态等等建立了新的活动状态集。问题是输出弧的数量变化很大，从两个或三个到几千个。所以计算线程的加载效率很低。

我尝试通过共享本地内存队列共享作业

int tx = threaIdx.x;

extern __shared__ int smem[];

int *stateSet_s = smem;                     //new active set
int *arcSet_s = &(smem[Q_LEN]);             //local shared queue
float *scores_s = (float*)&(smem[2*Q_LEN]);

__shared__ int arcCnt;
__shared__ int stateCnt;

if ( tx == 0 )
{
   arcCnt = 0;
   stateCnt = 0;
}

__syncthreads();

//load state index from compacted list of state indexes
int stateId = activeSetIn_g[gtx];

float srcCost = scores_g[ stateId ];
int startId = outputArcStartIds_g[stateId];

int nArcs = outputArcCounts_g[stateId]; //number of outgoing arcs to be propagated (2-3 to thousands)

/////////////////////////////////////////////
/// prepare arc set
/// !!!! that is the troubled code I think !!!!
/// bank conflicts? uncoalesced access?

int myPos = atomicAdd ( &arcCnt, nArcs );

while ( nArcs > 0 ) && ( myPos < Q_LEN ) )
{
    scores_s[myPos] = srcCost;
    arcSet_s[myPos] = startId + nArcs - 1;

    myPos++;
    nArcs--;
}

__syncthreads();

//////////////////////////////////////
/// parallel propagate arc set

if ( arcSet_s[tx] > 0 )
{
   FstArc arc = arcs_g[ arcSet_s[tx] ];
   float srcCost_ = scores_s[tx];

   DoSomeJob ( &srcCost_ );

   int *dst = &(transitionData_g[arc.dst]);

   int old = atomicMax( dst, FloatToInt ( srcCost_ ) );

   ////////////////////////////////
   //// new active set

   if ( old == ILZERO )
   {
      int pos = atomicAdd ( &stateCnt, 1 );
      stateSet_s[ pos ] = arc.dst;
   }
}

/////////////////////////////////////////////
/// transfer new active set from smem to gmem

__syncthreads();

__shared__ int gPos;

if ( tx == 0 )
{
   gPos = atomicAdd ( activeSetOutSz_g, stateCnt );
}

__syncthreads();

if ( tx < stateCnt )
{
    activeSetOut_g[gPos + tx] = stateSet_s[tx];
}

__syncthreads();

但它运行得非常慢，我的意思是如果没有使用活动集（活动集 = 所有状态）则更慢，尽管活动集占所有状态的 10% 至 15%。登记压力大幅上升，入住率很低，但我认为对此无能为力。

线程之间可能有更有效的工作共享方式吗？考虑一下 3.0 上的 warp-shuffle 操作，但我必须使用 2.x 设备。

4

1 回答 1

2

通常使用多个 CUDA 内核调用来解决工作负载不均和动态工作创建的问题。这可以通过使 CPU 循环如下来完成：

//CPU pseudocode
while ( job not done) {
    doYourComputationKernel();
    loadBalanceKernel();
}

doYourComputationKernel() 必须具有启发式方法，以了解何时是停止并将控制权发送回 CPU 以平衡工作负载的好时机。这可以通过对空闲块的数量使用全局计数器来完成。每次块完成其工作或无法创建更多工作时，此计数器都会增加。当空闲块的数量超过阈值时，所有块中的工作都会保存到全局内存中，并且所有块都完成。

loadBalanceKernel() 应该接收包含所有已保存工作的全局数组和每个块的另一个全局工作计数器数组。对后者进行reduce操作可以计算出作品的总数。有了这个，可以找到每个块的作品数量。最后，内核应该复制工作，以便每个块接收相同数量的元素。

循环继续，直到所有计算完成。有一篇关于这个的好论文：http: //gamma.cs.unc.edu/GPUCOL/。这个想法是平衡非常不均匀的连续碰撞检测的负载。

于 2013-08-30T19:45:59.967 回答