CUDA 块均值

Question

我正在尝试将 8*8 数组的块均值计算到 4*4 数组上。像这样：

我目前陷入了竞争条件类型的难题，每次我运行程序时，每个线程读取的值都不同。现在我唯一关心的是将所有块元素加在一起，稍后我会除掉我得到的总和。这是我的代码。

#include <stdio.h>
#include<math.h>

const int MAIN_SIZE = 8;
const int RESULT_SIZE = 4;

typedef int mainArray[MAIN_SIZE];
typedef int resultArray[RESULT_SIZE];

__global__ void computeMean(mainArray *main, resultArray *result) {
    int mColumn = blockIdx.x * blockDim.x + threadIdx.x;
    int mRow = blockIdx.y * blockDim.y + threadIdx.y;
    if (mRow >= MAIN_SIZE || mColumn >= MAIN_SIZE)
        return;

    // real calculation
    int rRow = std::floor(static_cast<float>(mRow / 2)),
        rColumn = std::floor(static_cast<float>(mColumn / 2));
    int x = result[rRow][rColumn] + main[mRow][mColumn];
    result[rRow][rColumn] += x;
    printf("Adding %d on %d %d at location %d %d; new value: %d\n", main[mRow][mColumn], mRow, mColumn, rRow, rColumn, result[rRow][rColumn]);
}

int main() {
    mainArray *hMain, *dMain;
    resultArray *hResult, *dResult;
    size_t mSize = MAIN_SIZE * MAIN_SIZE * sizeof(int*);
    size_t rSize = RESULT_SIZE * RESULT_SIZE * sizeof(int*);
    hMain = (mainArray *) malloc (mSize);
    hResult = (resultArray *) malloc (rSize);

    // populate arrays
    int k = 0;
    for(int i = 0; i < MAIN_SIZE; i++) {
        for(int j = 0; j < MAIN_SIZE; j++) {
            hMain[i][j] = ++k;
        }
    }
    memset(hResult, 0, rSize);

    printf("main\n");
    for(int i = 0; i < MAIN_SIZE; i++) {
        for(int j = 0; j < MAIN_SIZE; j++) {
            printf("%d  ", hMain[i][j]);
        }
        printf("\n");
    }

    printf("result\n");
    for(int i = 0; i < RESULT_SIZE; i++) {
        for(int j = 0; j < RESULT_SIZE; j++) {
            printf("%d  ", hResult[i][j]);
        }
        printf("\n");
    }

    // Allocate memory on device
    cudaMalloc(&dMain, mSize);
    cudaMalloc(&dResult, rSize);

    // Do memcopies to GPU
    cudaMemcpy(dMain, hMain, mSize, cudaMemcpyHostToDevice);
    cudaMemcpy(dResult, hResult, rSize, cudaMemcpyHostToDevice);

    dim3 block(1, 1);
    dim3 grid ((MAIN_SIZE + block.x - 1) / block.x, (MAIN_SIZE + block.y - 1) / block.y);
    computeMean<<<grid, block>>>(dMain, dResult);

    // Do memcopies back to host
    cudaMemcpy(hMain, dMain, mSize, cudaMemcpyDeviceToHost);
    cudaMemcpy(hResult, dResult, rSize, cudaMemcpyDeviceToHost);

    // validate
    if (cudaGetLastError() != cudaSuccess) {printf("cuda error\n"); return -1;}

    printf("success!\n");
    for(int i = 0; i < RESULT_SIZE; i++) {
        for(int j = 0; j < RESULT_SIZE; j++) {
            printf("%d  ", hResult[i][j]);
        }
        printf("\n");
    }

    free(hMain);
    free(hResult);
    cudaFree(dMain);
    cudaFree(dResult);
    return 0;
}

我目前是 CUDA 的新手，所以如果我从一开始就使用了错误的方法，请告诉我（我认为我的数组全错了，但我无法为二维分配动态 space一）。提前致谢。

Answer 1

这部分代码存在一些问题：

int rRow = std::floor(static_cast<float>(mRow / 2)),
    rColumn = std::floor(static_cast<float>(mColumn / 2));
int x = result[rRow][rColumn] + main[mRow][mColumn];
result[rRow][rColumn] += x;

因为 rRow 和 rColumn 是使用 mRow 和 mColumn 的整数除法找到的，所以很明显 threadIdx.x 为 0 的线程or 1 将产生相同的 rColumn 结果，并且还有许多其他重复实例。因此，您的线程是：

在没有任何命令或控制的情况下写入同一位置
从其他线程可能写入的位置读取

CUDA 不会为您解决这些危险。您必须采取特定的编程步骤来处理它们。（您还有一个逻辑错误，即您将主要数据添加到现有结果数据，然后将该总和再次添加到您的结果数据；我有理由相信这不是什么你想要的。）解决线程危险的两种典型方法是：

使用atomics
使用classical parallel reduction方法

