nccl 04 nvidia 官方小程序

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>

// seed_base
// void gen_512(seed, A, start_idx, N)
// gen_512(seed_base+0, A, 0*512, N);  gen_512(seed_base + 1, A+1*512); gen_512(seed_base + 2, A+2*512);     ...    gen_512(seed_base + n, A, , N); ...
#if 1





//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;
    }
}

//

//#define BUILD_MAIN

//
#ifdef BUILD_MAIN

void print_matrix(unsigned long m, unsigned long n, float* A, unsigned long lda)
{
    for(unsigned long i=m-7; i<m; i++){
        for(unsigned long j=n-7; j<n; j++){
            printf("%7.4f ", A[i + j*lda]);
        }
        printf("\n");
    }
}


#include <string.h>

int main(){

    unsigned long m = 312*4*32*32*128;
    unsigned long n = 36;
    unsigned long lda = m;

    float* A = nullptr;
printf("LL: 00\n");
    A = (float*)malloc(lda * n * sizeof(float));
    //memset(A, 0x7F, lda*n*sizeof(float));
printf("LL:: 01\n");
    parallel_mt19937 gen_rand;
    gen_rand.srand(2024);
    gen_rand.rand_float(A, lda*n);
printf("LL:: 02\n");
    print_matrix(m, n, A, lda);

    free(A);
    return 0;
}

#endif

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

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：/a/762189.html

如若内容造成侵权/违法违规/事实不符，请联系我们进行投诉反馈qq邮箱809451989@qq.com，一经查实，立即删除！

相关文章

什么是原始权益人？

什么是原始权益人？

摘要&#xff1a;每天学习一点金融小知识 原始权益人&#xff0c;在资产证券化&#xff08;ABS&#xff09;和公募REITs等金融产品中&#xff0c;指的是证券化基础资产的原始所有者&#xff0c;即金融产品的真正融资方。他们是按照相关规定及约定向资产支持专项计划转移其合法拥…
阅读更多...
RabbitMQ消息可靠性等机制详解（精细版三）

RabbitMQ消息可靠性等机制详解（精细版三）

目录 七 RabbitMQ的其他操作 7.1 消息的可靠性(发送可靠) 7.1.1 confim机制(保证发送可靠) 7.1.2 Return机制(保证发送可靠) 7.1.3 编写配置文件 7.1.4 开启Confirm和Return 7.2 手动Ack(保证接收可靠) 7.2.1 添加配置文件 7.2.2 手动ack 7.3 避免消息重复消费 7.3.…
阅读更多...
【数据结构】计数排序等排序

【数据结构】计数排序等排序

&#x1f4e2;博客主页&#xff1a;https://blog.csdn.net/2301_779549673 &#x1f4e2;欢迎点赞 &#x1f44d; 收藏 ⭐留言 &#x1f4dd; 如有错误敬请指正&#xff01; &#x1f4e2;本文由 JohnKi 原创&#xff0c;首发于 CSDN&#x1f649; &#x1f4e2;未来很长&#…
阅读更多...
企业软文投放为什么要选择包收录媒体？

企业软文投放为什么要选择包收录媒体？

如今这个信息时代企业想要有效的将品牌推广出去&#xff0c;那选择推广方式至关重要。软文投放作为一种常见的品牌推广方式&#xff0c;其效果往往取决于投放的媒体质量。而在众多媒体中&#xff0c;包收录媒体凭借其独特的优势&#xff0c;成为了企业软文投放的明智之选。 一…
阅读更多...
Nuxt3 的生命周期和钩子函数（七）

Nuxt3 的生命周期和钩子函数（七）

title: Nuxt3 的生命周期和钩子函数&#xff08;七&#xff09; date: 2024/6/30 updated: 2024/6/30 author: cmdragon excerpt: 摘要&#xff1a;文章阐述了Nuxt3中Nitro生命周期钩子的使用&#xff0c;如nitro:config自定义配置、nitro:init注册构建钩子、nitro:build:be…
阅读更多...
为什么企业应用开发，c++干不过java？

为什么企业应用开发，c++干不过java？

在开始前刚好我有一些资料&#xff0c;是我根据网友给的问题精心整理了一份「c的资料从专业入门到高级教程」&#xff0c; 点个关注在评论区回复“888”之后私信回复“888”&#xff0c;全部无偿共享给大家&#xff01;&#xff01;&#xff01; C/C这种东西&#xff0c;根本…
阅读更多...
15_软件程序设计基础

15_软件程序设计基础

目录 嵌入式软件开发原理 宿主机和目标机 交叉编译 交叉调试 嵌入式软件开发特点和挑战 开发工具 程序设计语言基本概念 解释和编译 常见程序设计语言 程序设计语言的基本成分 编译程序基本原理 嵌入式软件开发原理 宿主机和目标机 嵌入式软件开发不同于传统软件开…
阅读更多...
QT5:在窗口右上角显示图标

QT5:在窗口右上角显示图标

目录 一、环境与目标 二、实现逻辑&#xff08;纯代码&#xff09;与效果 三、参考代码 四、总结 一、环境与目标 qt版本&#xff1a;5.12.7 windows 11 下的 Qt Designer &#xff08;已搭建&#xff09; 目标&#xff1a;使用嵌套布局的方式将两个按钮显示在窗口右上角…
阅读更多...
C++专业面试真题（1）学习

C++专业面试真题（1）学习

