通过 cuda 内核缩减找到最小值和最大值

Question

这是我的内核代码

typedef unsigned char Npp8u;
...
    // Kernel Implementation
__device__ unsigned int min_device;
__device__ unsigned int max_device;


__global__ void findMax_Min(Npp8u * data, int numEl){
    int index = blockDim.x*blockIdx.x + threadIdx.x;
    int shared_index = threadIdx.x;

    __shared__ Npp8u data_shared_min[BLOCKDIM];
    __shared__ Npp8u data_shared_max[BLOCKDIM];

    // check index condition
    if(index < numEl){
        data_shared_min[shared_index] = data[index]; //pass values from global to shared memory
        __syncthreads();
        data_shared_max[shared_index] = data[index]; //pass values from global to shared memory


        for (unsigned int stride = BLOCKDIM/2; stride > 0; stride >>= 1) {
            if(threadIdx.x <  stride){
                if(data_shared_max[threadIdx.x] <  data_shared_max[threadIdx.x+stride]) data_shared_max[shared_index] = data_shared_max[shared_index+stride];
                if(data_shared_min[threadIdx.x]>  data_shared_min[threadIdx.x+stride]) data_shared_min[shared_index] = data_shared_min[shared_index+stride];
            }
            __syncthreads();
        }
        if(threadIdx.x == 0  ){
            atomicMin(&(min_device), (unsigned int)data_shared_min[threadIdx.x ]);
            //min_device =10;
            __syncthreads();
            atomicMax(&(max_device), (unsigned int)data_shared_max[threadIdx.x ]);
        }
    }else{
        data_shared_min[shared_index] = 9999;
    }
}

我有一个 512x512 的图像，我想找到最小和最大像素值。 data 是图像的一维版本。此代码适用于最大值但不适用于最小值。正如我从 matlab 中检查的那样，最大值为 202，最小值为 10，但它发现最小值为 0。这是我的内核代码和 memcpy 调用

int main(){
    // Host parameter declarations.
    Npp8u * imageHost;
    int   nWidth, nHeight, nMaxGray;

    // Load image to the host.
    std::cout << "Load PGM file." << std::endl;
    imageHost = LoadPGM("lena_before.pgm", nWidth, nHeight, nMaxGray);

    // Device parameter declarations.
    Npp8u    * imageDevice;
    unsigned int   max, min;
    size_t size = sizeof(Npp8u)*nWidth*nHeight;

    cudaMalloc((Npp8u**)&imageDevice, size);

    cudaMemcpy(imageDevice, imageHost, size, cudaMemcpyHostToDevice);

    int numPixels = nWidth*nHeight;

    dim3 numThreads(BLOCKDIM);
    dim3 numBlocks(numPixels/BLOCKDIM + (numPixels%BLOCKDIM == 0 ? 0 : 1));

    findMax_Min<<<numBlocks, numThreads>>>(imageDevice,numPixels);
    cudaMemcpyFromSymbol(&max,max_device, sizeof(max_device), 0, cudaMemcpyDeviceToHost);
    cudaMemcpyFromSymbol(&min,min_device, sizeof(min_device), 0, cudaMemcpyDeviceToHost);


    printf("Min value for image : %i\n", min);
    printf("Max value for image : %i\n", max);
...

另一个有趣的事情是在内核调用之后更改cudaMemcpy 的顺序也会导致故障并且值都被读取为零。我没有看到问题。有没有人看到被遮挡的部分？

Answer 1

您可能想要进行 cuda 错误检查。您可能还希望将min_device 初始化为一个较大的值，并将max_device 初始化为零。您的缩减方法还有其他与步幅相关的问题（当您将步幅添加到 threadIdx.x 时，奇数大小图像的最后一个块会发生什么，它可能超出共享内存中定义的图像范围），但我不认为这对 512x512 图像很重要。如果min_device 恰好从零开始，那么您的所有 atomicMin 操作将始终保持零。

您可以尝试像这样初始化min_device 和max_device：

__device__ unsigned int min_device = 9999;
__device__ unsigned int max_device = 0;

对于最后的 cudamemcpy 调用，您将 4 个字节（max_device 的大小）复制到一个单字节变量（Npp8u 最大值）中，同样地用于最小值。所以这是个问题。由于您使用的是指针，因此复制操作肯定会覆盖您不希望的内容。如果编译器按照您定义变量的方式顺序存储变量，则一个复制操作会覆盖另一个变量，我认为这可以解释您所看到的行为。如果您将min 和max 创建为unsigned int 数量，我认为这个问题会消失。

编辑：由于您没有显示实际的块尺寸，因此您可能仍然对减少有问题。您可能想要更改此行：

        if(threadIdx.x <  stride){

类似于：

        if((threadIdx.x <  stride) && ((index + stride)< numEl)){

这个或类似的东西应该可以纠正我在第一段中提到的危险。我猜你想用这条线来解释危险：

    data_shared_min[shared_index] = 9999;

但不能保证在共享内存中设置的数据元素被其他线程读取之前执行代码行。我也不知道当您将 9999 值分配给字节数量时会发生什么，但这可能不是您所期望的。