{"id":40335,"date":"2025-05-28T13:43:56","date_gmt":"2025-05-28T05:43:56","guid":{"rendered":"https:\/\/www.wsisp.com\/helps\/40335.html"},"modified":"2025-05-28T13:43:56","modified_gmt":"2025-05-28T05:43:56","slug":"%e6%91%a9%e5%b0%94%e7%ba%bf%e7%a8%8bs4000%e5%9b%bd%e4%ba%a7%e4%bf%a1%e5%88%9b%e8%ae%a1%e7%ae%97%e5%8d%a1%e6%80%a7%e8%83%bd%e5%ae%9e%e6%88%98-pytorch%e8%bd%ac%e8%af%91%ef%bc%8c","status":"publish","type":"post","link":"https:\/\/www.wsisp.com\/helps\/40335.html","title":{"rendered":"\u6469\u5c14\u7ebf\u7a0bS4000\u56fd\u4ea7\u4fe1\u521b\u8ba1\u7b97\u5361\u6027\u80fd\u5b9e\u6218\u2014\u2014Pytorch\u8f6c\u8bd1\uff0c\u591a\u5361P2P\u901a\u4fe1\u4e0eMUSA\u7f16\u7a0b"},"content":{"rendered":"

\u7b80\u4ecb<\/h2>\n

MTT S4000 \u662f\u57fa\u4e8e\u6469\u5c14\u7ebf\u7a0b\u66f2\u9662 GPU \u67b6\u6784\u6253\u9020\u7684\u5168\u529f\u80fd\u5143\u8ba1\u7b97\u5361，\u4e3a\u5343\u4ebf\u89c4\u6a21\u5927\u8bed\u8a00\u6a21\u578b\u7684\u8bad\u7ec3\u3001\u5fae\u8c03\u548c\u63a8\u7406\u8fdb\u884c\u4e86\u5b9a\u5236\u4f18\u5316，\u7ed3\u5408\u5148\u8fdb\u7684\u56fe\u5f62\u6e32\u67d3\u80fd\u529b\u3001\u89c6\u9891\u7f16\u89e3\u7801\u80fd\u529b\u548c\u8d85\u9ad8\u6e05 8K HDR \u663e\u793a\u80fd\u529b，\u52a9\u529b\u4eba\u5de5\u667a\u80fd\u3001\u56fe\u5f62\u6e32\u67d3\u3001\u591a\u5a92\u4f53\u3001\u79d1\u5b66\u8ba1\u7b97\u4e0e\u7269\u7406\u4eff\u771f\u7b49\u590d\u5408\u5e94\u7528\u573a\u666f\u7684\u8ba1\u7b97\u52a0\u901f\u3002<\/p>\n

MTT S4000 \u5168\u9762\u652f\u6301\u5927\u8bed\u8a00\u6a21\u578b\u7684\u9884\u8bad\u7ec3\u3001\u5fae\u8c03\u548c\u63a8\u7406\u670d\u52a1，MUSA \u8f6f\u4ef6\u6808\u4e13\u95e8\u9488\u5bf9\u5927\u89c4\u6a21\u96c6\u7fa4\u7684\u5206\u5e03\u5f0f\u8ba1\u7b97\u6027\u80fd\u8fdb\u884c\u4e86\u4f18\u5316，\u9002\u914d\u4e3b\u6d41\u5206\u5e03\u5f0f\u8ba1\u7b97\u52a0\u901f\u6846\u67b6， \u5305\u62ec DeepSpeed， Colossal AI，Megatron \u7b49，\u652f\u6301\u5343\u4ebf\u53c2\u6570\u5927\u8bed\u8a00\u6a21\u578b\u7684\u7a33\u5b9a\u9884\u8bad\u7ec3\u3002<\/p>\n

\u5b98\u65b9\u53c2\u6570\u5982\u4e0b<\/p>\n

\"\"<\/p>\n

\u8fd0\u884c\u73af\u5883<\/h2>\n

\u672c\u6b21\u8fd0\u884c\u73af\u5883\u4e3aAutoDL\u4e91\u4e2d\u7684\u955c\u50cf\u73af\u5883，\u7cfb\u7edf\u73af\u5883\u5982\u4e0b<\/p>\n

\"\"<\/p>\n

\u5e38\u7528\u547d\u4ee4<\/h2>\n

\u663e\u5361\u8fd0\u884c\u72b6\u6001<\/h3>\n

\u8f93\u5165\u5982\u4e0b\u547d\u4ee4<\/p>\n

mthreads-gmi<\/p>\n

\u5373\u53ef\u67e5\u770b\u5f53\u524d\u663e\u5361\u8fd0\u884c\u72b6\u6001<\/p>\n

\"\"<\/p>\n

\u67e5\u770b\u5f53\u524dGPU\u8be6\u7ec6\u4fe1\u606f<\/h3>\n

\u8f93\u5165<\/p>\n

musaInfo <\/p>\n

\u5373\u53ef<\/p>\n

\"\"<\/p>\n

\u67e5\u770b\u5f53\u524d\u8fd0\u884c\u73af\u5883\u7248\u672c<\/h3>\n

\u8f93\u5165<\/p>\n

musa_version_query<\/p>\n

\u5373\u53ef\u67e5\u770b\u5f53\u524d\u8fd0\u884c\u73af\u5883\u7248\u672c<\/p>\n

\"\"<\/p>\n

Pytorch\u90e8\u5206<\/h2>\n

\u8f6c\u4e49<\/h3>\n

\u6839\u636e\u5b98\u7f51\u4ecb\u7ecd，\u5bf9\u4e8epytorch\u4ee3\u7801，\u53ea\u9700\u8981\u6b63\u786eimport torch_musa\u7684\u62d3\u5c55\u63d2\u4ef6，\u5e76\u4e14\u5c06\u4ee3\u7801\u4e2d\u7684\u6240\u6709cuda->musa，\u5c06\u6240\u6709\u7684nccl->mccl\u5373\u53ef\u3002<\/p>\n

\"\"<\/p>\n

\u5b9e\u6d4b<\/h3>\n

\u4f5c\u8005\u4f7f\u7528\u8c46\u5305\u968f\u673a\u751f\u6210\u4e86\u4e00\u4e2a\u6d4b\u8bd5allreduce\u7684pytorch\u4ee3\u7801，\u4ee3\u7801\u5982\u4e0b，\u5728\u7ecf\u8fc7\u4e0a\u8ff0\u8f6c\u8bd1\u540e\u80fd\u6b63\u5e38\u8fd0\u884c<\/p>\n

import os
\nimport time
\nimport argparse
\nimport torch
\nimport torch_musa
\nimport torch.distributed as dist
\nfrom torch.nn.parallel import DistributedDataParallel as DDP<\/p>\n

def setup(rank, world_size):
\n os.environ['MASTER_ADDR'] = 'localhost'
\n os.environ['MASTER_PORT'] = '12355'<\/p>\n

# \u521d\u59cb\u5316MUSA\u5206\u5e03\u5f0f\u73af\u5883
\n dist.init_process_group("mccl", rank=rank, world_size=world_size)
\n torch.musa.set_device(rank)<\/p>\n

def cleanup():
\n dist.destroy_process_group()<\/p>\n

def run_benchmark(rank, world_size, sizes, num_iters=100, warmup=20):
\n setup(rank, world_size)<\/p>\n

for size in sizes:
\n # \u521b\u5efa\u968f\u673a\u5f20\u91cf（\u4f7f\u7528MUSA\u8bbe\u5907）
\n tensor = torch.rand(size, device=f'musa:{rank}')<\/p>\n

# \u9884\u70ed
\n for _ in range(warmup):
\n dist.all_reduce(tensor)
\n torch.musa.synchronize()<\/p>\n

# \u6d4b\u91cf\u65f6\u95f4
\n start_time = time.time()
\n for _ in range(num_iters):
\n dist.all_reduce(tensor)
\n torch.musa.synchronize()
\n end_time = time.time()<\/p>\n

# \u8ba1\u7b97\u7edf\u8ba1\u4fe1\u606f
\n total_time = end_time – start_time
\n avg_time = total_time \/ num_iters
\n size_mb = size * 4 \/ (1024 * 1024) # float32\u662f4\u5b57\u8282
\n bandwidth = (size_mb * world_size) \/ avg_time # MB\/s<\/p>\n

if rank == 0:
\n print(f"\u5f20\u91cf\u5927\u5c0f: {size:,} \u5143\u7d20 ({size_mb:.2f} MB)")
\n print(f"\u5e73\u5747\u8017\u65f6: {avg_time * 1000:.2f} ms")
\n print(f"\u5e26\u5bbd: {bandwidth \/ 1024:.2f} GB\/s")
\n print("-" * 50)<\/p>\n

cleanup()<\/p>\n

def main():
\n parser = argparse.ArgumentParser()
\n parser.add_argument('–sizes', type=int, nargs='+',
\n default=[1000, 10000, 100000, 1000000, 10000000, 100000000],
\n metavar='N',
\n help='\u6d4b\u8bd5\u7684\u5f20\u91cf\u5927\u5c0f\u5217\u8868')
\n parser.add_argument('–num-iters', type=int, default=100,
\n help='\u6bcf\u4e2a\u5927\u5c0f\u7684\u8fed\u4ee3\u6b21\u6570')
\n parser.add_argument('–warmup', type=int, default=20,
\n help='\u9884\u70ed\u8fed\u4ee3\u6b21\u6570')
\n args = parser.parse_args()<\/p>\n

world_size = torch.musa.device_count()
\n if world_size != 4:
\n raise ValueError("\u6b64\u811a\u672c\u9700\u89814\u4e2aMUSA GPU，\u4f46\u53d1\u73b0 {} \u4e2aGPU".format(world_size))<\/p>\n

import torch.multiprocessing as mp
\n mp.spawn(run_benchmark,
\n args=(world_size, args.sizes, args.num_iters, args.warmup),
\n nprocs=world_size,
\n join=True)<\/p>\n

if __name__ == "__main__":
\n main()<\/p>\n

\"\"\"\"<\/p>\n

MUSA\u7f16\u7a0b<\/h2>\n

p2p\u901a\u4fe1\u90e8\u5206<\/h3>\n

\u4ee3\u7801\u53c2\u8003<\/h4>\n

\u7b14\u8005\u6309\u7167\u82f1\u4f1f\u8fbecudasamples\u4ed3\u5e93\u4e2d\u7684p2pbandwidthtest \u4ee3\u7801，cuda-samples\/Samples\/5_Domain_Specific\/p2pBandwidthLatencyTest at master \u00b7 NVIDIA\/cuda-samples \u00b7 GitHub<\/p>\n

\u5e76\u4e14\u53c2\u8003\u76f8\u5e94\u7684musa event api\u4e0emublasapi<\/p>\n

https:\/\/docs.mthreads.com\/musa-sdk\/musa-sdk-doc-online\/api\/mcc_um.zh-CN<\/p>\n

\u7f16\u5199\u4e86\u4e00\u4e2a\u9002\u7528\u4e8e\u6469\u5c14\u7ebf\u7a0b\u7684p2p\u901a\u4fe1\u68c0\u6d4b\u9a8c\u8bc1\u7a0b\u5e8f<\/p>\n

\u4ee3\u7801\u90e8\u5206<\/h4>\n

#include <cstdio>
\n#include <vector>
\n#include <musa_runtime.h> \/\/ \u5047\u8bbe MUSA \u5934\u6587\u4ef6<\/p>\n

using namespace std;<\/p>\n

const char *sSampleName = "P2P (Peer-to-Peer) GPU Bandwidth Latency Test";<\/p>\n

typedef enum {
\n P2P_WRITE = 0,
\n P2P_READ = 1,
\n} P2PDataTransfer;<\/p>\n

typedef enum {
\n CE = 0,
\n SM = 1,
\n} P2PEngine;<\/p>\n

P2PEngine p2p_mechanism = CE; \/\/ \u9ed8\u8ba4\u4f7f\u7528 Copy Engine<\/p>\n

\/\/ \u9519\u8bef\u68c0\u67e5\u5b8f
\n#define musaCheckError() \\\\
\n { \\\\
\n musaError_t e = musaGetLastError(); \\\\
\n if (e != musaSuccess) { \\\\
\n printf("MUSA failure %s:%d: '%s'\\\\n", __FILE__, __LINE__, musaGetErrorString(e)); \\\\
\n exit(EXIT_FAILURE); \\\\
\n } \\\\
\n }<\/p>\n

\/\/ \u5ef6\u8fdf\u5185\u6838
\n__global__ void delay(volatile int *flag, unsigned long long timeout_clocks = 10000000)
\n{
\n \/\/ \u7b49\u5f85\u5e94\u7528\u7a0b\u5e8f\u901a\u77e5\u6211\u4eec\u5b83\u5df2\u7ecf\u5b8c\u6210\u4e86\u5b9e\u9a8c\u7684\u6392\u961f，\u6216\u8005\u8d85\u65f6\u5e76\u9000\u51fa，\u5141\u8bb8\u5e94\u7528\u7a0b\u5e8f\u7ee7\u7eed\u6267\u884c
\n long long int start_clock, sample_clock;
\n start_clock = clock64();<\/p>\n

while (!*flag) {
\n sample_clock = clock64();<\/p>\n

if (sample_clock – start_clock > timeout_clocks) {
\n break;
\n }
\n }
\n}<\/p>\n

\/\/ P2P \u590d\u5236\u5185\u6838
\n__global__ void copyp2p(int4 *__restrict__ dest, const int4 *__restrict__ src, size_t num_elems) {
\n size_t globalId = blockIdx.x * blockDim.x + threadIdx.x;
\n size_t gridSize = blockDim.x * gridDim.x;<\/p>\n

#pragma unroll 5 \/\/ \u79fb\u9664\u62ec\u53f7
\n for (size_t i = globalId; i < num_elems; i += gridSize) {
\n dest[i] = src[i];
\n }
\n}<\/p>\n

\/\/ \u6253\u5370\u5e2e\u52a9\u4fe1\u606f
\nvoid printHelp(void) {
\n printf("Usage: p2pBandwidthLatencyTest [OPTION]…\\\\n");
\n printf("Tests bandwidth\/latency of GPU pairs using P2P and without P2P\\\\n");
\n printf("\\\\n");
\n printf("Options:\\\\n");
\n printf("–help\\\\t\\\\tDisplay this help menu\\\\n");
\n printf("–p2p_read\\\\tUse P2P reads for data transfers between GPU pairs\\\\n");
\n printf("–sm_copy\\\\tUse SM intiated p2p transfers instead of Copy Engine\\\\n");
\n printf("–numElems=<NUM_OF_INT_ELEMS> Number of integer elements for p2p copy\\\\n");
\n}<\/p>\n

\/\/ \u68c0\u67e5P2P\u8bbf\u95ee
\nvoid checkP2Paccess(int numGPUs) {
\n for (int i = 0; i < numGPUs; i++) {
\n musaSetDevice(i);
\n musaCheckError();<\/p>\n

for (int j = 0; j < numGPUs; j++) {
\n if (i != j) {
\n int access;
\n musaDeviceCanAccessPeer(&access, i, j);
\n musaCheckError();
\n printf("Device=%d %s Access Peer Device=%d\\\\n", i, access ? "CAN" : "CANNOT", j);
\n }
\n }
\n }
\n printf("\\\\n***NOTE: Devices without P2P access fall back to normal memcpy.\\\\n");
\n}<\/p>\n

\/\/ \u6267\u884cP2P\u590d\u5236
\nvoid performP2PCopy(int *dest, int destDevice, int *src, int srcDevice,
\n size_t num_elems, int repeat, bool p2paccess,
\n musaStream_t streamToRun) {
\n int blockSize, numBlocks;
\n musaOccupancyMaxPotentialBlockSize(&numBlocks, &blockSize, copyp2p);
\n musaCheckError();<\/p>\n

if (p2p_mechanism == SM && p2paccess) {
\n for (int r = 0; r < repeat; r++) {
\n copyp2p<<<numBlocks, blockSize, 0, streamToRun>>>((int4*)dest, (int4*)src, num_elems\/4);
\n }
\n } else {
\n for (int r = 0; r < repeat; r++) {
\n musaMemcpyPeerAsync(dest, destDevice, src, srcDevice,
\n sizeof(int)*num_elems, streamToRun);
\n musaCheckError();
\n }
\n }
\n}<\/p>\n

\/\/ \u8f93\u51fa\u5e26\u5bbd\u77e9\u9635
\nvoid outputBandwidthMatrix(int numElems, int numGPUs, bool p2p, P2PDataTransfer p2p_method) {
\n int repeat = 5;
\n volatile int *flag = NULL;
\n vector<int *> buffers(numGPUs);
\n vector<int *> buffersD2D(numGPUs);
\n vector<musaEvent_t> start(numGPUs);
\n vector<musaEvent_t> stop(numGPUs);
\n vector<musaStream_t> stream(numGPUs);<\/p>\n

musaHostAlloc((void **)&flag, sizeof(*flag), musaHostAllocPortable);
\n musaCheckError();<\/p>\n

for (int d = 0; d < numGPUs; d++) {
\n musaSetDevice(d);
\n musaStreamCreateWithFlags(&stream[d], musaStreamNonBlocking);
\n musaMalloc(&buffers[d], numElems * sizeof(int));
\n musaMemset(buffers[d], 0, numElems * sizeof(int));
\n musaMalloc(&buffersD2D[d], numElems * sizeof(int));
\n musaMemset(buffersD2D[d], 0, numElems * sizeof(int));
\n musaCheckError();
\n musaEventCreate(&start[d]);
\n musaCheckError();
\n musaEventCreate(&stop[d]);
\n musaCheckError();
\n }<\/p>\n

vector<double> bandwidthMatrix(numGPUs * numGPUs);<\/p>\n

for (int i = 0; i < numGPUs; i++) {
\n musaSetDevice(i);<\/p>\n

for (int j = 0; j < numGPUs; j++) {
\n int access = 0;
\n if (p2p) {
\n musaDeviceCanAccessPeer(&access, i, j);
\n if (access) {
\n musaDeviceEnablePeerAccess(j, 0);
\n musaCheckError();
\n musaSetDevice(j);
\n musaDeviceEnablePeerAccess(i, 0);
\n musaCheckError();
\n musaSetDevice(i);
\n musaCheckError();
\n }
\n }<\/p>\n

musaStreamSynchronize(stream[i]);
\n musaCheckError();<\/p>\n

\/\/ \u963b\u585e\u6d41，\u76f4\u5230\u6240\u6709\u5de5\u4f5c\u6392\u961f\u5b8c\u6210
\n *flag = 0;
\n delay<<<1, 1, 0, stream[i]>>>(flag);
\n musaCheckError();
\n musaEventRecord(start[i], stream[i]);
\n musaCheckError();<\/p>\n

if (i == j) {
\n performP2PCopy(buffers[i], i, buffersD2D[i], i, numElems, repeat, access, stream[i]);
\n }
\n else {
\n if (p2p_method == P2P_WRITE) {
\n performP2PCopy(buffers[j], j, buffers[i], i, numElems, repeat, access, stream[i]);
\n }
\n else {
\n performP2PCopy(buffers[i], i, buffers[j], j, numElems, repeat, access, stream[i]);
\n }
\n }<\/p>\n

musaEventRecord(stop[i], stream[i]);
\n musaCheckError();<\/p>\n

\/\/ \u91ca\u653e\u6392\u961f\u7684\u4e8b\u4ef6
\n *flag = 1;
\n musaStreamSynchronize(stream[i]);
\n musaCheckError();<\/p>\n

float time_ms;
\n musaEventElapsedTime(&time_ms, start[i], stop[i]);
\n double time_s = time_ms \/ 1e3;<\/p>\n

double gb = numElems * sizeof(int) * repeat \/ (double)1e9;
\n if (i == j) {
\n gb *= 2;
\n }
\n bandwidthMatrix[i * numGPUs + j] = gb \/ time_s;
\n if (p2p && access) {
\n musaDeviceDisablePeerAccess(j);
\n musaSetDevice(j);
\n musaDeviceDisablePeerAccess(i);
\n musaSetDevice(i);
\n musaCheckError();
\n }
\n }
\n }<\/p>\n

printf(" D\\\\\\\\D");
\n for (int j = 0; j < numGPUs; j++) {
\n printf("%6d ", j);
\n }
\n printf("\\\\n");<\/p>\n

for (int i = 0; i < numGPUs; i++) {
\n printf("%6d ", i);
\n for (int j = 0; j < numGPUs; j++) {
\n printf("%6.02f ", bandwidthMatrix[i * numGPUs + j]);
\n }
\n printf("\\\\n");
\n }<\/p>\n

for (int d = 0; d < numGPUs; d++) {
\n musaSetDevice(d);
\n musaFree(buffers[d]);
\n musaFree(buffersD2D[d]);
\n musaCheckError();
\n musaEventDestroy(start[d]);
\n musaCheckError();
\n musaEventDestroy(stop[d]);
\n musaCheckError();
\n musaStreamDestroy(stream[d]);
\n musaCheckError();
\n }<\/p>\n

musaFreeHost((void *)flag);
\n musaCheckError();
\n}<\/p>\n

\/\/ \u8f93\u51fa\u53cc\u5411\u5e26\u5bbd\u77e9\u9635
\nvoid outputBidirectionalBandwidthMatrix(int numElems, int numGPUs, bool p2p) {
\n int repeat = 5;
\n volatile int *flag = NULL;
\n vector<int *> buffers(numGPUs);
\n vector<int *> buffersD2D(numGPUs);
\n vector<musaEvent_t> start(numGPUs);
\n vector<musaEvent_t> stop(numGPUs);
\n vector<musaStream_t> stream0(numGPUs);
\n vector<musaStream_t> stream1(numGPUs);<\/p>\n

musaHostAlloc((void **)&flag, sizeof(*flag), musaHostAllocPortable);
\n musaCheckError();<\/p>\n

for (int d = 0; d < numGPUs; d++) {
\n musaSetDevice(d);
\n musaMalloc(&buffers[d], numElems * sizeof(int));
\n musaMemset(buffers[d], 0, numElems * sizeof(int));
\n musaMalloc(&buffersD2D[d], numElems * sizeof(int));
\n musaMemset(buffersD2D[d], 0, numElems * sizeof(int));
\n musaCheckError();
\n musaEventCreate(&start[d]);
\n musaCheckError();
\n musaEventCreate(&stop[d]);
\n musaCheckError();
\n musaStreamCreateWithFlags(&stream0[d], musaStreamNonBlocking);
\n musaCheckError();
\n musaStreamCreateWithFlags(&stream1[d], musaStreamNonBlocking);
\n musaCheckError();
\n }<\/p>\n

vector<double> bandwidthMatrix(numGPUs * numGPUs);<\/p>\n

for (int i = 0; i < numGPUs; i++) {
\n musaSetDevice(i);<\/p>\n

for (int j = 0; j < numGPUs; j++) {
\n int access = 0;
\n if (p2p) {
\n musaDeviceCanAccessPeer(&access, i, j);
\n if (access) {
\n musaSetDevice(i);
\n musaDeviceEnablePeerAccess(j, 0);
\n musaCheckError();
\n musaSetDevice(j);
\n musaDeviceEnablePeerAccess(i, 0);
\n musaCheckError();
\n }
\n }<\/p>\n

musaSetDevice(i);
\n musaStreamSynchronize(stream0[i]);
\n musaStreamSynchronize(stream1[j]);
\n musaCheckError();<\/p>\n

\/\/ \u963b\u585e\u6d41，\u76f4\u5230\u6240\u6709\u5de5\u4f5c\u6392\u961f\u5b8c\u6210
\n *flag = 0;
\n musaSetDevice(i);
\n \/\/ \u65e0\u9700\u963b\u585e stream1，\u56e0\u4e3a\u5b83\u4f1a\u5728 stream0 \u7684\u4e8b\u4ef6\u4e0a\u963b\u585e
\n delay<<<1, 1, 0, stream0[i]>>>(flag);
\n musaCheckError();<\/p>\n

\/\/ \u5f3a\u5236 stream1 \u5728 stream0 \u5f00\u59cb\u4e4b\u524d\u4e0d\u542f\u52a8，\u4ee5\u786e\u4fdd stream0 \u4e0a\u7684\u4e8b\u4ef6\u5b8c\u5168\u6db5\u76d6\u6240\u6709\u64cd\u4f5c\u6240\u9700\u7684\u65f6\u95f4
\n musaEventRecord(start[i], stream0[i]);
\n musaStreamWaitEvent(stream1[j], start[i], 0);<\/p>\n

if (i == j) {
\n \/\/ \u5bf9\u4e8e GPU \u5185\u64cd\u4f5c，\u6267\u884c 2 \u6b21\u5185\u5b58\u590d\u5236 buffersD2D <-> buffers
\n performP2PCopy(buffers[i], i, buffersD2D[i], i, numElems, repeat, access, stream0[i]);
\n performP2PCopy(buffersD2D[i], i, buffers[i], i, numElems, repeat, access, stream1[i]);
\n }
\n else {
\n if (access && p2p_mechanism == SM) {
\n musaSetDevice(j);
\n }
\n performP2PCopy(buffers[i], i, buffers[j], j, numElems, repeat, access, stream1[j]);
\n if (access && p2p_mechanism == SM) {
\n musaSetDevice(i);
\n }
\n performP2PCopy(buffers[j], j, buffers[i], i, numElems, repeat, access, stream0[i]);
\n }<\/p>\n

\/\/ \u901a\u77e5 stream0 stream1 \u5df2\u5b8c\u6210，\u5e76\u8bb0\u5f55\u603b\u4e8b\u52a1\u7684\u65f6\u95f4
\n musaEventRecord(stop[j], stream1[j]);
\n musaStreamWaitEvent(stream0[i], stop[j], 0);
\n musaEventRecord(stop[i], stream0[i]);<\/p>\n

\/\/ \u91ca\u653e\u6392\u961f\u7684\u64cd\u4f5c
\n *flag = 1;
\n musaStreamSynchronize(stream0[i]);
\n musaStreamSynchronize(stream1[j]);
\n musaCheckError();<\/p>\n

float time_ms;
\n musaEventElapsedTime(&time_ms, start[i], stop[i]);
\n double time_s = time_ms \/ 1e3;<\/p>\n

double gb = 2.0 * numElems * sizeof(int) * repeat \/ (double)1e9;
\n if (i == j) {
\n gb *= 2;
\n }
\n bandwidthMatrix[i * numGPUs + j] = gb \/ time_s;
\n if (p2p && access) {
\n musaSetDevice(i);
\n musaDeviceDisablePeerAccess(j);
\n musaSetDevice(j);
\n musaDeviceDisablePeerAccess(i);
\n }
\n }
\n }<\/p>\n

printf(" D\\\\\\\\D");
\n for (int j = 0; j < numGPUs; j++) {
\n printf("%6d ", j);
\n }
\n printf("\\\\n");<\/p>\n

for (int i = 0; i < numGPUs; i++) {
\n printf("%6d ", i);
\n for (int j = 0; j < numGPUs; j++) {
\n printf("%6.02f ", bandwidthMatrix[i * numGPUs + j]);
\n }
\n printf("\\\\n");
\n }<\/p>\n

for (int d = 0; d < numGPUs; d++) {
\n musaSetDevice(d);
\n musaFree(buffers[d]);
\n musaFree(buffersD2D[d]);
\n musaCheckError();
\n musaEventDestroy(start[d]);
\n musaCheckError();
\n musaEventDestroy(stop[d]);
\n musaCheckError();
\n musaStreamDestroy(stream0[d]);
\n musaCheckError();
\n musaStreamDestroy(stream1[d]);
\n musaCheckError();
\n }<\/p>\n

musaFreeHost((void *)flag);
\n musaCheckError();
\n}<\/p>\n

\/\/ \u8f93\u51fa\u5ef6\u8fdf\u77e9\u9635
\nvoid outputLatencyMatrix(int numGPUs, bool p2p, P2PDataTransfer p2p_method) {
\n int repeat = 100;
\n int numElems = 4; \/\/ \u6267\u884c 1 \u4e2a int4 \u4f20\u8f93
\n volatile int *flag = NULL;
\n vector<int *> buffers(numGPUs);
\n vector<int *> buffersD2D(numGPUs); \/\/ \u7528\u4e8e D2D（\u5373 GPU \u5185\u590d\u5236）\u7684\u7f13\u51b2\u533a
\n vector<musaStream_t> stream(numGPUs);
\n vector<musaEvent_t> start(numGPUs);
\n vector<musaEvent_t> stop(numGPUs);<\/p>\n

musaHostAlloc((void **)&flag, sizeof(*flag), musaHostAllocPortable);
\n musaCheckError();<\/p>\n

for (int d = 0; d < numGPUs; d++) {
\n musaSetDevice(d);
\n musaStreamCreateWithFlags(&stream[d], musaStreamNonBlocking);
\n musaMalloc(&buffers[d], sizeof(int) * numElems);
\n musaMemset(buffers[d], 0, sizeof(int) * numElems);
\n musaMalloc(&buffersD2D[d], sizeof(int) * numElems);
\n musaMemset(buffersD2D[d], 0, sizeof(int) * numElems);
\n musaCheckError();
\n musaEventCreate(&start[d]);
\n musaCheckError();
\n musaEventCreate(&stop[d]);
\n musaCheckError();
\n }<\/p>\n

vector<double> gpuLatencyMatrix(numGPUs * numGPUs);
\n vector<double> cpuLatencyMatrix(numGPUs * numGPUs);<\/p>\n

for (int i = 0; i < numGPUs; i++) {
\n musaSetDevice(i);<\/p>\n

for (int j = 0; j < numGPUs; j++) {
\n int access = 0;
\n if (p2p) {
\n musaDeviceCanAccessPeer(&access, i, j);
\n if (access) {
\n musaDeviceEnablePeerAccess(j, 0);
\n musaCheckError();
\n musaSetDevice(j);
\n musaDeviceEnablePeerAccess(i, 0);
\n musaSetDevice(i);
\n musaCheckError();
\n }
\n }
\n musaStreamSynchronize(stream[i]);
\n musaCheckError();<\/p>\n

\/\/ \u963b\u585e\u6d41，\u76f4\u5230\u6240\u6709\u5de5\u4f5c\u6392\u961f\u5b8c\u6210
\n *flag = 0;
\n delay<<<1, 1, 0, stream[i]>>>(flag);
\n musaCheckError();
\n musaEventRecord(start[i], stream[i]);<\/p>\n

if (i == j) {
\n \/\/ \u6267\u884c GPU \u5185\u7684 D2D \u590d\u5236
\n performP2PCopy(buffers[i], i, buffersD2D[i], i, numElems, repeat, access, stream[i]);
\n }
\n else {
\n if (p2p_method == P2P_WRITE) {
\n performP2PCopy(buffers[j], j, buffers[i], i, numElems, repeat, access, stream[i]);
\n }
\n else {
\n performP2PCopy(buffers[i], i, buffers[j], j, numElems, repeat, access, stream[i]);
\n }
\n }<\/p>\n

musaEventRecord(stop[i], stream[i]);
\n \/\/ \u73b0\u5728\u5de5\u4f5c\u5df2\u7ecf\u6392\u961f\u5b8c\u6210，\u91ca\u653e\u6d41
\n *flag = 1;
\n musaStreamSynchronize(stream[i]);
\n musaCheckError();<\/p>\n

float gpu_time_ms;
\n musaEventElapsedTime(&gpu_time_ms, start[i], stop[i]);<\/p>\n

gpuLatencyMatrix[i * numGPUs + j] = gpu_time_ms * 1e3 \/ repeat;
\n if (p2p && access) {
\n musaDeviceDisablePeerAccess(j);
\n musaSetDevice(j);
\n musaDeviceDisablePeerAccess(i);
\n musaSetDevice(i);
\n musaCheckError();
\n }
\n }
\n }<\/p>\n

printf(" GPU");
\n for (int j = 0; j < numGPUs; j++) {
\n printf("%6d ", j);
\n }
\n printf("\\\\n");<\/p>\n

for (int i = 0; i < numGPUs; i++) {
\n printf("%6d ", i);
\n for (int j = 0; j < numGPUs; j++) {
\n printf("%6.02f ", gpuLatencyMatrix[i * numGPUs + j]);
\n }
\n printf("\\\\n");
\n }<\/p>\n

for (int d = 0; d < numGPUs; d++) {
\n musaSetDevice(d);
\n musaFree(buffers[d]);
\n musaFree(buffersD2D[d]);
\n musaCheckError();
\n musaEventDestroy(start[d]);
\n musaCheckError();
\n musaEventDestroy(stop[d]);
\n musaCheckError();
\n musaStreamDestroy(stream[d]);
\n musaCheckError();
\n }<\/p>\n

musaFreeHost((void *)flag);
\n musaCheckError();
\n}<\/p>\n

\/\/ \u4e3b\u51fd\u6570
\nint main(int argc, char **argv) {
\n int numGPUs, numElems = 40000000;
\n P2PDataTransfer p2p_method = P2P_WRITE;<\/p>\n

musaGetDeviceCount(&numGPUs);
\n musaCheckError();<\/p>\n

\/\/ \u5904\u7406\u547d\u4ee4\u884c\u53c2\u6570
\n for (int i = 1; i < argc; i++) {
\n if (strcmp(argv[i], "–help") == 0) {
\n printHelp();
\n return 0;
\n } else if (strcmp(argv[i], "–p2p_read") == 0) {
\n p2p_method = P2P_READ;
\n } else if (strcmp(argv[i], "–sm_copy") == 0) {
\n p2p_mechanism = SM;
\n } else if (strncmp(argv[i], "–numElems=", 11) == 0) {
\n numElems = atoi(argv[i] + 11);
\n }
\n }<\/p>\n

printf("[%s]\\\\n", sSampleName);<\/p>\n

\/\/ \u8f93\u51fa\u8bbe\u5907\u4fe1\u606f
\n for (int i = 0; i < numGPUs; i++) {
\n musaDeviceProp prop;
\n musaGetDeviceProperties(&prop, i);
\n printf("Device: %d, %s, pciBusID: %x, pciDeviceID: %x, pciDomainID:%x\\\\n",
\n i, prop.name, prop.pciBusID, prop.pciDeviceID, prop.pciDomainID);
\n }<\/p>\n

checkP2Paccess(numGPUs);<\/p>\n

\/\/ \u8f93\u51faP2P\u8fde\u63a5\u77e9\u9635
\n printf("P2P Connectivity Matrix\\\\n");
\n printf(" D\\\\\\\\D");
\n for (int j = 0; j < numGPUs; j++) {
\n printf("%6d", j);
\n }
\n printf("\\\\n");<\/p>\n

for (int i = 0; i < numGPUs; i++) {
\n printf("%6d\\\\t", i);
\n for (int j = 0; j < numGPUs; j++) {
\n if (i != j) {
\n int access;
\n musaDeviceCanAccessPeer(&access, i, j);
\n printf("%6d", (access) ? 1 : 0);
\n } else {
\n printf("%6d", 1);
\n }
\n }
\n printf("\\\\n");
\n }<\/p>\n

\/\/ \u8f93\u51fa\u5404\u79cd\u6d4b\u8bd5\u7ed3\u679c
\n printf("Unidirectional P2P=Disabled Bandwidth Matrix (GB\/s)\\\\n");
\n outputBandwidthMatrix(numElems, numGPUs, false, P2P_WRITE);
\n printf("Unidirectional P2P=Enabled Bandwidth (P2P Writes) Matrix (GB\/s)\\\\n");
\n outputBandwidthMatrix(numElems, numGPUs, true, P2P_WRITE);
\n if (p2p_method == P2P_READ) {
\n printf("Unidirectional P2P=Enabled Bandwidth (P2P Reads) Matrix (GB\/s)\\\\n");
\n outputBandwidthMatrix(numElems, numGPUs, true, p2p_method);
\n }
\n printf("Bidirectional P2P=Disabled Bandwidth Matrix (GB\/s)\\\\n");
\n outputBidirectionalBandwidthMatrix(numElems, numGPUs, false);
\n printf("Bidirectional P2P=Enabled Bandwidth Matrix (GB\/s)\\\\n");
\n outputBidirectionalBandwidthMatrix(numElems, numGPUs, true);<\/p>\n

printf("P2P=Disabled Latency Matrix (us)\\\\n");
\n outputLatencyMatrix(numGPUs, false, P2P_WRITE);
\n printf("P2P=Enabled Latency (P2P Writes) Matrix (us)\\\\n");
\n outputLatencyMatrix(numGPUs, true, P2P_WRITE);
\n if (p2p_method == P2P_READ) {
\n printf("P2P=Enabled Latency (P2P Reads) Matrix (us)\\\\n");
\n outputLatencyMatrix(numGPUs, true, p2p_method);
\n }<\/p>\n

printf("\\\\nNOTE: Results may vary when GPU Boost is enabled.\\\\n");<\/p>\n

return 0;
\n}<\/p>\n

\u7f16\u8bd1<\/h4>\n

\u53c2\u8003mcc\u7f16\u8bd1\u624b\u518c，\u6b64\u65f6\u4ee3\u7801\u4e2d\u5f15\u7528\u7684\u5e93\u4e3amusa_runtime,\u5219\u7f16\u8bd1\u662f-l\u53c2\u6570\u540e\u8ddfmusart<\/p>\n

mcc p2p.mu -o p2p -lmusart<\/p>\n

\u7ed3\u679c<\/h4>\n

\u53ef\u4ee5\u770b\u5230p2p\u5df2\u7ecf\u6b63\u786e\u5f00\u542f，\u4f46\u662f\u5ef6\u8fdf\u6d4b\u8bd5\u6709\u95ee\u9898，\u540e\u7eed\u6539\u8fdb<\/p>\n

\"\"<\/p>\n

\u57fa\u4e8emusa\u7f16\u7a0b\u7684allreduce\u6d4b\u8bd5<\/h3>\n

\u4ee3\u7801\u53c2\u8003<\/h4>\n

\u4e3b\u8981\u53c2\u8003\u4e86NCCLtest\u4e2d\u7684allreduce\u90e8\u5206\u903b\u8f91<\/p>\n

GitHub – NVIDIA\/nccl-tests: NCCL Tests<\/p>\n

\u5e76\u4e14\u53c2\u8003\u4e86mublas api\u8bbe\u8ba1 https:\/\/docs.mthreads.com\/musa-sdk\/musa-sdk-doc-online\/api\/mublas_api<\/p>\n

\u4ee3\u7801\u90e8\u5206<\/h4>\n

#include <stdio.h>
\n#include <stdlib.h>
\n#include <math.h>
\n#include "musa_runtime.h"
\n#include "mccl.h"
\n#include <inttypes.h> \/\/ \u5fc5\u987b\u5305\u542b\u6b64\u5934\u6587\u4ef6<\/p>\n

\/\/ \u5b8f\u5b9a\u4e49（\u6240\u6709\u6807\u8bc6\u7b26\u5728\u6b64\u5904\u58f0\u660e）
\n#define MIN_SIZE_B 16ULL \/\/ \u6700\u5c0f\u6d4b\u8bd5\u5c3a\u5bf8（16\u5b57\u8282）
\n#define MAX_SIZE_B (4096ULL * 1024ULL * 1024ULL) \/\/ \u6700\u5927\u6d4b\u8bd5\u5c3a\u5bf8（4096MB）
\n#define STEP_FACTOR 2ULL \/\/ \u5c3a\u5bf8\u589e\u957f\u56e0\u5b50（\u6bcf\u6b21\u7ffb\u500d）
\n#define WARMUP_ITERS 5 \/\/ \u70ed\u8eab\u8fed\u4ee3\u6b21\u6570
\n#define TEST_ITERS 20 \/\/ \u6d4b\u8bd5\u8fed\u4ee3\u6b21\u6570
\n#define ROOT_RANK -1 \/\/ \u6839\u8282\u70b9（-1\u8868\u793a\u5168\u5f52\u7ea6）
\n#define DATA_TYPE mcclFloat \/\/ \u6570\u636e\u7c7b\u578b
\n#define REDUCTION_OP mcclSum \/\/ \u5f52\u7ea6\u64cd\u4f5c
\n#define FLOAT_SIZE sizeof(float) \/\/ float\u7c7b\u578b\u5b57\u8282\u6570（4\u5b57\u8282）<\/p>\n

\/\/ \u9519\u8bef\u68c0\u67e5\u5b8f
\n#define MUSACHECK(cmd) do { \\\\
\n musaError_t err = cmd; \\\\
\n if (err != musaSuccess) { \\\\
\n printf("MUSA Error at %s:%d: %s\\\\n", __FILE__, __LINE__, musaGetErrorString(err)); \\\\
\n exit(EXIT_FAILURE); \\\\
\n } \\\\
\n} while(0)<\/p>\n

#define MCCLCHECK(cmd) do { \\\\
\n mcclResult_t res = cmd; \\\\
\n if (res != mcclSuccess) { \\\\
\n printf("MCCL Error at %s:%d: %s\\\\n", __FILE__, __LINE__, mcclGetErrorString(res)); \\\\
\n exit(EXIT_FAILURE); \\\\
\n } \\\\
\n} while(0)<\/p>\n

