c ++ CUDA sum of some array elements

В блоке потоков у вас есть несколько потоков (по y), каждый из которых пытается обновить одно и то же местоположение. CUDA не разбирается в этом автоматически.

Типичным подходом здесь было бы использование атомики, чтобы несколько потоков, добавляемых в одно и то же место, происходили упорядоченным образом.

Поскольку вы используете двойные количества, CUDA обеспечивает только собственное атомарное добавление для устройств с вычислительной мощностью 6.x и выше. В противном случае вы должны использовать метод, приведенный в руководстве по программированию, чтобы обеспечить операцию двойного атомарного добавления. (В качестве альтернативы вы можете переключиться с double на float).

Вот пример вашего измененного кода, работающего на устройстве cc 7.0:

$ cat t1812.cu
#include <stdio.h>

struct Node {
    double xyz[3];
};

cudaError_t addWithCuda(Node* c, Node* a, unsigned int* order);


unsigned int node_in_face[24] = {0, 3, 2, 1,   ///< Face 0
                                   0, 1, 5, 4,   ///< Face 1 //-V112
                                   1, 2, 6, 5,   ///< Face 2
                                   2, 3, 7, 6,   ///< Face 3
                                   3, 0, 4, 7,   ///< Face 4 //-V112
                                   4, 5, 6, 7};  ///< Face 5 //-V112

__global__ void centroid(Node* nodes_of_value, Node* nodes_centroid, unsigned int* order) {
  int n_centroid = threadIdx.x;
  int n_node     = threadIdx.y;
  int z          = threadIdx.z;
  int n_value    = blockIdx.x;

  atomicAdd(&(nodes_centroid[n_centroid].xyz[z]),  nodes_of_value[n_value * 8 + order[n_centroid * 4 + n_node]].xyz[z]);
  __syncthreads();

}

int main()
{
    const int arraySize = 8;
    Node nodes[arraySize] = {{0., 0., 0.},
                             {1., 0., 0.},
                             {1., 1., 0.},
                             {0., 1., 0.},
                             {0., 0., 1.},
                             {1., 0., 1.},
                             {1., 1., 1.},
                             {0., 1., 1.}};  // Координаты каждой точки
    Node centroids[6] = {{0., 0., 0.},
                         {0., 0., 0.},
                         {0., 0., 0.},
                         {0., 0., 0.},
                         {0., 0., 0.},
                         {0., 0., 0.}};  // 6 пустых точек, куда нужно просуммировать значения

    // Add vectors in parallel.
    cudaError_t cudaStatus = addWithCuda(nodes, centroids, node_in_face);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }

    printf(
        "{0 = %f, %f, %f}\n"
        "{1 = %f, %f, %f}\n"
        "{2 = %f, %f, %f}\n"
        "{3 = %f, %f, %f}\n"
        "{4 = %f, %f, %f}\n"
        "{5 = %f, %f, %f}\n",
        centroids[0].xyz[0], centroids[0].xyz[1], centroids[0].xyz[2],
        centroids[1].xyz[0], centroids[1].xyz[1], centroids[1].xyz[2],
        centroids[2].xyz[0], centroids[2].xyz[1], centroids[2].xyz[2],
        centroids[3].xyz[0], centroids[3].xyz[1], centroids[3].xyz[2],
        centroids[4].xyz[0], centroids[4].xyz[1], centroids[4].xyz[2],
        centroids[5].xyz[0], centroids[5].xyz[1], centroids[5].xyz[2]);

    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(Node* nodes, Node* centroids, unsigned int* order)
{
    Node* dev_nodes               = 0;
    Node* dev_centroids           = 0;
    unsigned int* dev_order       = 0;

    cudaError_t cudaStatus;

    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
    }

    // Allocate GPU buffers for three vectors (two input, one output)    .
    cudaStatus = cudaMalloc((void**)&dev_nodes, 8 * sizeof(Node));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
    }

    cudaStatus = cudaMalloc((void**)&dev_centroids, 6 * sizeof(Node));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
    }

    cudaStatus = cudaMalloc((void**)&dev_order, 24 * sizeof(unsigned int));
    if(cudaStatus != cudaSuccess) {
      fprintf(stderr, "cudaMalloc failed!");
    }

    // Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpy(dev_nodes, nodes, 8 * sizeof(Node), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
    }

    cudaStatus = cudaMemcpy(dev_order, order, 24 * sizeof(unsigned int), cudaMemcpyHostToDevice);
    if(cudaStatus != cudaSuccess) {
      fprintf(stderr, "cudaMemcpy failed!");
    }

    // Launch a kernel on the GPU with one thread for each element.

    dim3 threadsPerBlock(6, 4, 3);
    centroid<<<1, threadsPerBlock >>>(dev_nodes, dev_centroids, dev_order);

    // Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
    }

    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
    }

    // Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(centroids, dev_centroids, 6 * sizeof(Node), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
    }

    cudaFree(dev_nodes);
    cudaFree(dev_centroids);

    return cudaStatus;
}
$ nvcc -o t1812 t1812.cu -arch=sm_70
$ ./t1812
{0 = 2.000000, 2.000000, 0.000000}
{1 = 2.000000, 0.000000, 2.000000}
{2 = 4.000000, 2.000000, 2.000000}
{3 = 2.000000, 4.000000, 2.000000}
{4 = 0.000000, 2.000000, 2.000000}
{5 = 2.000000, 2.000000, 4.000000}
$

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions