当前位置: 首页 > news >正文

debug mccl 02 —— 环境搭建及初步调试

1, 搭建nccl 调试环境

下载 nccl 源代码

 git clone --recursive https://github.com/NVIDIA/nccl.git


只debug host代码,故将设备代码的编译标志改成 -O3

 

(base) hipper@hipper-G21:~/let_debug_nccl/nccl$ git diff 
diff --git a/makefiles/common.mk b/makefiles/common.mk
index a037cf3..ee2aa8e 100644
--- a/makefiles/common.mk
+++ b/makefiles/common.mk
@@ -82,7 +82,8 @@ ifeq ($(DEBUG), 0)NVCUFLAGS += -O3CXXFLAGS  += -O3 -gelse
-NVCUFLAGS += -O0 -G -g
+#NVCUFLAGS += -O0 -G -g
+NVCUFLAGS += -O3CXXFLAGS  += -O0 -g -ggdb3endif

修改后变成如下:

nccl$ vim makefiles/common.mk

ifeq ($(DEBUG), 0)
NVCUFLAGS += -O3
CXXFLAGS  += -O3 -g
else
#NVCUFLAGS += -O0 -G -g
NVCUFLAGS += -O3
CXXFLAGS  += -O0 -g -ggdb3
endif

构建 nccl shared library:

机器上是几张sm_85 的卡,故:

$ cd nccl
$ make -j src.build  DEBUG=1       NVCC_GENCODE="-gencode=arch=compute_80,code=sm_80"

到此即可,不需要安装nccl,直接过来使用即可;

2, 创建调试APP


在nccl所在的目录中创建app文件夹:

$ mkdir app$ cd app$ vim sp_md_nccl.cpp$ vim Makefile

sp_md_nccl.cpp:

#include <stdlib.h>
#include <stdio.h>
#include "cuda_runtime.h"
#include "nccl.h"
#include <time.h>
#include <sys/time.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)void  get_seed(long long &seed)
{struct timeval tv;gettimeofday(&tv, NULL);seed = (long long)tv.tv_sec * 1000*1000 + tv.tv_usec;//only second and usecond;printf("useconds:%lld\n", seed);
}void  init_vector(float* A, int n)
{long long seed = 0;get_seed(seed);srand(seed);for(int i=0; i<n; i++){A[i] = (rand()%100)/100.0f;}
}void print_vector(float* A, float size)
{for(int i=0; i<size; i++)printf("%.2f ", A[i]);printf("\n");
}void vector_add_vector(float* sum, float* A, int n)
{for(int i=0; i<n; i++){sum[i] += A[i];}
}int main(int argc, char* argv[])
{ncclComm_t comms[4];printf("ncclComm_t is a pointer type, sizeof(ncclComm_t)=%lu\n", sizeof(ncclComm_t));//managing 4 devices//int nDev = 4;int nDev = 2;//int size = 32*1024*1024;int size = 16*16;int devs[4] = { 0, 1, 2, 3 };float** sendbuff_host = (float**)malloc(nDev * sizeof(float*));float** recvbuff_host = (float**)malloc(nDev * sizeof(float*));for(int dev=0; dev<nDev; dev++){sendbuff_host[dev] = (float*)malloc(size*sizeof(float));recvbuff_host[dev] = (float*)malloc(size*sizeof(float));init_vector(sendbuff_host[dev], size);init_vector(recvbuff_host[dev], size);}//sigma(sendbuff_host[i]); i = 0, 1, ..., nDev-1float* result = (float*)malloc(size*sizeof(float));memset(result, 0, size*sizeof(float));for(int dev=0; dev<nDev; dev++){vector_add_vector(result, sendbuff_host[dev], size);printf("sendbuff_host[%d]=\n", dev);print_vector(sendbuff_host[dev], size);}printf("result=\n");print_vector(result, size);//allocating and initializing device buffersfloat** 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(sendbuff + i, size * sizeof(float)));CUDACHECK(cudaMalloc(recvbuff + i, size * sizeof(float)));CUDACHECK(cudaMemcpy(sendbuff[i], sendbuff_host[i], size*sizeof(float), cudaMemcpyHostToDevice));CUDACHECK(cudaMemcpy(recvbuff[i], recvbuff_host[i], size*sizeof(float), cudaMemcpyHostToDevice));CUDACHECK(cudaStreamCreate(s+i));}//initializing NCCLNCCLCHECK(ncclCommInitAll(comms, nDev, devs));//calling NCCL communication API. Group API is required when using//multiple devices per threadNCCLCHECK(ncclGroupStart());printf("blocked ncclAllReduce will be calleded\n");fflush(stdout);for (int i = 0; i < nDev; ++i)NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, ncclFloat, ncclSum, comms[i], s[i]));printf("blocked ncclAllReduce is calleded nDev =%d\n", nDev);fflush(stdout);NCCLCHECK(ncclGroupEnd());//synchronizing on CUDA streams to wait for completion of NCCL operationfor (int i = 0; i < nDev; ++i) {CUDACHECK(cudaSetDevice(i));CUDACHECK(cudaStreamSynchronize(s[i]));}for (int i = 0; i < nDev; ++i) {CUDACHECK(cudaSetDevice(i));CUDACHECK(cudaMemcpy(recvbuff_host[i], recvbuff[i], size*sizeof(float), cudaMemcpyDeviceToHost));}for (int i = 0; i < nDev; ++i) {CUDACHECK(cudaSetDevice(i));CUDACHECK(cudaStreamSynchronize(s[i]));}for(int i=0; i<nDev; i++) {printf("recvbuff_dev2host[%d]=\n", i);print_vector(recvbuff_host[i], size);}//free device buffersfor (int i = 0; i < nDev; ++i) {CUDACHECK(cudaSetDevice(i));CUDACHECK(cudaFree(sendbuff[i]));CUDACHECK(cudaFree(recvbuff[i]));}//finalizing NCCLfor(int i = 0; i < nDev; ++i)ncclCommDestroy(comms[i]);printf("Success \n");return 0;
}