\/\/ \u5e26\u5bbd\u8ba1\u7b97\u51fd\u6570
\nvoid calculate_bandwidth(size_t count, int type_size, double time_sec, double* alg_bw, double* bus_bw, int nranks) {
\n if (time_sec <= 0 || count == 0) {
\n *alg_bw = 0.0;
\n *bus_bw = 0.0;
\n return;
\n }
\n double data_size_gb = (double)(count * type_size) \/ 1e9;
\n *alg_bw = data_size_gb \/ time_sec;
\n double factor = (nranks > 1) ? (2.0 * (nranks – 1)) \/ nranks : 1.0;
\n *bus_bw = *alg_bw * factor;
\n}<\/p>\n

int main(int argc, char* argv[]) {
\n int nDev = 4; \/\/ \u8bbe\u5907\u6570\u91cf
\n int devs[4] = {0, 1, 2, 3}; \/\/ \u8bbe\u5907ID\u5217\u8868
\n mcclComm_t comms[4]; \/\/ MCCL\u901a\u4fe1\u5668
\n musaStream_t streams[4]; \/\/ \u6d41\u6570\u7ec4
\n float** sendbuff = NULL; \/\/ \u53d1\u9001\u7f13\u51b2\u533a
\n float** recvbuff = NULL; \/\/ \u63a5\u6536\u7f13\u51b2\u533a
\n size_t current_size_b = MIN_SIZE_B; \/\/ \u5f53\u524d\u6d4b\u8bd5\u5c3a\u5bf8（\u5b57\u8282）
\n double alg_bw, bus_bw; \/\/ \u7b97\u6cd5\u5e26\u5bbd\u548c\u603b\u7ebf\u5e26\u5bbd
\n int test_wrong = 0; \/\/ \u9519\u8bef\u8ba1\u6570<\/p>\n

\/\/ \u521d\u59cb\u5316MCCL\u901a\u4fe1\u5668
\n MCCLCHECK(mcclCommInitAll(comms, nDev, devs));<\/p>\n

\/\/ \u5206\u914d\u8bbe\u5907\u5185\u5b58\u5e76\u521b\u5efa\u6d41
\n sendbuff = (float**)malloc(nDev * sizeof(float*));
\n recvbuff = (float**)malloc(nDev * sizeof(float*));
\n for (int i = 0; i < nDev; ++i) {
\n MUSACHECK(musaSetDevice(i));
\n MUSACHECK(musaMalloc(&sendbuff[i], MAX_SIZE_B)); \/\/ \u5206\u914d\u6700\u5927\u5c3a\u5bf8\u5185\u5b58
\n MUSACHECK(musaMalloc(&recvbuff[i], MAX_SIZE_B));
\n MUSACHECK(musaStreamCreate(&streams[i])); \/\/ \u521b\u5efa\u72ec\u7acb\u6d41
\n }<\/p>\n

\/\/ \u6253\u5370\u7ed3\u679c\u8868\u5934
\n printf("| %10s | %10s | %5s | %4s | %14s | %13s | %13s | %13s | %5s |\\\\n",
\n "size (B)", "count", "type", "root", "warmup_time (us)", "test_time (us)", "alg_bw (GB\/s)", "bus_bw (GB\/s)", "#wrong");
\n printf("|————|————|——-|——|——————|—————-|—————|—————|——–|\\\\n");
\n \/\/ \u5c3a\u5bf8\u5faa\u73af\u6d4b\u8bd5
\n while (current_size_b <= MAX_SIZE_B) {
\n size_t element_count = current_size_b \/ FLOAT_SIZE; \/\/ \u5143\u7d20\u6570\u91cf<\/p>\n

\/\/ \u8df3\u8fc7\u975e\u5bf9\u9f50\u5c3a\u5bf8
\n if (current_size_b % FLOAT_SIZE != 0) {
\n current_size_b *= STEP_FACTOR;
\n continue;
\n }<\/p>\n

\/\/ \u521d\u59cb\u5316\u8bbe\u5907\u6570\u636e（\u901a\u8fc7\u4e3b\u673a\u5185\u5b58\u6b63\u786e\u8d4b\u503c\u4e3a1.0f）
\n for (int i = 0; i < nDev; ++i) {
\n MUSACHECK(musaSetDevice(i));
\n float* host_buf = (float*)malloc(current_size_b);
\n for (size_t j = 0; j < element_count; ++j) host_buf[j] = 1.0f;
\n MUSACHECK(musaMemcpy(sendbuff[i], host_buf, current_size_b, musaMemcpyHostToDevice));
\n free(host_buf);
\n MUSACHECK(musaMemset(recvbuff[i], 0, current_size_b));
\n }<\/p>\n

\/\/ \u70ed\u8eab\u8fed\u4ee3（\u5305\u542b\u6d41\u540c\u6b65）
\n for (int warmup = 0; warmup < WARMUP_ITERS; ++warmup) {
\n MCCLCHECK(mcclGroupStart());
\n for (int i = 0; i < nDev; ++i) {
\n MCCLCHECK(mcclAllReduce(
\n sendbuff[i], recvbuff[i],
\n element_count, DATA_TYPE, REDUCTION_OP,
\n comms[i], streams[i]
\n ));
\n }
\n MCCLCHECK(mcclGroupEnd());
\n for (int i = 0; i < nDev; ++i) {
\n MUSACHECK(musaSetDevice(i));
\n MUSACHECK(musaStreamSynchronize(streams[i]));
\n }
\n }<\/p>\n

\/\/ \u4e8b\u4ef6\u8ba1\u65f6（\u4ec5\u5728\u4e3b\u8bbe\u59070\u64cd\u4f5c）
\n musaEvent_t start, stop;
\n MUSACHECK(musaSetDevice(0));
\n MUSACHECK(musaEventCreate(&start));
\n MUSACHECK(musaEventCreate(&stop));
\n MUSACHECK(musaEventRecord(start, streams[0]));<\/p>\n

\/\/ \u6d4b\u8bd5\u8fed\u4ee3（\u5305\u542b\u5b8c\u6574Group\u64cd\u4f5c）
\n MCCLCHECK(mcclGroupStart());
\n for (int iter = 0; iter < TEST_ITERS; ++iter) {
\n for (int i = 0; i < nDev; ++i) {
\n MUSACHECK(musaSetDevice(i));
\n MCCLCHECK(mcclAllReduce(
\n sendbuff[i], recvbuff[i],
\n element_count, DATA_TYPE, REDUCTION_OP,
\n comms[i], streams[i]
\n ));
\n }
\n }
\n MCCLCHECK(mcclGroupEnd());<\/p>\n

MUSACHECK(musaEventRecord(stop, streams[0]));
\n MUSACHECK(musaEventSynchronize(stop));<\/p>\n

\/\/ \u8ba1\u7b97\u5e73\u5747\u65f6\u95f4
\n float total_time_ms;
\n MUSACHECK(musaEventElapsedTime(&total_time_ms, start, stop));
\n double avg_time_us = (total_time_ms \/ TEST_ITERS) * 1000;<\/p>\n

\/\/ \u8ba1\u7b97\u5e26\u5bbd
\n calculate_bandwidth(element_count, FLOAT_SIZE, avg_time_us \/ 1e6, &alg_bw, &bus_bw, nDev);<\/p>\n

\/\/ \u9a8c\u8bc1\u7ed3\u679c（\u5141\u8bb8\u6d6e\u70b9\u7cbe\u5ea6\u8bef\u5dee）
\n test_wrong = 0;
\n float expected = (float)nDev;
\n for (int i = 0; i < nDev; ++i) {
\n MUSACHECK(musaSetDevice(i));
\n float* h_recv = (float*)malloc(current_size_b);
\n MUSACHECK(musaMemcpy(h_recv, recvbuff[i], current_size_b, musaMemcpyDeviceToHost));
\n for (size_t j = 0; j < element_count; ++j) {
\n if (fabs(h_recv[j] – expected) > 1e-6) test_wrong++;
\n }
\n free(h_recv);
\n }<\/p>\n

\/\/ \u6253\u5370\u7ed3\u679c
\n printf("| %10" PRIu64 " | %10" PRIu64 " | %4s | %4d | %16.3f | %14.3f | %13.3f | %13.3f | %6d |\\\\n",
\n (uint64_t)current_size_b, (uint64_t)element_count, "float", ROOT_RANK, 0.0, avg_time_us, alg_bw, bus_bw, test_wrong);<\/p>\n

\/\/ \u9500\u6bc1\u4e8b\u4ef6
\n MUSACHECK(musaSetDevice(0));
\n MUSACHECK(musaEventDestroy(start));
\n MUSACHECK(musaEventDestroy(stop));<\/p>\n

\/\/ \u589e\u5927\u6d4b\u8bd5\u5c3a\u5bf8
\n current_size_b *= STEP_FACTOR;
\n }<\/p>\n

\/\/ \u91ca\u653e\u8d44\u6e90
\n for (int i = 0; i < nDev; ++i) {
\n MUSACHECK(musaSetDevice(i));
\n MUSACHECK(musaFree(sendbuff[i]));
\n MUSACHECK(musaFree(recvbuff[i]));
\n MUSACHECK(musaStreamDestroy(streams[i]));
\n mcclCommDestroy(comms[i]);
\n }
\n free(sendbuff);
\n free(recvbuff);<\/p>\n

printf("AllReduce Test Completed Successfully\\\\n");
\n return 0;
\n} <\/p>\n

