总和未正确存储

Question

这是我第二次尝试 CUDA 程序和并行开发。

我试图找出点积两个数组的原因。

例子：

给定

A = [1 2 3] 
B = [4 5 6]

C[0] = (1)(4) + (1)(5) + (1)(6)  
C[1] = (2)(4) + (2)(5) + (2)(6)  
C[2] = (3)(4) + (3)(5) + (3)(6)  

我初始化 2 个数组 A 和 B ，随机填充一个范围内的元素，然后将 A 中的每个元素与 B 中的每个元素相乘，并将乘积的总和存储在标识为 @987654324 的第三个数组中@。我已将数组 A、B 和 C 的大小设置为 100。

这给了我 10'000 次乘法，我使用 100 个块和 128 个线程（由于 warp 大小）并行化这些乘法。

这是我的内核函数：

__global__ void kernel(float *a, float *b, float *c, const int N) {
    if( threadIdx.x < N ) 
        c[blockIdx.x] += a[blockIdx.x] * b[threadIdx.x];
}

这是我的推理，因为聚合必须在C 中累积，它与A 中的枢轴索引具有相同的索引，因此我可以重用blockidx.x 并且“应该”正常工作；但事实并非如此。

我怀疑C 索引在线程更改时被清除或不共享，但我真的不确定这是我寻求建议的原因。

这是完整的代码，我已经明确避免使用 HANDLE_ERROR 函数包装器来缩短

#include <stdio.h>
#include <cuda.h>
#include <time.h>

#define M 100

__global__ void kernel(float *a, float *b, float *c, const int N) {
    if(threadIdx.x < N) 
        c[blockIdx.x] += a[blockIdx.x] * b[threadIdx.x];
}

void init_array(float*, const int);
void fill_array(float*, const int, const float); 
void print_array(float*, const int, *char);

int main (void) {
    srand( time(NULL) );

    float a[M], b[M], c[M] = { 0.0 };
    float *dev_a, *dev_b, *dev_c;
    const int S = sizeof(float) * M;

    init_array(a, M);
    init_array(b, M);

    print_array(a, M, "a");
    print_array(b, M, "b");
    print_array(c, M, "c");

    cudaMalloc((void**)&dev_a, S);
    cudaMalloc((void**)&dev_b, S);
    cudaMalloc((void**)&dev_c, S);

    cudaMemcpy(dev_a, a, S, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, S, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_c, c, S, cudaMemcpyHostToDevice);

    kernel<<<M, M + 28>>>(dev_a, dev_b, dev_c, M);

    cudaMemcpy(c, dev_c, S, cudaMemcpyDeviceToHost);

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    print_array(c, M, "c");

    return 0;
}

void init_array(float *a, const int N) {
   int i;  
   for(i=0; i<N; i++)
       a[i] = rand() % M + 1;
}

void fill_array(float *a, const int N, const float v) {
   int i;  
   for(i=0; i<N; i++)
       a[i] = v;
}

void print_array(float *a, const int N, char *d) {
   int i;  
   for(i=0; i<N; i++)
       printf("\n%s[%d]: %f",d, i, a[i]);
}

Answer 1

将结果累加到 C 中时，您会遇到互斥。您不能让多个线程更新同一个数组索引。解决它的一种方法是使用原子指令，在你的情况下 atomicAdd(..)。

这不起作用的原因是块 0 中的线程 0 和 1 更新了 C 数组中的相同位置。你得到一个竞争条件。

Answer 2

使用 cuBLAS 会简单得多：cublasSdot() 例程可以满足您的需求（即点积两个向量）。

这并不能帮助您学习如何编写并行代码或在 CUDA 中工作，但它会为您提供良好的性能，并且会针对不同的 GPU 进行优化。最佳做法是尽可能使用库，除非有充分的理由不这样做。

另一个答案指出您需要使用原子或其他方法来避免竞争。一种更有效的方法是让每个线程计算部分结果，进行逐块归约（有关此示例，请参见 SDK），最后让每个块中的一个线程在全局内存中进行原子加法以累积来自不同块的结果。

Answer 3

您的代码将多个值写入全局内存中的同一位置，而无需同步和考虑。您可以使用某种关键部分来修复它。这是执行您想要的代码：

#include <stdio.h>
#include <cuda.h>
#include <time.h>

#define M 5

__global__ void kernel(float *a, float *b, float *c, const int N) {
    __shared__ int lock;
    lock=0u;
    if(threadIdx.x < N) {
        __syncthreads();
        float mult = a[blockIdx.x] * b[threadIdx.x];

        int leaveLoop = 0;
        while (leaveLoop==0) {
            if (atomicExch(&lock, 1u) == 0u) {
                //critical section
                c[blockIdx.x] += mult;
                leaveLoop = 1;
                atomicExch(&lock,0u);
            }
        }
        __syncthreads();
    }
}

void init_array(float *a, const int N) {
    int i;  
    for(i=0; i<N; i++)
        a[i] = rand() % M + 1;
}

void fill_array(float *a, const int N, const float v) {
    int i;  
    for(i=0; i<N; i++)
        a[i] = v;
}

void print_array(float *a, const int N, char *d) {
    int i;  
    for(i=0; i<N; i++)
        printf("\n%s[%d]: %f",d, i, a[i]);
}

int main (void) {
    srand( time(NULL) );

    float a[M], b[M], c[M] = { 0.0 };
    float *dev_a, *dev_b, *dev_c;
    const int S = sizeof(float) * M;

    init_array(a, M);
    init_array(b, M);

    print_array(a, M, "a");
    print_array(b, M, "b");
    print_array(c, M, "c");

    cudaMalloc((void**)&dev_a, S);
    cudaMalloc((void**)&dev_b, S);
    cudaMalloc((void**)&dev_c, S);

    cudaMemcpy(dev_a, a, S, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, S, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_c, c, S, cudaMemcpyHostToDevice);

    kernel<<<M, M + 28>>>(dev_a, dev_b, dev_c, M);

    cudaMemcpy(c, dev_c, S, cudaMemcpyDeviceToHost);

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    print_array(c, M, "c");

    return 0;
}

为了实现聚合数组乘法之和的目标，我建议您尽可能使用@Tom 提出的建议。我编写此代码是因为有时使用您的代码而不是外部库更容易。