Makefile:


INC := -I /usr/local/cuda/include -I ../nccl/build/include
LD_FLAGS := -L ../nccl/build/lib -lnccl -L /usr/local/cuda/lib64 -lcudartEXE := singleProc_multiDev_ncclall: $(EXE)%: %.cppg++ -g -ggdb3  $<  -o  $@  $(INC)  $(LD_FLAGS).PHONY: clean
clean: -rm -rf $(EXE)

export LD_LIBRARY_PATH=../nccl/build/lib

3, 开始调试


$ cuda-gdb sp_md_nccl(cuda-gdb) start (cuda-gdb) rbreak doLauches(cuda-gdb) c(cuda-gdb) p ncclGroupCommHead->tasks.collQueue.head->op 

初步实现了可debug的效果:

现在想要搞清楚在程序调用 ncclAllReduce(..., ncclSum,  ... ) 后,是如何映射到 cudaLaunchKernel调用到了正确的 cuda kernel 函数的。

在doLaunches函数中,作如下debug动作:

查看 doLaunches(ncclComm*) 的函数参数,即,gropu.cc中的变量:ncclGroupCommHead的某个成员的成员的值:op

其结果如下:

(cuda-gdb) p ncclGroupCommHead                           
$5 = (ncclComm *) 0x5555563231e0
(cuda-gdb) p ncclGroupCommHead->tasks.collQueue.head->op 
$6 = {op = ncclDevSum, proxyOp = ncclSum, scalarArgIsPtr = false, scalarArg = 256}
(cuda-gdb) 

不过这依然只停留在了 ncclSum的这个枚举类型上,还没锁定对应的cudaKernel。

接下来继续努力 ...

相关文章:

debug mccl 02 —— 环境搭建及初步调试

