Tesla T4 P2P测试

Tesla T4 P2P测试

  • 一.测试环境
  • 二.测试步骤
    • 1.获取设备信息
    • 2.查看PCIE拓扑结构
    • 3.选择9B、9E这二张
    • 4.查看逻辑设备ID
    • 5.设置环境变量(需要用逻辑设备ID,通过UUID跟smi看到的物理ID关联)
    • 6.不同地址的原子操作
    • 2.P2P与非P2P的性能差异
    • 3.GPU带宽测试

Tesla T4 P2P测试

  • 通过物理ID找到逻辑ID
  • NCU P2P相关的Metrics
  • PCIE、DRAM相关的Metrics

一.测试环境

二.测试步骤

1.获取设备信息

nvidia-smi -L
GPU 0: NVIDIA A100 80GB PCIe (UUID: GPU-c91a8013-0877-df61-53b9-016fabcd5f82)
GPU 1: NVIDIA A30 (UUID: GPU-f67790b5-bb58-614d-4190-3598a99f925e)
GPU 2: Tesla T4 (UUID: GPU-e95bfaa3-bf41-7aeb-d1e7-4f4f98ac3a63)
GPU 3: Tesla T4 (UUID: GPU-7b470c8f-cfe3-81d2-1dd8-2e36c2552d0e)
GPU 4: Tesla T4 (UUID: GPU-d59282d2-060d-270e-1c0e-f50e936ffede)
GPU 5: NVIDIA GeForce RTX 3080 Ti (UUID: GPU-9a131b18-28de-d6a1-01e9-76a133e21680)
GPU 6: NVIDIA A30 (UUID: GPU-49daa3a5-490b-569c-f1d2-79d98c6d3a02)
GPU 7: Tesla T4 (UUID: GPU-1f45f1e1-1e10-7d2d-f25a-e79ac17ddfa1)nvidia-smi  -q | grep "Bus Id"
Bus Id                            : 00000000:34:00.0
Bus Id                            : 00000000:35:00.0
Bus Id                            : 00000000:36:00.0 Tesla T4
Bus Id                            : 00000000:37:00.0 Tesla T4
Bus Id                            : 00000000:9B:00.0 Tesla T4 d59282d2
Bus Id                            : 00000000:9C:00.0
Bus Id                            : 00000000:9D:00.0
Bus Id                            : 00000000:9E:00.0 Tesla T4 1f45f1e1

2.查看PCIE拓扑结构

lstopo --ignore Core --ignore Misc --ignore PU

在这里插入图片描述

3.选择9B、9E这二张

4.查看逻辑设备ID

tee devinfo.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>
#include <chrono>
#include <thread>#define CUDA_CHECK(call) \do { \cudaError_t error = call; \if (error != cudaSuccess) { \fprintf(stderr, "CUDA error in file '%s' in line %i: %s.\n", __FILE__, __LINE__, cudaGetErrorString(error)); \exit(EXIT_FAILURE); \} \} while (0)int main() {int device_count=0;CUDA_CHECK(cudaGetDeviceCount(&device_count));for(int deviceid=0; deviceid<device_count;deviceid++){CUDA_CHECK(cudaSetDevice(deviceid));  cudaDeviceProp deviceProp;CUDA_CHECK(cudaGetDeviceProperties(&deviceProp, deviceid));printf("Index,%d UUID:GPU-%02x%02x%02x%02x-%02x%02x-%02x%02x-%02x%02x-%02x%02x%02x%02x%02x%02x\n",deviceid,(unsigned char)deviceProp.uuid.bytes[0], (unsigned char)deviceProp.uuid.bytes[1],(unsigned char)deviceProp.uuid.bytes[2], (unsigned char)deviceProp.uuid.bytes[3],(unsigned char)deviceProp.uuid.bytes[4], (unsigned char)deviceProp.uuid.bytes[5],(unsigned char)deviceProp.uuid.bytes[6], (unsigned char)deviceProp.uuid.bytes[7],(unsigned char)deviceProp.uuid.bytes[8], (unsigned char)deviceProp.uuid.bytes[9],(unsigned char)deviceProp.uuid.bytes[10],(unsigned char)deviceProp.uuid.bytes[11],(unsigned char)deviceProp.uuid.bytes[12],(unsigned char)deviceProp.uuid.bytes[13],(unsigned char)deviceProp.uuid.bytes[14],(unsigned char)deviceProp.uuid.bytes[15]);}return 0;
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -o devinfo devinfo.cu -I /usr/local/cuda/include -L /usr/local/cuda/lib64  -lcuda
unset CUDA_VISIBLE_DEVICES && ./devinfo
  • 输出