进程有多少种状态&#xff0c;如何转换 创建&#xff1a;一个进程启动&#xff0c;首先进入创建状态&#xff0c;需要获取系统资源创建进程管理科PCB完成资源分配。就绪态&#xff1a;在创建完成后&#xff0c;进程已经准备好&#xff0c;处于就绪状态&#xff0c;但是还未获得…
阅读更多...
深度势能生成器(DP-GEN)入门讲解

深度势能生成器(DP-GEN)入门讲解

文章目录 1.原子间相互作用1.为什么研究原子间相互作用2.研究原子间相互作用的传统方法 2.深度学习研究原子间相互作用1.深度势能平滑模型(DeepPot-se)2.Deep Potential 模型训练3.同步学习→充足采样&筛选样本 3.DP-GEN操作及运行1.DP-GEN主流程2.DP-GEN基本命令3.生成初始…
阅读更多...
Sui创始团队在竞速环节中的快问快答

Sui创始团队在竞速环节中的快问快答

在Sui Basecamp活动期间&#xff0c;Sui区块链的最初贡献者在Oracle红牛赛车模拟器上展示了他们的技术能力&#xff0c;在驾驶圈时回答了有关Sui的问题。 Evan Cheng&#xff08;又名Revvin’ Evan&#xff09;在解释Mysticeti创下区块链最终性记录的同时保持着他的驾驶线路。…
阅读更多...
【深度好文】LLMOps揭秘：AI工作流程的高效管理之道！

【深度好文】LLMOps揭秘：AI工作流程的高效管理之道！

可以关注我的公众号&#xff1a;Halo咯咯 01。 概述 将大型语言模型&#xff08;LLMs&#xff09;的强大能力与机器学习运维&#xff08;MLOps&#xff09;的有序结构相结合&#xff0c;团队能够以更高效的方式工作&#xff0c;而非仅仅增加劳动强度。团队的焦点可以专注于开…
阅读更多...
Redis分布式集群部署

Redis分布式集群部署

目录 一. 原理简述 二. 集群配置​​​​​​​ 2.1 环境准备 2.2 编译安装一个redis 2.3 创建集群 2.4 写入数据测试 实验一&#xff1a; 实验二&#xff1a; 实验三&#xff1a; 实验四&#xff1a; 添加节点 自动分配槽位 提升节点为master&#xff1a; 实验…
阅读更多...
Spring Security 认证流程

Spring Security 认证流程

Spring Scurity是spring生态下用于认证和授权的框架&#xff0c;具有高度的灵活性和可扩展行&#xff0c;本节主要对Spring Security的认证过程中进行概括性的介绍&#xff0c;主要介绍在该过程中&#xff0c;会涉及到哪些组件以及每个组件所承担的职责&#xff0c;希望大家可以…
阅读更多...
Java [ 基础 ] 方法引用 ✨

Java [ 基础 ] 方法引用 ✨

✨探索Java基础✨ Java基础&#xff1a;方法引用 方法引用是Java 8中引入的一种新特性&#xff0c;它使得代码更加简洁和易读。方法引用提供了一种可以直接引用已有方法作为Lambda表达式的替代方案。本文将深入介绍方法引用的基本概念、使用方法、具体实例及其在实际开发中的…
阅读更多...
Open3D 点云的旋转与平移

Open3D 点云的旋转与平移

目录 一、概述 1.1旋转 1.2平移 二、代码实现 2.1实现旋转 2.2实现平移 2.3组合变换 三、实现效果 3.1原始点云 3.2变换后点云 一、概述 在Open3D中&#xff0c;点云的旋转和平移是通过几何变换来实现的。几何变换可以应用于点云对象&#xff0c;使其在空间中移动或旋…
阅读更多...
MobPush iOS端海外推送最佳实现

MobPush iOS端海外推送最佳实现

推送注册 在AppDelegate里进行SDK初始化&#xff08;也可以在Info.plist文件中进行AppKey&#xff0c;AppSecret的配置&#xff09;并对通知功能进行注册以及设置推送的环境和切换海外服务器等&#xff0c;参考如下步骤代码&#xff1a; <span style"background-colo…
阅读更多...
叮！云原生虚拟数仓 PieCloudDB Database 动态包裹已送达

叮！云原生虚拟数仓 PieCloudDB Database 动态包裹已送达

第一部分 PieCloudDB Database 最新动态 支持动态配置查询簇 PieCloudDB 最新内核版本 v2.14.0 新增动态配置查询簇功能。PieCloudDB 动态配置查询簇功能实现可伸缩的并行化查询&#xff0c;可提升单个查询并行使用底层资源的能力&#xff0c;同时加快查询响应速度。 动态配…
阅读更多...
redis,memcached,nginx网络组件

redis,memcached,nginx网络组件

课程目标&#xff1a; 1.网络模块要处理哪些事情 2.reactor是怎么处理这些事情的 3.reactor怎么封装 4.网络模块与业务逻辑的关系 5.怎么优化reactor? io函数 函数调用 都有两个作用&#xff1a;io检测 是否就绪 io操作 1. int clientfd accept(listenfd, &addr, &l…
阅读更多...
navicat Lite 版

navicat Lite 版

navicat Lite 版&#xff1a; Navicat 出了一个 Navicat Premium 的Lite版。 官方现在链接&#xff1a;https://www.navicat.com.cn/download/navicat-premium-lite#windows 从官网可以看到现在能够下载最新版本 17&#xff0c;支持各种平台
阅读更多...
最新文章

Copyright @ 2022~2023