1, 搭建nccl 调试环境 下载 nccl 源代码 git clone --recursive https://github.com/NVIDIA/nccl.git 只debug host代码&#xff0c;故将设备代码的编译标志改成 -O3 (base) hipperhipper-G21:~/let_debug_nccl/nccl$ git diff diff --git a/makefiles/common.mk b/makefiles/…...

ros python 接收GPS RTK 串口消息再转发 ros 主题消息

代码是一个ROS(Robot Operating System)节点,用于从GPS设备读取RTK(实时动态)数据并通过ROS主题发布。 步骤: 导入必要的模块: rospy 是ROS的Python库,用于ROS的节点、发布者和订阅者。serial 用于串行通信。NavSatFix 和 NavSatStatus 是从GPS接收的NMEA 0183标准消息…...

2024年网络安全竞赛-页面信息发现任务解析

页面信息发现任务说明: 服务器场景:win20230305(关闭链接)在渗透机中对服务器信息收集,将获取到的服务器网站端口作为Flag值提交;访问服务器网站页面,找到主页面中的Flag值信息,将Flag值提交;访问服务器网站页面,找到主页面中的脚本信息,并将Flag值提交;访问服务器…...

ARM DMA使用整理

Direct Memory Access&#xff0c; 直接存储访问。同SPI,IIC,USART等一样&#xff0c;属于MCU的一个外设&#xff0c;用于在不需要MCU介入的情况下进行数据传输。可以将数据从外设传输到flash&#xff0c;也可以将数据从flash传输到外设&#xff0c;或者flash内部数据移动。 它…...

移动通信原理与关键技术学习(第四代蜂窝移动通信系统)

前言&#xff1a;LTE 标准于2008 年底完成了第一个版本3GPP Release 8的制定工作。另一方面&#xff0c;ITU 于2007 年召开了世界无线电会议WRC07&#xff0c;开始了B3G 频谱的分配&#xff0c;并于2008 年完成了IMT-2000&#xff08;即3G&#xff09;系统的演进——IMT-Advanc…...

光明源@智慧厕所技术:优化生活,提升卫生舒适度

在当今数字科技飞速发展的时代&#xff0c;我们的日常生活正在经历一场革命&#xff0c;而这场革命的其中一个前沿领域就是智慧厕所技术。这项技术不仅仅是对传统卫生间的一次升级&#xff0c;更是对我们生活品质的全方位提升。从智能感应到数据分析&#xff0c;从环保设计到舒…...

【Bootstrap学习 day13】

Bootstrap5 下拉菜单 下拉菜单通常用于导航标题内&#xff0c;在用户鼠标悬停或单击触发元素时显示相关链接列表。 基础的下拉列表 <div class"dropdown"><button type"button" class"btn btn-primary dropdown-toggle" data-bs-toggl…...

Shell:常用命令之dirname与basename

一、介绍 1、dirname命令用于去除文件名中的非目录部分&#xff0c;删除最后一个“\”后面的路径&#xff0c;显示父目录。 语法&#xff1a;dirname [选项] 参数 2、basename命令用于打印目录或者文件的基本名称&#xff0c;显示最后的目录名或文件名。 语法&#xff1a;basen…...

Linux-v4l2框架

框架图 从上图不难看出&#xff0c;v4l2_device作为顶层管理者&#xff0c;一方面通过嵌入到一个video_device中&#xff0c;暴露video设备节点给用户空间进行控制&#xff1b;另一方面&#xff0c;video_device内部会创建一个media_entity作为在media controller中的抽象体&a…...

VPC网络架构下的网络上数据采集

介绍 想象这样一个场景&#xff0c;一开始在公司里&#xff0c;所有的部门的物理机、POD都在一个经典网络内&#xff0c;它们可以通过 IP 访问彼此&#xff0c;没有任何限制。因此有很多系统基于此设计了很多点对点 IP 直连的访问&#xff0c;比如中心控制服务器 S 会主动访问物…...

模拟算法(模拟算法 == 依葫芦画瓢)万字

