CUDA 中的浮点最小值/最大值比 CPU 版本慢。为什么？

Question

我编写了一个内核，用于使用归约计算大约 100,000 个浮点数的数组的最小值和最大值（参见下面的代码）。我使用线程块将 1024 个值的块缩减为单个值（在共享内存中），然后在 CPU 上的块之间进行最终缩减。

然后我将其与仅在 CPU 上进行的串行计算进行了比较。 CUDA版本耗时2.2ms，CPU版本耗时0.21ms。为什么CUDA版本慢很多？数组大小是否不足以利用并行性，或者我的代码没有以某种方式优化？

这是 Udacity 并行编程课程练习的一部分。我通过他们的网站运行这个，所以我不知道确切的硬件是什么，但他们声称代码在实际的 GPU 上运行。

这是 CUDA 代码：

__global__ void min_max_kernel(const float* const d_logLuminance,
                            const size_t length,
                            float* d_min_logLum,
                            float* d_max_logLum) {
    // Shared working memory
    extern __shared__ float sh_logLuminance[];

    int blockWidth = blockDim.x;
    int x = blockDim.x * blockIdx.x + threadIdx.x;

    float* min_logLuminance = sh_logLuminance;
    float* max_logLuminance = sh_logLuminance + blockWidth;

    // Copy this block's chunk of the data to shared memory
    // We copy twice so we compute min and max at the same time
    if (x < length) {
        min_logLuminance[threadIdx.x] = d_logLuminance[x];
        max_logLuminance[threadIdx.x] = min_logLuminance[threadIdx.x];
    }
    else {
        // Pad if we're out of range
        min_logLuminance[threadIdx.x] = FLT_MAX;
        max_logLuminance[threadIdx.x] = -FLT_MAX;
    }

    __syncthreads();

    // Reduce
    for (int s = blockWidth/2; s > 0; s /= 2) {
        if (threadIdx.x < s) {
            if (min_logLuminance[threadIdx.x + s] < min_logLuminance[threadIdx.x]) {
                min_logLuminance[threadIdx.x] = min_logLuminance[threadIdx.x + s];
            }

            if (max_logLuminance[threadIdx.x + s] > max_logLuminance[threadIdx.x]) {
                max_logLuminance[threadIdx.x] = max_logLuminance[threadIdx.x + s];
            }
        }

        __syncthreads();
    }

    // Write to global memory
    if (threadIdx.x == 0) {
        d_min_logLum[blockIdx.x] = min_logLuminance[0];
        d_max_logLum[blockIdx.x] = max_logLuminance[0];
    }
}

size_t get_num_blocks(size_t inputLength, size_t threadsPerBlock) {
    return inputLength / threadsPerBlock +
        ((inputLength % threadsPerBlock == 0) ? 0 : 1);
}

/*
* Compute min, max over the data by first reducing on the device, then
* doing the final reducation on the host.
*/
void compute_min_max(const float* const d_logLuminance,
                    float& min_logLum,
                    float& max_logLum,
                    const size_t numRows,
                    const size_t numCols) {
    // Compute min, max
    printf("\n=== computing min/max ===\n");
    const size_t blockWidth = 1024;
    const size_t numPixels = numRows * numCols;
    size_t numBlocks = get_num_blocks(numPixels, blockWidth);

    printf("Num min/max blocks = %d\n", numBlocks);

    float* d_min_logLum;
    float* d_max_logLum;
    int alloc_size = sizeof(float) * numBlocks;
    checkCudaErrors(cudaMalloc(&d_min_logLum, alloc_size));
    checkCudaErrors(cudaMalloc(&d_max_logLum, alloc_size));

    min_max_kernel<<<numBlocks, blockWidth, sizeof(float) * blockWidth * 2>>>
        (d_logLuminance, numPixels, d_min_logLum, d_max_logLum);

    float* h_min_logLum = (float*) malloc(alloc_size);
    float* h_max_logLum = (float*) malloc(alloc_size);
    checkCudaErrors(cudaMemcpy(h_min_logLum, d_min_logLum, alloc_size, cudaMemcpyDeviceToHost));
    checkCudaErrors(cudaMemcpy(h_max_logLum, d_max_logLum, alloc_size, cudaMemcpyDeviceToHost));

    min_logLum = FLT_MAX;
    max_logLum = -FLT_MAX;

    // Reduce over the block results
    // (would be a bit faster to do it on the GPU, but it's just 96 numbers)
    for (int i = 0; i < numBlocks; i++) {
        if (h_min_logLum[i] < min_logLum) {
            min_logLum = h_min_logLum[i];
        }
        if (h_max_logLum[i] > max_logLum) {
            max_logLum = h_max_logLum[i];
        }
    }

    printf("min_logLum = %.2f\nmax_logLum = %.2f\n", min_logLum, max_logLum);

    checkCudaErrors(cudaFree(d_min_logLum));
    checkCudaErrors(cudaFree(d_max_logLum));
    free(h_min_logLum);
    free(h_max_logLum);
}

这是主机版本：

void compute_min_max_on_host(const float* const d_logLuminance, size_t numPixels) {
    int alloc_size = sizeof(float) * numPixels;
    float* h_logLuminance = (float*) malloc(alloc_size);
    checkCudaErrors(cudaMemcpy(h_logLuminance, d_logLuminance, alloc_size, cudaMemcpyDeviceToHost));
    float host_min_logLum = FLT_MAX;
    float host_max_logLum = -FLT_MAX;
    printf("HOST ");
    for (int i = 0; i < numPixels; i++) {
        if (h_logLuminance[i] < host_min_logLum) {
            host_min_logLum = h_logLuminance[i];
        }
        if (h_logLuminance[i] > host_max_logLum) {
            host_max_logLum = h_logLuminance[i];
        }
    }
    printf("host_min_logLum = %.2f\nhost_max_logLum = %.2f\n",
        host_min_logLum, host_max_logLum);
    free(h_logLuminance);
}

Answer 1

1. 正如@talonmies 建议的那样，较大尺寸的行为可能会有所不同； 100,000 确实不算多：其中大部分都适合现代 CPU 内核的整体 L1 缓存；其中一半适合单核的二级缓存。
2. 通过 PCI express 传输需要时间；在你的情况下，它可能有两倍的时间，因为你不使用固定内存。
3. 您没有将计算和 PCI express I/O 重叠（并不是说仅对 100,000 个元素有意义）
4. 您的内核相当慢，原因不止一个；其中最重要的是共享内存的广泛使用，其中大部分是不必要的

更一般地说：始终使用 nvvp（或 nvprof 获取文本信息以供进一步分析）来分析您的代码。