0

我对网格中的控制线程块有疑问。

我的来源是图像上的递归作业。但在处理过程中，许多块满足结束条件的次数约为 8 次。只有几个块导致执行再次循环超过 16 次。所以我想跳过满足执行结束条件的块。

有可能的？

__global__ main(){
/* previous */
int *blockMap;
cudaMalloc((void**)&blockMap, sizeof(int) * nXBlockNum * nYBlockNum);
cudaMemset((void**)&blockMap, 0, sizeof(int) * nXBlockNum * nYBlockNum);

kernel<<<nblocks, nthreads>>>(inputimage, outputbuffer, blockmap);
/* after */}

__global__
kernel(byte* inputeimage, byte* outputbuffer, int* blockmap) {
    __shared__ int *skipFlag;


    if((blockDim.x * threadIdx.y + threadIdx.x) == 0)
    {
        *skipFlag = g_bMap[blockIdx.y * gridDim.x + blockIdx.x];
    }

    if(*skipFlag == 0)
    {
              /* recursive job */
    }
}
4

2 回答 2

1

是的，你可以这样做，但你展示的内核代码并不完全是这样做的。假设您希望每个 block 都有一个整数标志，那么代码应该类似于：

__global__
kernel(byte* inputeimage, byte* outputbuffer, int* blockmap) {
    __shared__ int skipFlag;


    if (threadIdx.x == 0)
    {
        skipFlag = g_bMap[blockIdx.x];
    }
    __syncthreads();

    if(skipFlag == 0)
    {
              /* recursive job */
    }
}

这里每个块中的第一个线程从全局加载该特定块的标志并将其存储到共享内存整数变量中。在逐块同步之后，每个线程都可以读取该值并相应地对其进行操作。

于 2013-03-15T16:16:08.023 回答
0

我不太明白你的问题，但这种方式似乎是正确的。每个块skipFlag在它们的共享内存中都有一个唯一的，得到的true那些不会执行其余的代码。

而且，也许__syncthreads()两个ifs之间的一个好主意。

于 2013-03-15T16:06:07.960 回答