为简单起见，我将提供一个演示第一种方法的重写代码。

$ cat t1324.cu
#include <stdio.h>
#include<math.h>

const int MAIN_SIZE = 8;
const int RESULT_SIZE = 4;

typedef int mainArray[MAIN_SIZE];
typedef int resultArray[RESULT_SIZE];

__global__ void computeMean(mainArray *main, resultArray *result) {
    int mColumn = blockIdx.x * blockDim.x + threadIdx.x;
    int mRow = blockIdx.y * blockDim.y + threadIdx.y;
    if (mRow >= MAIN_SIZE || mColumn >= MAIN_SIZE)
        return;

    // real calculation
    int rRow = std::floor(static_cast<float>(mRow / 2)),
        rColumn = std::floor(static_cast<float>(mColumn / 2));
    //int x = result[rRow][rColumn] + main[mRow][mColumn];
    //result[rRow][rColumn] += x;
    atomicAdd(&(result[rRow][rColumn]), main[mRow][mColumn]);
    //printf("Adding %d on %d %d at location %d %d; new value: %d\n", main[mRow][mColumn], mRow, mColumn, rRow, rColumn, result[rRow][rColumn]);
}

int main() {
    mainArray *hMain, *dMain;
    resultArray *hResult, *dResult;
    size_t mSize = MAIN_SIZE * MAIN_SIZE * sizeof(int*);
    size_t rSize = RESULT_SIZE * RESULT_SIZE * sizeof(int*);
    hMain = (mainArray *) malloc (mSize);
    hResult = (resultArray *) malloc (rSize);

    // populate arrays
    //int k = 0;
    for(int i = 0; i < MAIN_SIZE; i++) {
        for(int j = 0; j < MAIN_SIZE; j++) {
            hMain[i][j] = 1; //++k;
        }
    }
    memset(hResult, 0, rSize);

    printf("main\n");
    for(int i = 0; i < MAIN_SIZE; i++) {
        for(int j = 0; j < MAIN_SIZE; j++) {
            printf("%d  ", hMain[i][j]);
        }
        printf("\n");
    }

    printf("result\n");
    for(int i = 0; i < RESULT_SIZE; i++) {
        for(int j = 0; j < RESULT_SIZE; j++) {
            printf("%d  ", hResult[i][j]);
        }
        printf("\n");
    }

    // Allocate memory on device
    cudaMalloc(&dMain, mSize);
    cudaMalloc(&dResult, rSize);

    // Do memcopies to GPU
    cudaMemcpy(dMain, hMain, mSize, cudaMemcpyHostToDevice);
    cudaMemcpy(dResult, hResult, rSize, cudaMemcpyHostToDevice);

    dim3 block(1, 1);
    dim3 grid ((MAIN_SIZE + block.x - 1) / block.x, (MAIN_SIZE + block.y - 1) / block.y);
    computeMean<<<grid, block>>>(dMain, dResult);

    // Do memcopies back to host
    cudaMemcpy(hMain, dMain, mSize, cudaMemcpyDeviceToHost);
    cudaMemcpy(hResult, dResult, rSize, cudaMemcpyDeviceToHost);

    // validate
    if (cudaGetLastError() != cudaSuccess) {printf("cuda error\n"); return -1;}

    printf("success!\n");
    for(int i = 0; i < RESULT_SIZE; i++) {
        for(int j = 0; j < RESULT_SIZE; j++) {
            printf("%d  ", hResult[i][j]);
        }
        printf("\n");
    }

    free(hMain);
    free(hResult);
    cudaFree(dMain);
    cudaFree(dResult);
    return 0;
}
$ nvcc -arch=sm_35 -o t1324 t1324.cu
$ cuda-memcheck ./t1324
========= CUDA-MEMCHECK
main
1  1  1  1  1  1  1  1
1  1  1  1  1  1  1  1
1  1  1  1  1  1  1  1
1  1  1  1  1  1  1  1
1  1  1  1  1  1  1  1
1  1  1  1  1  1  1  1
1  1  1  1  1  1  1  1
1  1  1  1  1  1  1  1
result
0  0  0  0
0  0  0  0
0  0  0  0
0  0  0  0
success!
4  4  4  4
4  4  4  4
4  4  4  4
4  4  4  4
========= ERROR SUMMARY: 0 errors
$

其他一些注意事项：

我更改了初始化数据，以便快速识别正确的输出。
您的代码：
```
int rRow = std::floor(static_cast<float>(mRow / 2)),
```
我相信不是按照你的想法去做。 mRow/2 是你写的整数除法。随后转换为 float 然后使用 floor 没有效果，我不认为。它不会伤害我能看到的任何东西（我很确定你想要在这里进行整数除法），所以我保持原样。如果你想要浮点除法，你需要首先将你的两个整数操作数之一转换为浮点数。您编写的代码不会那样做。（它投射结果。）

CUDA 块均值

CUDA Block Mean

cuda

race-condition

multidimensional-array