1，代码重新编辑

为了地毯式地检查结果的正确性，这里修改了代码

主要步骤为

step1:  data_p指向的空间中，分别生成随机数；

step2:  分别拷贝到gpu的sendbuff的显存中；

step3:  通过nccl_all_reduce sum；

step4:  取回 recvbuff中的数据；

step5:  将data_p中的数据 allreduce sum；

step6: 对比 recvbuff中的数据与 data_p中的数据的一致性；

#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <omp.h>


//NN is tile length
#define NN 312
#define MM 156
#define MATRIX_A 0xB5026F5AA96619E9ULL
#define UM 0xFFFFFFFF80000000ULL /* Most significant 33 bits */
#define LM 0x7FFFFFFFULL /* Least significant 31 bits */

class thread_mt199937{

public:
void srand(unsigned long long seed)
{
    init_genrand64(seed);
}

/* generates a random number on [0, 2^63-1]-interval */
long long genrand64_int63(void)
{
    return (long long)(genrand64_int64() >> 1);
}

/* generates a random number on [0,1]-real-interval */
double genrand64_real1(void)
{
    return (genrand64_int64() >> 11) * (1.0/9007199254740991.0);
}

/* generates a random number on [0,1)-real-interval */
double genrand64_real2(void)
{
    return (genrand64_int64() >> 11) * (1.0/9007199254740992.0);
}

/* generates a random number on (0,1)-real-interval */
double genrand64_real3(void)
{
    return ((genrand64_int64() >> 12) + 0.5) * (1.0/4503599627370496.0);
}


private:
/* The array for the state vector */
unsigned long long mt[NN];
/* mti==NN+1 means mt[NN] is not initialized */
int mti=NN+1;

/* initializes mt[NN] with a seed */
void init_genrand64(unsigned long long seed)
{
    mt[0] = seed;
    for (mti=1; mti<NN; mti++)
        mt[mti] =  (6364136223846793005ULL * (mt[mti-1] ^ (mt[mti-1] >> 62)) + mti);
}

/* initialize by an array with array-length */
/* init_key is the array for initializing keys */
/* key_length is its length */
void init_by_array64(unsigned long long init_key[],
		     unsigned long long key_length)
{
    unsigned long long i, j, k;
    init_genrand64(19650218ULL);
    i=1; j=0;
    k = (NN>key_length ? NN : key_length);
    for (; k; k--) {
        mt[i] = (mt[i] ^ ((mt[i-1] ^ (mt[i-1] >> 62)) * 3935559000370003845ULL))
          + init_key[j] + j; /* non linear */
        i++; j++;
        if (i>=NN) { mt[0] = mt[NN-1]; i=1; }
        if (j>=key_length) j=0;
    }
    for (k=NN-1; k; k--) {
        mt[i] = (mt[i] ^ ((mt[i-1] ^ (mt[i-1] >> 62)) * 2862933555777941757ULL))
          - i; /* non linear */
        i++;
        if (i>=NN) { mt[0] = mt[NN-1]; i=1; }
    }

    mt[0] = 1ULL << 63; /* MSB is 1; assuring non-zero initial array */
}

/* generates a random number on [0, 2^64-1]-interval */
unsigned long long genrand64_int64(void)
{
    int i;
    unsigned long long x;
    static unsigned long long mag01[2]={0ULL, MATRIX_A};

    if (mti >= NN) { /* generate NN words at one time */

        /* if init_genrand64() has not been called, */
        /* a default initial seed is used     */
        if (mti == NN+1)
            init_genrand64(5489ULL);

        for (i=0;i<NN-MM;i++) {
            x = (mt[i]&UM)|(mt[i+1]&LM);
            mt[i] = mt[i+MM] ^ (x>>1) ^ mag01[(int)(x&1ULL)];
        }
        for (;i<NN-1;i++) {
            x = (mt[i]&UM)|(mt[i+1]&LM);
            mt[i] = mt[i+(MM-NN)] ^ (x>>1) ^ mag01[(int)(x&1ULL)];
        }
        x = (mt[NN-1]&UM)|(mt[0]&LM);
        mt[NN-1] = mt[MM-1] ^ (x>>1) ^ mag01[(int)(x&1ULL)];

        mti = 0;
        //printf("LL::\n");
    }

    x = mt[mti++];

    x ^= (x >> 29) & 0x5555555555555555ULL;
    x ^= (x << 17) & 0x71D67FFFEDA60000ULL;
    x ^= (x << 37) & 0xFFF7EEE000000000ULL;
    x ^= (x >> 43);

    return x;
}

};


#endif

//#define NN 312

// multi thread accelerating， to be singleton
class parallel_mt19937{
public:
//    int _base_seed;
void srand(unsigned long long seed)
{
    _base_seed = seed;
}
void rand_float(float* A, unsigned long len);

private:
    unsigned long long _base_seed = 0;
};

//int pmt19937::_base_seed = 0;

void parallel_mt19937::rand_float(float* A, unsigned long len)
{
    #pragma omp parallel
    {
        unsigned long block_dim = omp_get_num_threads();
        unsigned long thid = omp_get_thread_num();

        if(thid == block_dim -1)
            std::cout << "Here are " << thid << " threads generating random number ..."<<std::endl;
        unsigned long tile_count = (len+NN-1)/NN;

        thread_mt199937 *t_mt_p = new thread_mt199937();// to be singleton

        for(unsigned long tile_id=thid; tile_id<tile_count; tile_id+=block_dim){
            //each tile has a specific seed: (_base_seed + tile_id) to smt19937, to keep consistence
            unsigned long tile_seed = _base_seed + tile_id;
            t_mt_p->srand(tile_seed);

            //if(thid == 35)                std::cout << "Hello from thread " << thid << std::endl;
            unsigned long tile_idx_start = tile_id*NN;
            unsigned long tile_idx_end = (((tile_id+1)*NN) <= len)? (tile_id+1)*NN: len;

            for(unsigned long idx = tile_idx_start; idx<tile_idx_end; idx++)
                A[idx] = float(t_mt_p->genrand64_real2());
        }
        delete t_mt_p;
    }
}



void init_data_float(float** data_p, int nDev, int size,  unsigned long seed)
{
  for(int idx=0; idx<nDev; idx++){
    *(data_p+idx) = nullptr;
    *(data_p+idx) = (float*)malloc(size*sizeof(float));
  }

  parallel_mt19937 gen_rand;

  for(int idx =0; idx<nDev; idx++){
    gen_rand.srand(seed+idx);
    gen_rand.rand_float(*(data_p + idx), size);
  }

}

void host_all_reduce_origin(float** sendbuff, int nDev, int size)
{

  for(int idx=1; idx<nDev; idx++){
#pragma omp paralell for
    for(int i=0; i<size; i++)
      sendbuff[0][i] += sendbuff[idx][i];
  }
}

void check_equal(float** result_recv_data_buff, float* sendbuff_0, int nDev, int size)
{

  for(int idx=0; idx<nDev; idx++){
#pragma omp paralell for
    for(int i=0; i<size; i++){
      if(result_recv_data_buff[idx][i] != sendbuff_0[i])
        printf("ERROR: %7.4f != %7.4f  idx=%d, i=%d\n", result_recv_data_buff[idx][i], sendbuff_0[i], idx, i);
        return;
    }
  }
}

void free_buff(float** buffArray, int n)
{
  for(int i=0; i<n; i++){
    if(buffArray[i] != nullptr){
      free(buffArray[i]);
    }
  }
}

//

#include <stdlib.h>
#include <stdio.h>
#include "cuda_runtime.h"
#include "nccl.h"





#define CUDACHECK(cmd) do {                         \
  cudaError_t err = cmd;                            \
  if (err != cudaSuccess) {                         \
    printf("Failed: Cuda error %s:%d '%s'\n",       \
        __FILE__,__LINE__,cudaGetErrorString(err)); \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)


#define NCCLCHECK(cmd) do {                         \
  ncclResult_t res = cmd;                           \
  if (res != ncclSuccess) {                         \
    printf("Failed, NCCL error %s:%d '%s'\n",       \
        __FILE__,__LINE__,ncclGetErrorString(res)); \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)