Index,0 UUID:GPU-c91a8013-0877-df61-53b9-016fabcd5f82
Index,1 UUID:GPU-9a131b18-28de-d6a1-01e9-76a133e21680
Index,2 UUID:GPU-f67790b5-bb58-614d-4190-3598a99f925e
Index,3 UUID:GPU-49daa3a5-490b-569c-f1d2-79d98c6d3a02
Index,4 UUID:GPU-e95bfaa3-bf41-7aeb-d1e7-4f4f98ac3a63
Index,5 UUID:GPU-7b470c8f-cfe3-81d2-1dd8-2e36c2552d0e
Index,6 UUID:GPU-d59282d2-060d-270e-1c0e-f50e936ffede  #选中
Index,7 UUID:GPU-1f45f1e1-1e10-7d2d-f25a-e79ac17ddfa1  #选中

5.设置环境变量(需要用逻辑设备ID,通过UUID跟smi看到的物理ID关联)

export CUDA_VISIBLE_DEVICES=6,7

6.不同地址的原子操作

tee p2p.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>
#include <chrono>
#include <thread>#define CUDA_CHECK(call) \do { \cudaError_t error = call; \if (error != cudaSuccess) { \fprintf(stderr, "CUDA error in file '%s' in line %i: %s.\n", __FILE__, __LINE__, cudaGetErrorString(error)); \exit(EXIT_FAILURE); \} \} while (0)template<int mode>
__global__ void dummyKernel(float *data) {int idx = threadIdx.x + blockIdx.x * blockDim.x;for(int i=0;i<102400;i++){atomicAdd(&data[idx], idx*i);}
}template <typename F>
void TIMEIT(F const &f,cudaStream_t &stream,cudaEvent_t &start_ev,cudaEvent_t&stop_ev)
{ CUDA_CHECK(cudaDeviceSynchronize());auto start = std::chrono::high_resolution_clock::now();cudaEventRecord(start_ev, stream); f(stream); cudaEventRecord(stop_ev, stream); CUDA_CHECK(cudaEventSynchronize(stop_ev)); auto end = std::chrono::high_resolution_clock::now(); std::chrono::duration<double> diff = end - start; float milliseconds = 0; cudaEventElapsedTime(&milliseconds, start_ev, stop_ev); printf("E2E:%8.4fms Kernel:%8.4fms\n",diff.count()*1000,milliseconds);
}int main() {int devID0 = 0, devID1 = 1;int device_count=0;CUDA_CHECK(cudaGetDeviceCount(&device_count));for(int deviceid=0; deviceid<2;deviceid++){CUDA_CHECK(cudaSetDevice(deviceid));  cudaDeviceProp deviceProp;CUDA_CHECK(cudaGetDeviceProperties(&deviceProp, deviceid));std::cout << "-----------------------------------" << std::endl;std::cout << "Device Index: " << deviceid << std::endl;std::cout << "Compute Capability:"<<deviceProp.major<<"."<<deviceProp.minor<<std::endl;std::cout << "Device name: " << deviceProp.name << std::endl;std::cout << "Max threads per block: " << deviceProp.maxThreadsPerBlock << std::endl;std::cout << "Shared memory per block: " << deviceProp.sharedMemPerBlock << " bytes" << std::endl;std::cout << "Max blocks per SM: " << deviceProp.maxBlocksPerMultiProcessor << std::endl;std::cout << "asyncEngineCount: " << deviceProp.asyncEngineCount << std::endl;std::cout << "directManagedMemAccessFromHost: " << deviceProp.directManagedMemAccessFromHost << std::endl;std::cout << "unifiedAddressing: " << deviceProp.unifiedAddressing << std::endl;std::cout << "canMapHostMemory: " << deviceProp.canMapHostMemory << std::endl;std::cout << "Number of SMs: " << deviceProp.multiProcessorCount << std::endl;printf("UUID:GPU-%02x%02x%02x%02x-%02x%02x-%02x%02x-%02x%02x-%02x%02x%02x%02x%02x%02x\n",(unsigned char)deviceProp.uuid.bytes[0], (unsigned char)deviceProp.uuid.bytes[1],(unsigned char)deviceProp.uuid.bytes[2], (unsigned char)deviceProp.uuid.bytes[3],(unsigned char)deviceProp.uuid.bytes[4], (unsigned char)deviceProp.uuid.bytes[5],(unsigned char)deviceProp.uuid.bytes[6], (unsigned char)deviceProp.uuid.bytes[7],(unsigned char)deviceProp.uuid.bytes[8], (unsigned char)deviceProp.uuid.bytes[9],(unsigned char)deviceProp.uuid.bytes[10],(unsigned char)deviceProp.uuid.bytes[11],(unsigned char)deviceProp.uuid.bytes[12],(unsigned char)deviceProp.uuid.bytes[13],(unsigned char)deviceProp.uuid.bytes[14],(unsigned char)deviceProp.uuid.bytes[15]);}std::cout << "-----------------------------------" << std::endl;int p2p_value=0;CUDA_CHECK(cudaDeviceGetP2PAttribute(&p2p_value,cudaDevP2PAttrAccessSupported,devID0,devID1));std::cout << "cudaDevP2PAttrAccessSupported: " << p2p_value << std::endl;#define block_size (32*4)size_t dataSize = block_size * sizeof(float);float *data0_dev, *data1_dev;CUDA_CHECK(cudaSetDevice(devID0));CUDA_CHECK(cudaMalloc(&data0_dev, dataSize));CUDA_CHECK(cudaSetDevice(devID1));CUDA_CHECK(cudaMalloc(&data1_dev, dataSize));CUDA_CHECK(cudaMemset(data1_dev, 0, dataSize));// 启用P2Pint canAccessPeer=0;CUDA_CHECK(cudaDeviceCanAccessPeer(&canAccessPeer, devID0, devID1));if (canAccessPeer) {CUDA_CHECK(cudaSetDevice(devID1));cudaStream_t stream;cudaStreamCreate(&stream);cudaEvent_t start_ev, stop_ev;cudaEventCreate(&start_ev);cudaEventCreate(&stop_ev);CUDA_CHECK(cudaDeviceEnablePeerAccess(devID0, 0));//让devID1可以访问devID0的设备内存TIMEIT([&data0_dev](cudaStream_t &stream)-> void {dummyKernel<1><<<1, block_size,0,stream>>>(data0_dev);},stream,start_ev,stop_ev);TIMEIT([&data1_dev](cudaStream_t &stream)-> void {dummyKernel<3><<<1, block_size,0,stream>>>(data1_dev);},stream,start_ev,stop_ev);CUDA_CHECK(cudaDeviceDisablePeerAccess(devID0));}else{printf("%s %d canAccessPeer=0\n",__FILE__,__LINE__);}CUDA_CHECK(cudaFree(data0_dev));CUDA_CHECK(cudaFree(data1_dev));return 0;
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -o p2p p2p.cu -I /usr/local/cuda/include -L /usr/local/cuda/lib64  -lcuda
./p2p/usr/local/cuda/bin/ncu --metrics \
lts__t_requests_aperture_device.sum,\
lts__t_sectors_aperture_device.sum,\
lts__t_requests_aperture_peer.sum,\
lts__t_requests_srcnode_gpc_aperture_peer.sum,\
lts__t_requests_srcunit_l1_aperture_peer.sum,\
lts__t_requests_srcunit_tex_aperture_peer.sum,\
lts__t_sectors_aperture_peer.sum,\
lts__t_sectors_srcnode_gpc_aperture_peer.sum,\
lts__t_sectors_srcunit_l1_aperture_peer.sum,\
lts__t_sectors_srcunit_tex_aperture_peer.sum,\
dram__bytes_read.sum,\
dram__bytes_write.sum,\
dram__bytes_read.sum.per_second,\
pcie__read_bytes.sum,\
pcie__write_bytes.sum,\
pcie__read_bytes.sum.per_second,\
pcie__write_bytes.sum.per_second,\
dram__bytes_write.sum.per_second ./p2p
  • 输出
-----------------------------------
Device Index: 0
Compute Capability:7.5
Device name: Tesla T4
Max threads per block: 1024
Shared memory per block: 49152 bytes
Max blocks per SM: 16
asyncEngineCount: 3
directManagedMemAccessFromHost: 0
unifiedAddressing: 1
canMapHostMemory: 1
Number of SMs: 40
UUID:GPU-d59282d2-060d-270e-1c0e-f50e936ffede
-----------------------------------
Device Index: 1
Compute Capability:7.5
Device name: Tesla T4
Max threads per block: 1024
Shared memory per block: 49152 bytes
Max blocks per SM: 16
asyncEngineCount: 3
directManagedMemAccessFromHost: 0
unifiedAddressing: 1
canMapHostMemory: 1
Number of SMs: 40
UUID:GPU-1f45f1e1-1e10-7d2d-f25a-e79ac17ddfa1
-----------------------------------
cudaDevP2PAttrAccessSupported: 1
E2E:132.1300ms Kernel:132.1245ms  #GPU1通过P2P对GPU0的设备内存进行原子操作
E2E:  3.5552ms Kernel:  3.5444ms  #GPU1对本地DRAM进行原子操作void dummyKernel<(int)1>(float *), 2024-Sep-25 17:06:47, Context 2, Stream 34
Section: Command line profiler metrics
---------------------------------------------------------------------- --------------- ------------------------------
dram__bytes_read.sum                                                             Kbyte                          52.86
dram__bytes_read.sum.per_second                                           Kbyte/second                         405.11
dram__bytes_write.sum                                                            Kbyte                           1.25
dram__bytes_write.sum.per_second                                          Kbyte/second                           9.56
lts__t_requests_aperture_device.sum                                            request                          17775
lts__t_requests_aperture_peer.sum                                              request                         409600 #4个warp,每个102400 次合并访问 地址范围了4*32*4
lts__t_requests_srcnode_gpc_aperture_peer.sum                                  request                         409600
lts__t_requests_srcunit_l1_aperture_peer.sum                                   request                              0
lts__t_requests_srcunit_tex_aperture_peer.sum                                  request                         409600
lts__t_sectors_aperture_device.sum                                              sector                          93043
lts__t_sectors_aperture_peer.sum                                                sector                        1638400 
lts__t_sectors_srcnode_gpc_aperture_peer.sum                                    sector                        1638400 #合并访问一次请求4个sector 1638400*32=50MB
lts__t_sectors_srcunit_l1_aperture_peer.sum                                     sector                              0
lts__t_sectors_srcunit_tex_aperture_peer.sum                                    sector                        1638400
pcie__read_bytes.sum                                                             Mbyte                          59.10
pcie__read_bytes.sum.per_second                                           Mbyte/second                         452.93
pcie__write_bytes.sum                                                            Mbyte                          72.18 #实际PCIE读写加起来超过50MB
pcie__write_bytes.sum.per_second                                          Mbyte/second                         553.17
---------------------------------------------------------------------- --------------- ------------------------------void dummyKernel<(int)3>(float *), 2024-Sep-25 17:06:48, Context 2, Stream 34
Section: Command line profiler metrics
---------------------------------------------------------------------- --------------- ------------------------------
dram__bytes_read.sum                                                             Kbyte                           6.21
dram__bytes_read.sum.per_second                                           Mbyte/second                           1.77
dram__bytes_write.sum                                                             byte                            224
dram__bytes_write.sum.per_second                                          Kbyte/second                          63.93
lts__t_requests_aperture_device.sum                                            request                         414530
lts__t_requests_aperture_peer.sum                                              request                              0
lts__t_requests_srcnode_gpc_aperture_peer.sum                                  request                              0
lts__t_requests_srcunit_l1_aperture_peer.sum                                   request                              0
lts__t_requests_srcunit_tex_aperture_peer.sum                                  request                              0
lts__t_sectors_aperture_device.sum                                              sector                        1643605
lts__t_sectors_aperture_peer.sum                                                sector                              0
lts__t_sectors_srcnode_gpc_aperture_peer.sum                                    sector                              0
lts__t_sectors_srcunit_l1_aperture_peer.sum                                     sector                              0
lts__t_sectors_srcunit_tex_aperture_peer.sum                                    sector                              0
pcie__read_bytes.sum                                                             Kbyte                           3.58
pcie__read_bytes.sum.per_second                                           Mbyte/second                           1.02
pcie__write_bytes.sum                                                            Kbyte                           3.07
pcie__write_bytes.sum.per_second                                          Kbyte/second                         876.74
---------------------------------------------------------------------- --------------- ------------------------------

2.P2P与非P2P的性能差异

tee p2p.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>
#include <chrono>
#include <thread>#define CUDA_CHECK(call) \do { \cudaError_t error = call; \if (error != cudaSuccess) { \fprintf(stderr, "CUDA error in file '%s' in line %i: %s.\n", __FILE__, __LINE__, cudaGetErrorString(error)); \exit(EXIT_FAILURE); \} \} while (0)template<int mode>
__global__ void dummyKernel(float *input_data,float *output_data) {int idx = threadIdx.x + blockIdx.x * blockDim.x;output_data[idx]=input_data[idx];
}template <typename F>
void TIMEIT(F const &f,cudaStream_t &stream,cudaEvent_t &start_ev,cudaEvent_t&stop_ev)
{ CUDA_CHECK(cudaDeviceSynchronize());auto start = std::chrono::high_resolution_clock::now();cudaEventRecord(start_ev, stream); f(stream); cudaEventRecord(stop_ev, stream); CUDA_CHECK(cudaEventSynchronize(stop_ev)); auto end = std::chrono::high_resolution_clock::now(); std::chrono::duration<double> diff = end - start; float milliseconds = 0; cudaEventElapsedTime(&milliseconds, start_ev, stop_ev); printf("E2E:%7.2fms Kernel:%7.2fms\n",diff.count()*1000,milliseconds);
}int main() {int devID0 = 0, devID1 = 1;int device_count=0;CUDA_CHECK(cudaGetDeviceCount(&device_count));for(int deviceid=0; deviceid<2;deviceid++){CUDA_CHECK(cudaSetDevice(deviceid));  cudaDeviceProp deviceProp;CUDA_CHECK(cudaGetDeviceProperties(&deviceProp, deviceid));std::cout << "-----------------------------------" << std::endl;std::cout << "Device Index: " << deviceid << std::endl;std::cout << "Compute Capability:"<<deviceProp.major<<"."<<deviceProp.minor<<std::endl;std::cout << "Device name: " << deviceProp.name << std::endl;std::cout << "Max threads per block: " << deviceProp.maxThreadsPerBlock << std::endl;std::cout << "Shared memory per block: " << deviceProp.sharedMemPerBlock << " bytes" << std::endl;std::cout << "Max blocks per SM: " << deviceProp.maxBlocksPerMultiProcessor << std::endl;std::cout << "asyncEngineCount: " << deviceProp.asyncEngineCount << std::endl;std::cout << "directManagedMemAccessFromHost: " << deviceProp.directManagedMemAccessFromHost << std::endl;std::cout << "unifiedAddressing: " << deviceProp.unifiedAddressing << std::endl;std::cout << "Number of SMs: " << deviceProp.multiProcessorCount << std::endl;}std::cout << "-----------------------------------" << std::endl;int p2p_value=0;CUDA_CHECK(cudaDeviceGetP2PAttribute(&p2p_value,cudaDevP2PAttrAccessSupported,devID0,devID1));std::cout << "cudaDevP2PAttrAccessSupported: " << p2p_value << std::endl;#define block_size 1024#define block_count 1000000Lsize_t dataSize = block_count*block_size * sizeof(float);float *data0_dev, *data1_dev,*data1_dev_ex;CUDA_CHECK(cudaSetDevice(devID0));CUDA_CHECK(cudaMalloc(&data0_dev, dataSize));CUDA_CHECK(cudaSetDevice(devID1));CUDA_CHECK(cudaMalloc(&data1_dev, dataSize));CUDA_CHECK(cudaMalloc(&data1_dev_ex, dataSize));char *host;CUDA_CHECK(cudaMallocHost(&host,dataSize));printf("Init Done(%.2f)GB..\n",dataSize/1024.0/1024.0/1024.0);// 启用P2Pint canAccessPeer=0;CUDA_CHECK(cudaDeviceCanAccessPeer(&canAccessPeer, devID0, devID1));if (canAccessPeer) {CUDA_CHECK(cudaSetDevice(devID1));cudaStream_t stream;cudaStreamCreate(&stream);cudaEvent_t start_ev, stop_ev;cudaEventCreate(&start_ev);cudaEventCreate(&stop_ev);CUDA_CHECK(cudaDeviceEnablePeerAccess(devID0, 0));//让devID1可以访问devID0的设备内存TIMEIT([&](cudaStream_t &stream)-> void {cudaMemcpyAsync(host,data1_dev,dataSize,cudaMemcpyHostToDevice,stream);},stream,start_ev,stop_ev);TIMEIT([&](cudaStream_t &stream)-> void {dummyKernel<1><<<block_count, block_size,0,stream>>>(data0_dev,data1_dev);},stream,start_ev,stop_ev);TIMEIT([&](cudaStream_t &stream)-> void {dummyKernel<2><<<block_count, block_size,0,stream>>>(data1_dev_ex,data1_dev);},stream,start_ev,stop_ev);CUDA_CHECK(cudaDeviceDisablePeerAccess(devID0));}else{printf("%s %d canAccessPeer=0\n",__FILE__,__LINE__);}CUDA_CHECK(cudaFreeHost(host));CUDA_CHECK(cudaFree(data0_dev));CUDA_CHECK(cudaFree(data1_dev));CUDA_CHECK(cudaFree(data1_dev_ex));return 0;
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -o p2p p2p.cu -I /usr/local/cuda/include -L /usr/local/cuda/lib64  -lcuda
./p2p/usr/local/NVIDIA-Nsight-Compute/ncu --metrics \
dram__bytes_read.sum.pct_of_peak_sustained_elapsed,\
dram__bytes_write.sum.pct_of_peak_sustained_elapsed,\
dram__bytes_read.sum.per_second,\
pcie__read_bytes.sum.per_second,\
pcie__write_bytes.sum.per_second,\
dram__bytes_write.sum.per_second ./p2p
  • 输出
-----------------------------------
Device Index: 0
Compute Capability:7.5
Device name: Tesla T4
Max threads per block: 1024
Shared memory per block: 49152 bytes
Max blocks per SM: 16
asyncEngineCount: 3
directManagedMemAccessFromHost: 0
unifiedAddressing: 1
Number of SMs: 40
-----------------------------------
Device Index: 1
Compute Capability:7.5
Device name: Tesla T4
Max threads per block: 1024
Shared memory per block: 49152 bytes
Max blocks per SM: 16
asyncEngineCount: 3
directManagedMemAccessFromHost: 0
unifiedAddressing: 1
Number of SMs: 40
-----------------------------------
cudaDevP2PAttrAccessSupported: 1
Init Done(3.81)GB..
E2E: 325.70ms Kernel: 325.71ms  # GPU1 cudaMemcpyHostToDevice  11.697GB/s
E2E: 307.29ms Kernel: 307.31ms  # GPU1 通过P2P从GPU0的设备内存读取数据(比H2D快) 12.39GB/s
E2E:  37.90ms Kernel:  37.89ms  # GPU1 DRAM内 D2D的拷贝 2*100.55GB/svoid dummyKernel<1>(float *, float *) (1000000, 1, 1)x(1024, 1, 1), Context 2, Stream 34, Device 7, CC 7.5
Section: Command line profiler metrics
--------------------------------------------------- ----------- ------------
Metric Name                                         Metric Unit Metric Value
--------------------------------------------------- ----------- ------------
dram__bytes_read.sum.pct_of_peak_sustained_elapsed            %         0.01
dram__bytes_read.sum.per_second                         Mbyte/s        37.35
dram__bytes_write.sum.pct_of_peak_sustained_elapsed           %         4.72
dram__bytes_write.sum.per_second                        Gbyte/s        15.10
pcie__read_bytes.sum.per_second                         Gbyte/s        15.01
pcie__write_bytes.sum.per_second                        Gbyte/s         3.34
--------------------------------------------------- ----------- ------------void dummyKernel<2>(float *, float *) (1000000, 1, 1)x(1024, 1, 1), Context 2, Stream 34, Device 7, CC 7.5
Section: Command line profiler metrics
--------------------------------------------------- ----------- ------------
Metric Name                                         Metric Unit Metric Value
--------------------------------------------------- ----------- ------------
dram__bytes_read.sum.pct_of_peak_sustained_elapsed            %        37.34  #同时读写时,利用率加起来75%
dram__bytes_read.sum.per_second                         Gbyte/s       119.41
dram__bytes_write.sum.pct_of_peak_sustained_elapsed           %        37.81
dram__bytes_write.sum.per_second                        Gbyte/s       120.89  #加起来239GB/s,跟后面的带宽测试一致
pcie__read_bytes.sum.per_second                         Mbyte/s        22.03
pcie__write_bytes.sum.per_second                        Mbyte/s         7.42
--------------------------------------------------- ----------- ------------

3.GPU带宽测试

git clone https://www.github.com/nvidia/cuda-samples
cd cuda-samples/Samples/1_Utilities/deviceQuery
make clean && make
./deviceQuery
cd ../bandwidthTest/
make clean && make
./bandwidthTest --device=0
  • 输出
Running on...Device 0: Tesla T4Quick ModeHost to Device Bandwidth, 1 Device(s)PINNED Memory TransfersTransfer Size (Bytes)        Bandwidth(GB/s)32000000                     12.8Device to Host Bandwidth, 1 Device(s)PINNED Memory TransfersTransfer Size (Bytes)        Bandwidth(GB/s)32000000                     13.1Device to Device Bandwidth, 1 Device(s)PINNED Memory TransfersTransfer Size (Bytes)        Bandwidth(GB/s)32000000                     239.4
Result = PASS

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

如若内容造成侵权/违法违规/事实不符，请联系长河编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

多个ECU测试方案-IP地址相同-DoIP刷新-环境测试耐久测试

多个ECU测试方案-IP地址相同-DoIP刷新-环境测试耐久测试

情况1&#xff1a;只有一个ECU进行测试 - 接口模块只需要使用一个车载以太网转换器&#xff1b; 情况2&#xff1a;多ECU同时测试&#xff0c;但ECU IP地址不一样&#xff0c;上位机多个网口 - 上位机测试软件&#xff0c;需要通过PC的不同网卡&#xff0c;访问各个ECU&#…
阅读更多...
微信小程序开发第九课

微信小程序开发第九课

一 后端上线 1.1 购买云服务器 1.2 安装python3.9 # 阿里云的centos上有python环境- python2.7.5 pip- python3.6.8 pip3-咱们项目开发&#xff0c;在3.9上开发的&#xff0c;需要使用3.9的解释器来运行# 可以使用yum 安装&#xff0c;不能指定版本&#xff08;yum i…
阅读更多...
CorePress Pro 网站加载慢 WordPress

CorePress Pro 网站加载慢 WordPress

一般来说是你用了「CorePress天气模块」 解决方案&#xff1a;这个插件从你右侧边栏里删掉就可以了&#xff08;上方的图中已经是删掉后的效果了&#xff09; 寻找加载时间长的原因&#xff1a; 谷歌浏览器F12->网络->打开录制->ShiftF5 得出结论&#xff1a;和风天气…
阅读更多...
Android 车载应用开发指南 - CarService 详解（下）

Android 车载应用开发指南 - CarService 详解（下）

车载应用正在改变人们的出行体验。从导航到娱乐、从安全到信息服务&#xff0c;车载应用的开发已成为汽车智能化发展的重要组成部分。而对于开发者来说&#xff0c;如何将自己的应用程序无缝集成到车载系统中&#xff0c;利用汽车的硬件和服务能力&#xff0c;是一个极具挑战性…
阅读更多...
计算机网络：物理层 --- 基本概念、编码与调制

计算机网络：物理层 --- 基本概念、编码与调制

目录 一. 物理层的基本概念 二. 数据通信系统的模型 三. 编码 3.1 基本概念 3.2 不归零制编码 3.3 归零制编码 3.4 曼切斯特编码 3.5 差分曼切斯特编码 ​编辑 四. 调制 4.1 调幅 4.2 调频 4.3 调相 4.4 混合调制 今天我们讲的是物理…
阅读更多...
【JavaEE】——线程的安全问题和解决方式

【JavaEE】——线程的安全问题和解决方式

阿华代码&#xff0c;不是逆风&#xff0c;就是我疯&#xff0c;你们的点赞收藏是我前进最大的动力&#xff01;&#xff01;希望本文内容能够帮助到你&#xff01; 目录 一&#xff1a;问题引入 二&#xff1a;问题深入 1&#xff1a;举例说明 2&#xff1a;图解双线程计算…
阅读更多...
机器学习04-逻辑回归（python）-02原理与损失函数

机器学习04-逻辑回归（python）-02原理与损失函数

​​​​​​​ 1. 逻辑回归概念 逻辑回归&#xff08;Logistic Regression&#xff09; 是一种 分类模型&#xff0c;主要用于解决 二分类问题&#xff08;即分成两类&#xff0c;如是否通过、是否患病等&#xff09;。逻辑回归的目标是根据输入的特征预测一个 概率&#xff0…
阅读更多...
AI大模型项目实战v0.2: 结合个人知识库

AI大模型项目实战v0.2: 结合个人知识库

前言 在AI大模型项目实战v0.1版本中&#xff0c;我们实现了一个最简单的基于纯LLM的问答机器人Tbot。 今天升级到v0.2版本&#xff0c;结合个人知识库。 本系列每个版本&#xff0c;都将提供完整的代码文档&#xff0c;获取方法见文末。 下面开启我们的v0.2版本之旅。 v0.2 Tb…
阅读更多...
Icarus翼星求生教你使用服务器开服

Icarus翼星求生教你使用服务器开服

1、购买后登录服务器&#xff08;百度莱卡云游戏面板&#xff09; 登录面板的信息在绿色的登陆面板按键下方&#xff0c;不是你的莱卡云账号 进入控制面板后会出现正在安装的界面&#xff0c;大约10分钟左右就能安装完成 2、创建端口 点击目录上的网络&#xff0c;再次页面下点…
阅读更多...
中伟视界：AI算法如何精准识别井下与传送带上堆料，提升矿山安全生产效率，减少事故风险

中伟视界：AI算法如何精准识别井下与传送带上堆料，提升矿山安全生产效率，减少事故风险

传送带堆料分为两种情况&#xff0c;一种是传送带的井下堆料检测AI算法&#xff0c;一种是传送带上面的堆料检测AI算法&#xff0c;传送带井下堆料检测AI算法是在带式输送机的漏煤下方井下安装摄像仪&#xff0c;通过视频分析检测井下堆煤情况&#xff0c;当洒煤堆积到一定程度…
阅读更多...
【Git入门】使用 Git 进行项目管理：Word Count 程序开发与托管

【Git入门】使用 Git 进行项目管理：Word Count 程序开发与托管

在软件开发过程中&#xff0c;版本控制工具是不可或缺的。Git 作为一款强大的分布式版本控制工具&#xff0c;为开发者提供了高效的代码管理和协作方式。本博客将介绍如何下载安装 Git 版本管理工具&#xff0c;并使用 Git 和 GitHub 平台进行一个名为 Word Count 的项目开发与…
阅读更多...
二分

二分

LeetCode34 在排序数组中查找元素的第一个和最后一个位置&#xff08;二分模板题&#xff0c;左闭右开写法&#xff09; /** lc appleetcode.cn id34 langcpp** [34] 在排序数组中查找元素的第一个和最后一个位置*/// lc codestart #include<iostream> using namespace s…
阅读更多...
Python发送邮件教程：如何实现自动化发信？

Python发送邮件教程：如何实现自动化发信？

Python发送邮件有哪些方法&#xff1f;如何利用python发送邮件&#xff1f; 无论是工作汇报、客户通知还是个人提醒&#xff0c;邮件都能快速传递信息。Python发送邮件的自动化功能就显得尤为重要。AokSend将详细介绍如何使用Python发送邮件&#xff0c;实现自动化发信&#x…
阅读更多...
逆向推理+ChatGPT，让论文更具说服力

逆向推理+ChatGPT，让论文更具说服力

学境思源&#xff0c;一键生成论文初稿&#xff1a; AcademicIdeas - 学境思源AI论文写作 使用ChatGPT辅助“逆向推理”技巧&#xff0c;可以显著提升论文的质量和说服力。逆向推理从结论出发&#xff0c;倒推所需的证据和论点&#xff0c;确保整个论证过程逻辑严密且无漏洞。…
阅读更多...
Spring Cloud ：Hystrix实现优雅的服务容错

Spring Cloud ：Hystrix实现优雅的服务容错

目录 Hystrix概述&#xff1a;第一个Hystrix程序步骤1&#xff1a;创建父工程hystrix-1步骤2&#xff1a;改造服务提供者步骤3&#xff1a;改造服务消费者为Hystrix客户端&#xff08;1&#xff09;添加Hystrix依赖&#xff08;2&#xff09;添加EnableHystrix注解&#xff08;…
阅读更多...
编程练习2 数据单元的变量替换

编程练习2 数据单元的变量替换

示例1: 1,2<A>00 示例2: 1,2<A>00,3<A>00 示例3: <B>12,1,2<B>1 示例4: <B<12,1 输出依次如下&#xff1a; #include<iostream> #include<vector> #include<string>using namespace std;/* 字符分割函数 将传入…
阅读更多...
人工智能-大语言模型-微调技术-LoRA及背后原理简介

人工智能-大语言模型-微调技术-LoRA及背后原理简介

1. 《LORA: LOW-RANK ADAPTATION OF LARGE LANGUAGE MODELS》 LORA: 大型语言模型的低秩适应 摘要&#xff1a; 随着大规模预训练模型的发展&#xff0c;全参数微调变得越来越不可行。本文提出了一种名为LoRA&#xff08;低秩适应&#xff09;的方法&#xff0c;通过在Transf…
阅读更多...
STM32 使用 CubeMX 实现按键外部中断

STM32 使用 CubeMX 实现按键外部中断

目录 问题背景知识参考需要改什么注意尽量不要在中断函数使用 循环函数做延时中断函数中延时方法调试 问题 我想实现按钮触发紧急停止类似功能&#xff0c;需要使用按键中断功能。 背景知识 GPIO 点亮 LED。stm32cubemx hal学习记录&#xff1a;GPIO输入输出。STM32—HAL库 …
阅读更多...
活动系统开发之采用设计模式与非设计模式的区别-后台功能总结

活动系统开发之采用设计模式与非设计模式的区别-后台功能总结

1、数据库ER图 2、后台功能字段 题目功能字段 数据列表 编号题目名称选项数量状态 1启用0禁用创建时间修改时间保存 题目名称选项集 选项内容是否正确答案 1正确0错误启禁用删除素材图库功能字段 数据列表 编号原文件名称文件类型文件大小加密后文件名文件具体路径上传类型状态…
阅读更多...
从零开始学习Python

从零开始学习Python

目录 从零开始学习Python 引言 环境搭建 安装Python解释器 选择IDE 基础语法 注释 变量和数据类型 变量命名规则 数据类型 运算符 算术运算符 比较运算符 逻辑运算符 输入和输出 控制流 条件语句 循环语句 for循环 while循环 循环控制语句 函数和模块 定…
阅读更多...
最新文章
推荐文章

Copyright @ 2022~2023