\u7f16\u8bd1<\/h4>\n

\u56e0\u4e3a\u4ee3\u7801\u7528\u4e86musa_runtime\u4e0emccl\u4e24\u4e2a\u5e93，\u56e0\u6b64\u7f16\u8bd1\u9009\u9879\u4e5f\u4f1a\u6709\u6240\u6539\u53d8<\/p>\n

mcc allreduce.mu -o allreduce -lmusart -lmccl<\/p>\n

\u7ed3\u679c<\/h4>\n

\u4e0d\u77e5\u9053\u4e3a\u4ec0\u4e48\u7ed3\u679c\u6d4b\u51fa\u6765\u548c\u7528pytorch\u6d4b\u51fa\u6765\u7ed3\u679c\u76f8\u5dee\u4e0d\u5c0f，\u76ee\u6d4b\u662f\u56e0\u4e3amusa event\u6253\u70b9\u8ba1\u65f6\u51fd\u6570\u6ca1\u4f7f\u7528\u6b63\u786e（\u5728p2p\u6d4b\u8bd5\u7684\u81ea\u4ea4\u4e2d\u4e5f\u6709\u4f53\u73b0，\u4e0d\u7ba1\u4ec0\u4e48\u60c5\u51b5\u90fd\u662f50us\u5de6\u53f3），\u8fd9\u4e2a\u9700\u8981\u540e\u7eed\u518d\u770b\u4e0b<\/p>\n