模拟算法 基本思想引入算法题替换所有的问号提莫攻击Z字形变换外观数列数青蛙 基本思想 模拟算法 依葫芦画瓢解题思维要么通俗易懂&#xff0c;要么就是找规律&#xff0c;主要难度在于将思路转换为代码。 特点&#xff1a;相对于其他算法思维&#xff0c;思路比较简单&#x…...

QtApplets-SystemInfo

QtApplets-SystemInfo ​ 今天是2024年1月3日09:18:44&#xff0c;这也是2024年的第一篇博客&#xff0c;今天我们主要两件事&#xff0c;第一件&#xff0c;获取系统CPU使用率&#xff0c;第二件&#xff0c;获取系统内存使用情况。 ​ 这里因为写博客的这个本本的环境配置不…...

vue3防抖函数封装与使用,以指令的形式使用

utils/debounce.js /*** 防抖函数* param {*} fn 函数* param {*} delay 暂停时间* returns */ export function debounce(fn, delay 500) {let timer nullreturn function (...args) {// console.log(arguments);// const args Array.from(arguments)if (timer) {clearTim…...

Hive学习(13)lag和lead函数取偏移量

hive里面lag函数 在数据处理和分析中&#xff0c;窗口函数是一种重要的技术&#xff0c;用于在数据集中执行聚合和分析操作。Hive作为一种大数据处理框架&#xff0c;也提供了窗口函数的支持。在Hive中&#xff0c;Lag函数是一种常用的窗口函数&#xff0c;可以用于计算前一行…...

Centos Unable to verify the graphical display setup

ERROR: Unable to verify the graphical display setup. 在Linux下安装Oracle时 运行 ./runInstaller 报错 ERROR: Unable to verify the graphical display setup. This application requires X display. Make sure that xdpyinfo exist under PATH variable. No X11 DISPL…...

Java 说一下 synchronized 底层实现原理?

Java 说一下 synchronized 底层实现原理&#xff1f; synchronized 是 Java 中用于实现同步的关键字&#xff0c;它保证了多个线程对共享资源的互斥访问。底层实现涉及到对象头的 Mark Word 和锁升级过程。 synchronized 可以用于方法上或代码块上&#xff0c;分别对应于方法…...

nginx访问路径匹配方法

目录 一&#xff1a;匹配方法 二&#xff1a;location使用: 三&#xff1a;rewrite使用 一&#xff1a;匹配方法 location和rewrite是两个用于处理请求的重要模块&#xff0c;它们都可以根据请求的路径进行匹配和处理。 二&#xff1a;location使用: 1&#xff1a;简单匹配…...

偌依 项目部署及上线步骤

准备实验环境&#xff0c;准备3台机器 1.作为前端服务器&#xff0c;mysql,redis服务器--同时临时作为代码打包服务器 192.168.2.65 nginx-server 2.作为后端服务器 192.168.2.66 java-server-1 192.168.2.67 java-server-2 安装nginx/mysql #安装nginx [rootweb-nginx ~]…...

PHP特性知识点扫盲 - 上篇

概述 之前在分析thinkphp源码的时候&#xff0c;对依赖注入等等php高级的特性一直想做一个梳理和总结&#xff0c;一直没有时间&#xff0c;好不容易抽一点时间对技术的盲点做一个扫盲和总结。 特性 1.命名空间 命名空间是在PHP5.3中引入&#xff0c;是一个很重要的工具&am…...

Docker一键极速安装Nacos,并配置数据库!

1 部署方式 1.1 DockerHub javaedgeJavaEdgedeMac-mini ~ % docker run --name nacos \ -e MODEstandalone \ -e JVM_XMS128m \ -e JVM_XMX128m \ -e JVM_XMN64m \ -e JVM_MS64m \ -e JVM_MMS64m \ -p 8848:8848 \ -d nacos/nacos-server:v2.2.3 a624c64a1a25ad2d15908a67316d…...

多模态2025:技术路线“神仙打架”,视频生成冲上云霄

文&#xff5c;魏琳华 编&#xff5c;王一粟 一场大会&#xff0c;聚集了中国多模态大模型的“半壁江山”。 智源大会2025为期两天的论坛中&#xff0c;汇集了学界、创业公司和大厂等三方的热门选手&#xff0c;关于多模态的集中讨论达到了前所未有的热度。其中&#xff0c;…...

基于距离变化能量开销动态调整的WSN低功耗拓扑控制开销算法matlab仿真

目录 1.程序功能描述 2.测试软件版本以及运行结果展示 3.核心程序 4.算法仿真参数 5.算法理论概述 6.参考文献 7.完整程序 1.程序功能描述 通过动态调整节点通信的能量开销&#xff0c;平衡网络负载&#xff0c;延长WSN生命周期。具体通过建立基于距离的能量消耗模型&am…...

Zustand 状态管理库:极简而强大的解决方案

Zustand 是一个轻量级、快速和可扩展的状态管理库&#xff0c;特别适合 React 应用。它以简洁的 API 和高效的性能解决了 Redux 等状态管理方案中的繁琐问题。 核心优势对比 基本使用指南 1. 创建 Store // store.js import create from zustandconst useStore create((set)…...

day52 ResNet18 CBAM

在深度学习的旅程中&#xff0c;我们不断探索如何提升模型的性能。今天&#xff0c;我将分享我在 ResNet18 模型中插入 CBAM&#xff08;Convolutional Block Attention Module&#xff09;模块&#xff0c;并采用分阶段微调策略的实践过程。通过这个过程&#xff0c;我不仅提升…...

Axios请求超时重发机制

Axios 超时重新请求实现方案 在 Axios 中实现超时重新请求可以通过以下几种方式&#xff1a; 1. 使用拦截器实现自动重试 import axios from axios;// 创建axios实例 const instance axios.create();// 设置超时时间 instance.defaults.timeout 5000;// 最大重试次数 cons…...

【JavaWeb】Docker项目部署

引言 之前学习了Linux操作系统的常见命令&#xff0c;在Linux上安装软件&#xff0c;以及如何在Linux上部署一个单体项目&#xff0c;大多数同学都会有相同的感受&#xff0c;那就是麻烦。 核心体现在三点&#xff1a; 命令太多了&#xff0c;记不住 软件安装包名字复杂&…...

Web 架构之 CDN 加速原理与落地实践

文章目录 一、思维导图二、正文内容&#xff08;一&#xff09;CDN 基础概念1. 定义2. 组成部分 &#xff08;二&#xff09;CDN 加速原理1. 请求路由2. 内容缓存3. 内容更新 &#xff08;三&#xff09;CDN 落地实践1. 选择 CDN 服务商2. 配置 CDN3. 集成到 Web 架构 &#xf…...

管理学院权限管理系统开发总结

文章目录 &#x1f393; 管理学院权限管理系统开发总结 - 现代化Web应用实践之路&#x1f4dd; 项目概述&#x1f3d7;️ 技术架构设计后端技术栈前端技术栈 &#x1f4a1; 核心功能特性1. 用户管理模块2. 权限管理系统3. 统计报表功能4. 用户体验优化 &#x1f5c4;️ 数据库设…...

七、数据库的完整性

七、数据库的完整性 主要内容 7.1 数据库的完整性概述 7.2 实体完整性 7.3 参照完整性 7.4 用户定义的完整性 7.5 触发器 7.6 SQL Server中数据库完整性的实现 7.7 小结 7.1 数据库的完整性概述 数据库完整性的含义 正确性 指数据的合法性 有效性 指数据是否属于所定…...

苹果AI眼镜:从“工具”到“社交姿态”的范式革命——重新定义AI交互入口的未来机会

在2025年的AI硬件浪潮中,苹果AI眼镜(Apple Glasses)正在引发一场关于“人机交互形态”的深度思考。它并非简单地替代AirPods或Apple Watch,而是开辟了一个全新的、日常可接受的AI入口。其核心价值不在于功能的堆叠,而在于如何通过形态设计打破社交壁垒,成为用户“全天佩戴…...