int main(int argc, char* argv[])
{
  ncclComm_t comms[4];


  //managing 4 devices
  //int nDev = 4;
  int nDev = 2;
  int size = 32*1024*1024*16;
  int devs[4] = { 0, 1, 2, 3 };

  float** data_p = (float**)malloc(nDev*sizeof(float*));
  init_data_float(data_p, nDev, size, 2024);



  //allocating and initializing device buffers
  float** sendbuff = (float**)malloc(nDev * sizeof(float*));
  float** recvbuff = (float**)malloc(nDev * sizeof(float*));
  cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);


  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(i));
    CUDACHECK(cudaMalloc((void**)sendbuff + i, size * sizeof(float)));
    CUDACHECK(cudaMalloc((void**)recvbuff + i, size * sizeof(float)));

    //CUDACHECK(cudaMemset(sendbuff[i], 1, size * sizeof(float)));
    CUDACHECK(cudaMemcpy(sendbuff[i], data_p[i], size*sizeof(float), cudaMemcpyHostToDevice));
    //CUDACHECK(cudaMemset(recvbuff[i], 0, size * sizeof(float)));
    CUDACHECK(cudaStreamCreate(s+i));
  }

printf("LL:: 04\n");

  //initializing NCCL
  NCCLCHECK(ncclCommInitAll(comms, nDev, devs));
printf("LL:: 05\n");

   //calling NCCL communication API. Group API is required when using
   //multiple devices per thread
  NCCLCHECK(ncclGroupStart());
  for (int i = 0; i < nDev; ++i)
    NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, ncclFloat, ncclSum,
        comms[i], s[i]));
  NCCLCHECK(ncclGroupEnd());
printf("LL:: 06\n");
/*LL::
In a sum allreduce operation between k ranks,
each rank will provide an array in of N values,
and receive identical results in array out of N values,
where out[i] = in0[i]+in1[i]+…+in(k-1)[i]
*/
  //synchronizing on CUDA streams to wait for completion of NCCL operation
  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(i));
    CUDACHECK(cudaStreamSynchronize(s[i]));
  }
printf("LL:: 07\n");
/*check:
each recvbuff[0] = sendbuff_0[0]+ sendbuff_1[0];
0x01010101 + 0x01010101 = 0x02020202;
*/
  float** result_recv_data_buff;
  result_recv_data_buff = (float**)malloc(nDev*sizeof(float*));
  init_data_float(result_recv_data_buff, nDev, size, 2024+nDev);

printf("LL:: 08\n");
  for(int idx=0; idx<nDev; idx++){
    cudaMemcpy(result_recv_data_buff[idx], recvbuff[idx], size*sizeof(float), cudaMemcpyDeviceToHost);
  }

printf("LL:: 09\n");
  host_all_reduce_origin(data_p, nDev, size);

printf("LL:: 10\n");
// ditanshi check
  check_equal(result_recv_data_buff, data_p[0], nDev, size);

printf("LL:: 11\n");

  free_buff(data_p, nDev);
  free_buff(result_recv_data_buff, nDev);
  //free device buffers
  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(i));
    CUDACHECK(cudaFree(sendbuff[i]));
    CUDACHECK(cudaFree(recvbuff[i]));
  }


  //finalizing NCCL
  for(int i = 0; i < nDev; ++i)
      ncclCommDestroy(comms[i]);


  printf("Success \n");
  return 0;
}

Makefile

EXE := ex_1_1_SingleProcessSingleThreadMultipleDevices

all: $(EXE)

INC := -I /usr/local/cuda/include -I /home/hipper/ex_nccl_20240701/local/include/
LD_FLAGS := -L /usr/local/cuda/lib64 -L /home/hipper/ex_nccl_20240701/local/lib -lcudart -lnccl -fopenmp



%: %.cpp
	g++ -g $< -o $@ $(INC) $(LD_FLAGS)



.PHONY: clean
clean:
	-rm -rf $(EXE)

2，编译运行

3，问题

显存占用了 7GB 和 8GB，实际数据应该只有2GB，recvbuff 2GB，总共4GB