\"\"<\/p>\n

\"\"<\/p>\n","protected":false},"excerpt":{"rendered":"

\u6587\u7ae0\u6d4f\u89c8\u9605\u8bfb1.1k\u6b21\uff0c\u70b9\u8d5e16\u6b21\uff0c\u6536\u85cf12\u6b21\u3002MTT S4000 \u662f\u57fa\u4e8e\u6469\u5c14\u7ebf\u7a0b\u66f2\u9662 GPU \u67b6\u6784\u6253\u9020\u7684\u5168\u529f\u80fd\u5143\u8ba1\u7b97\u5361\uff0c\u4e3a\u5343\u4ebf\u89c4\u6a21\u5927\u8bed\u8a00\u6a21\u578b\u7684\u8bad\u7ec3\u3001\u5fae\u8c03\u548c\u63a8\u7406\u8fdb\u884c\u4e86\u5b9a\u5236\u4f18\u5316\uff0c\u7ed3\u5408\u5148\u8fdb\u7684\u56fe\u5f62\u6e32\u67d3\u80fd\u529b\u3001\u89c6\u9891\u7f16\u89e3\u7801\u80fd\u529b\u548c\u8d85\u9ad8\u6e05 8K HDR \u663e\u793a\u80fd\u529b\uff0c\u52a9\u529b\u4eba\u5de5\u667a\u80fd\u3001\u56fe\u5f62\u6e32\u67d3\u3001\u591a\u5a92\u4f53\u3001\u79d1\u5b66\u8ba1\u7b97\u4e0e\u7269\u7406\u4eff\u771f\u7b49\u590d\u5408\u5e94\u7528\u573a\u666f\u7684\u8ba1\u7b97\u52a0\u901f\u3002MTT S4000 \u5168\u9762\u652f\u6301\u5927\u8bed\u8a00\u6a21\u578b\u7684\u9884\u8bad\u7ec3\u3001\u5fae\u8c03\u548c\u63a8\u7406\u670d\u52a1\uff0cMUSA \u8f6f\u4ef6\u6808\u4e13\u95e8\u9488\u5bf9\u5927\u89c4\u6a21\u96c6\u7fa4\u7684\u5206\u5e03\u5f0f\u8ba1\u7b97\u6027\u80fd\u8fdb\u884c\u4e86\u4f18\u5316\uff0c\u9002\u914d\u4e3b\u6d41\u5206\u5e03\u5f0f\u8ba1\u7b97\u52a0\u901f\u6846\u67b6\uff0c \u5305\u62ec DeepSpeed\uff0c Colossal AI\uff0cMegatron \u7b49\uff0c\u652f_\u6469\u5c14\u7ebf\u7a0b\u4fe1\u521b<\/p>\n","protected":false},"author":2,"featured_media":40324,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[1],"tags":[66,1458,81,152,50,3694,43,86],"topic":[],"class_list":["post-40335","post","type-post","status-publish","format-standard","has-post-thumbnail","hentry","category-server","tag-ai","tag-hpc","tag-python","tag-pytorch","tag-50","tag-3694","tag-43","tag-86"],"yoast_head":"\n\u6469\u5c14\u7ebf\u7a0bS4000\u56fd\u4ea7\u4fe1\u521b\u8ba1\u7b97\u5361\u6027\u80fd\u5b9e\u6218\u2014\u2014Pytorch\u8f6c\u8bd1\uff0c\u591a\u5361P2P\u901a\u4fe1\u4e0eMUSA\u7f16\u7a0b - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3<\/title>\n<meta name=\"robots\" content=\"index, follow, max-snippet:-1, max-image-preview:large, max-video-preview:-1\" \/>\n<link rel=\"canonical\" href=\"https:\/\/www.wsisp.com\/helps\/40335.html\" \/>\n<meta property=\"og:locale\" content=\"zh_CN\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"\u6469\u5c14\u7ebf\u7a0bS4000\u56fd\u4ea7\u4fe1\u521b\u8ba1\u7b97\u5361\u6027\u80fd\u5b9e\u6218\u2014\u2014Pytorch\u8f6c\u8bd1\uff0c\u591a\u5361P2P\u901a\u4fe1\u4e0eMUSA\u7f16\u7a0b - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\" \/>\n<meta property=\"og:description\" content=\"\u6587\u7ae0\u6d4f\u89c8\u9605\u8bfb1.1k\u6b21\uff0c\u70b9\u8d5e16\u6b21\uff0c\u6536\u85cf12\u6b21\u3002MTT S4000 \u662f\u57fa\u4e8e\u6469\u5c14\u7ebf\u7a0b\u66f2\u9662 GPU \u67b6\u6784\u6253\u9020\u7684\u5168\u529f\u80fd\u5143\u8ba1\u7b97\u5361\uff0c\u4e3a\u5343\u4ebf\u89c4\u6a21\u5927\u8bed\u8a00\u6a21\u578b\u7684\u8bad\u7ec3\u3001\u5fae\u8c03\u548c\u63a8\u7406\u8fdb\u884c\u4e86\u5b9a\u5236\u4f18\u5316\uff0c\u7ed3\u5408\u5148\u8fdb\u7684\u56fe\u5f62\u6e32\u67d3\u80fd\u529b\u3001\u89c6\u9891\u7f16\u89e3\u7801\u80fd\u529b\u548c\u8d85\u9ad8\u6e05 8K HDR \u663e\u793a\u80fd\u529b\uff0c\u52a9\u529b\u4eba\u5de5\u667a\u80fd\u3001\u56fe\u5f62\u6e32\u67d3\u3001\u591a\u5a92\u4f53\u3001\u79d1\u5b66\u8ba1\u7b97\u4e0e\u7269\u7406\u4eff\u771f\u7b49\u590d\u5408\u5e94\u7528\u573a\u666f\u7684\u8ba1\u7b97\u52a0\u901f\u3002MTT S4000 \u5168\u9762\u652f\u6301\u5927\u8bed\u8a00\u6a21\u578b\u7684\u9884\u8bad\u7ec3\u3001\u5fae\u8c03\u548c\u63a8\u7406\u670d\u52a1\uff0cMUSA \u8f6f\u4ef6\u6808\u4e13\u95e8\u9488\u5bf9\u5927\u89c4\u6a21\u96c6\u7fa4\u7684\u5206\u5e03\u5f0f\u8ba1\u7b97\u6027\u80fd\u8fdb\u884c\u4e86\u4f18\u5316\uff0c\u9002\u914d\u4e3b\u6d41\u5206\u5e03\u5f0f\u8ba1\u7b97\u52a0\u901f\u6846\u67b6\uff0c \u5305\u62ec DeepSpeed\uff0c Colossal AI\uff0cMegatron \u7b49\uff0c\u652f_\u6469\u5c14\u7ebf\u7a0b\u4fe1\u521b\" \/>\n<meta property=\"og:url\" content=\"https:\/\/www.wsisp.com\/helps\/40335.html\" \/>\n<meta property=\"og:site_name\" content=\"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\" \/>\n<meta property=\"article:published_time\" content=\"2025-05-28T05:43:56+00:00\" \/>\n<meta property=\"og:image\" content=\"https:\/\/www.wsisp.com\/helps\/wp-content\/uploads\/2025\/05\/20250528054352-6836a29866a81.png\" \/>\n<meta name=\"author\" content=\"admin\" \/>\n<meta name=\"twitter:card\" content=\"summary_large_image\" \/>\n<meta name=\"twitter:label1\" content=\"\u4f5c\u8005\" \/>\n\t<meta name=\"twitter:data1\" content=\"admin\" \/>\n\t<meta name=\"twitter:label2\" content=\"\u9884\u8ba1\u9605\u8bfb\u65f6\u95f4\" \/>\n\t<meta name=\"twitter:data2\" content=\"15 \u5206\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\/\/schema.org\",\"@graph\":[{\"@type\":\"WebPage\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/40335.html\",\"url\":\"https:\/\/www.wsisp.com\/helps\/40335.html\",\"name\":\"\u6469\u5c14\u7ebf\u7a0bS4000\u56fd\u4ea7\u4fe1\u521b\u8ba1\u7b97\u5361\u6027\u80fd\u5b9e\u6218\u2014\u2014Pytorch\u8f6c\u8bd1\uff0c\u591a\u5361P2P\u901a\u4fe1\u4e0eMUSA\u7f16\u7a0b - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\",\"isPartOf\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#website\"},\"datePublished\":\"2025-05-28T05:43:56+00:00\",\"dateModified\":\"2025-05-28T05:43:56+00:00\",\"author\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41\"},\"breadcrumb\":{\"@id\":\"https:\/\/www.wsisp.com\/helps\/40335.html#breadcrumb\"},\"inLanguage\":\"zh-Hans\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\/\/www.wsisp.com\/helps\/40335.html\"]}]},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/40335.html#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"\u9996\u9875\",\"item\":\"https:\/\/www.wsisp.com\/helps\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"\u6469\u5c14\u7ebf\u7a0bS4000\u56fd\u4ea7\u4fe1\u521b\u8ba1\u7b97\u5361\u6027\u80fd\u5b9e\u6218\u2014\u2014Pytorch\u8f6c\u8bd1\uff0c\u591a\u5361P2P\u901a\u4fe1\u4e0eMUSA\u7f16\u7a0b\"}]},{\"@type\":\"WebSite\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/#website\",\"url\":\"https:\/\/www.wsisp.com\/helps\/\",\"name\":\"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3\",\"description\":\"\u9999\u6e2f\u670d\u52a1\u5668_\u9999\u6e2f\u4e91\u670d\u52a1\u5668\u8d44\u8baf_\u670d\u52a1\u5668\u5e2e\u52a9\u6587\u6863_\u670d\u52a1\u5668\u6559\u7a0b\",\"potentialAction\":[{\"@type\":\"SearchAction\",\"target\":{\"@type\":\"EntryPoint\",\"urlTemplate\":\"https:\/\/www.wsisp.com\/helps\/?s={search_term_string}\"},\"query-input\":\"required name=search_term_string\"}],\"inLanguage\":\"zh-Hans\"},{\"@type\":\"Person\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41\",\"name\":\"admin\",\"image\":{\"@type\":\"ImageObject\",\"inLanguage\":\"zh-Hans\",\"@id\":\"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/image\/\",\"url\":\"https:\/\/gravatar.wp-china-yes.net\/avatar\/?s=96&d=mystery\",\"contentUrl\":\"https:\/\/gravatar.wp-china-yes.net\/avatar\/?s=96&d=mystery\",\"caption\":\"admin\"},\"sameAs\":[\"http:\/\/wp.wsisp.com\"],\"url\":\"https:\/\/www.wsisp.com\/helps\/author\/admin\"}]}<\/script>\n<!-- \/ Yoast SEO plugin. -->","yoast_head_json":{"title":"\u6469\u5c14\u7ebf\u7a0bS4000\u56fd\u4ea7\u4fe1\u521b\u8ba1\u7b97\u5361\u6027\u80fd\u5b9e\u6218\u2014\u2014Pytorch\u8f6c\u8bd1\uff0c\u591a\u5361P2P\u901a\u4fe1\u4e0eMUSA\u7f16\u7a0b - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","robots":{"index":"index","follow":"follow","max-snippet":"max-snippet:-1","max-image-preview":"max-image-preview:large","max-video-preview":"max-video-preview:-1"},"canonical":"https:\/\/www.wsisp.com\/helps\/40335.html","og_locale":"zh_CN","og_type":"article","og_title":"\u6469\u5c14\u7ebf\u7a0bS4000\u56fd\u4ea7\u4fe1\u521b\u8ba1\u7b97\u5361\u6027\u80fd\u5b9e\u6218\u2014\u2014Pytorch\u8f6c\u8bd1\uff0c\u591a\u5361P2P\u901a\u4fe1\u4e0eMUSA\u7f16\u7a0b - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","og_description":"\u6587\u7ae0\u6d4f\u89c8\u9605\u8bfb1.1k\u6b21\uff0c\u70b9\u8d5e16\u6b21\uff0c\u6536\u85cf12\u6b21\u3002MTT S4000 \u662f\u57fa\u4e8e\u6469\u5c14\u7ebf\u7a0b\u66f2\u9662 GPU \u67b6\u6784\u6253\u9020\u7684\u5168\u529f\u80fd\u5143\u8ba1\u7b97\u5361\uff0c\u4e3a\u5343\u4ebf\u89c4\u6a21\u5927\u8bed\u8a00\u6a21\u578b\u7684\u8bad\u7ec3\u3001\u5fae\u8c03\u548c\u63a8\u7406\u8fdb\u884c\u4e86\u5b9a\u5236\u4f18\u5316\uff0c\u7ed3\u5408\u5148\u8fdb\u7684\u56fe\u5f62\u6e32\u67d3\u80fd\u529b\u3001\u89c6\u9891\u7f16\u89e3\u7801\u80fd\u529b\u548c\u8d85\u9ad8\u6e05 8K HDR \u663e\u793a\u80fd\u529b\uff0c\u52a9\u529b\u4eba\u5de5\u667a\u80fd\u3001\u56fe\u5f62\u6e32\u67d3\u3001\u591a\u5a92\u4f53\u3001\u79d1\u5b66\u8ba1\u7b97\u4e0e\u7269\u7406\u4eff\u771f\u7b49\u590d\u5408\u5e94\u7528\u573a\u666f\u7684\u8ba1\u7b97\u52a0\u901f\u3002MTT S4000 \u5168\u9762\u652f\u6301\u5927\u8bed\u8a00\u6a21\u578b\u7684\u9884\u8bad\u7ec3\u3001\u5fae\u8c03\u548c\u63a8\u7406\u670d\u52a1\uff0cMUSA \u8f6f\u4ef6\u6808\u4e13\u95e8\u9488\u5bf9\u5927\u89c4\u6a21\u96c6\u7fa4\u7684\u5206\u5e03\u5f0f\u8ba1\u7b97\u6027\u80fd\u8fdb\u884c\u4e86\u4f18\u5316\uff0c\u9002\u914d\u4e3b\u6d41\u5206\u5e03\u5f0f\u8ba1\u7b97\u52a0\u901f\u6846\u67b6\uff0c \u5305\u62ec DeepSpeed\uff0c Colossal AI\uff0cMegatron \u7b49\uff0c\u652f_\u6469\u5c14\u7ebf\u7a0b\u4fe1\u521b","og_url":"https:\/\/www.wsisp.com\/helps\/40335.html","og_site_name":"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","article_published_time":"2025-05-28T05:43:56+00:00","og_image":[{"url":"https:\/\/www.wsisp.com\/helps\/wp-content\/uploads\/2025\/05\/20250528054352-6836a29866a81.png"}],"author":"admin","twitter_card":"summary_large_image","twitter_misc":{"\u4f5c\u8005":"admin","\u9884\u8ba1\u9605\u8bfb\u65f6\u95f4":"15 \u5206"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"WebPage","@id":"https:\/\/www.wsisp.com\/helps\/40335.html","url":"https:\/\/www.wsisp.com\/helps\/40335.html","name":"\u6469\u5c14\u7ebf\u7a0bS4000\u56fd\u4ea7\u4fe1\u521b\u8ba1\u7b97\u5361\u6027\u80fd\u5b9e\u6218\u2014\u2014Pytorch\u8f6c\u8bd1\uff0c\u591a\u5361P2P\u901a\u4fe1\u4e0eMUSA\u7f16\u7a0b - \u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","isPartOf":{"@id":"https:\/\/www.wsisp.com\/helps\/#website"},"datePublished":"2025-05-28T05:43:56+00:00","dateModified":"2025-05-28T05:43:56+00:00","author":{"@id":"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41"},"breadcrumb":{"@id":"https:\/\/www.wsisp.com\/helps\/40335.html#breadcrumb"},"inLanguage":"zh-Hans","potentialAction":[{"@type":"ReadAction","target":["https:\/\/www.wsisp.com\/helps\/40335.html"]}]},{"@type":"BreadcrumbList","@id":"https:\/\/www.wsisp.com\/helps\/40335.html#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"\u9996\u9875","item":"https:\/\/www.wsisp.com\/helps"},{"@type":"ListItem","position":2,"name":"\u6469\u5c14\u7ebf\u7a0bS4000\u56fd\u4ea7\u4fe1\u521b\u8ba1\u7b97\u5361\u6027\u80fd\u5b9e\u6218\u2014\u2014Pytorch\u8f6c\u8bd1\uff0c\u591a\u5361P2P\u901a\u4fe1\u4e0eMUSA\u7f16\u7a0b"}]},{"@type":"WebSite","@id":"https:\/\/www.wsisp.com\/helps\/#website","url":"https:\/\/www.wsisp.com\/helps\/","name":"\u7f51\u7855\u4e92\u8054\u5e2e\u52a9\u4e2d\u5fc3","description":"\u9999\u6e2f\u670d\u52a1\u5668_\u9999\u6e2f\u4e91\u670d\u52a1\u5668\u8d44\u8baf_\u670d\u52a1\u5668\u5e2e\u52a9\u6587\u6863_\u670d\u52a1\u5668\u6559\u7a0b","potentialAction":[{"@type":"SearchAction","target":{"@type":"EntryPoint","urlTemplate":"https:\/\/www.wsisp.com\/helps\/?s={search_term_string}"},"query-input":"required name=search_term_string"}],"inLanguage":"zh-Hans"},{"@type":"Person","@id":"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/358e386c577a3ab51c4493330a20ad41","name":"admin","image":{"@type":"ImageObject","inLanguage":"zh-Hans","@id":"https:\/\/www.wsisp.com\/helps\/#\/schema\/person\/image\/","url":"https:\/\/gravatar.wp-china-yes.net\/avatar\/?s=96&d=mystery","contentUrl":"https:\/\/gravatar.wp-china-yes.net\/avatar\/?s=96&d=mystery","caption":"admin"},"sameAs":["http:\/\/wp.wsisp.com"],"url":"https:\/\/www.wsisp.com\/helps\/author\/admin"}]}},"_links":{"self":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/posts\/40335","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/users\/2"}],"replies":[{"embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/comments?post=40335"}],"version-history":[{"count":0,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/posts\/40335\/revisions"}],"wp:featuredmedia":[{"embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/media\/40324"}],"wp:attachment":[{"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/media?parent=40335"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/categories?post=40335"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/tags?post=40335"},{"taxonomy":"topic","embeddable":true,"href":"https:\/\/www.wsisp.com\/helps\/wp-json\/wp\/v2\/topic?post=40335"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}