测试NV GPU SM的时钟是否一致

2024-08-23 20:36
文章标签 时钟 是否 测试 gpu 一致 sm nv

本文主要是介绍测试NV GPU SM的时钟是否一致,希望对大家解决编程问题提供一定的参考价值,需要的开发者们随着小编来一起学习吧!

测试NV GPU SM的时钟是否一致

  • 操作步骤

测试NV GPU SM的时钟是否一致

操作步骤

tee sm_clock_benchmark.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>#define CHECK_CUDA(call)                                           \do {                                                           \cudaError_t err = call;                                    \if (err != cudaSuccess) {                                  \std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__;  \std::cerr << " code=" << err << " (" << cudaGetErrorString(cudaGetLastError()) << ")" << std::endl; \}                                                          \} while (0)__global__ void kernel(unsigned long long*output_ts,unsigned int*output_smid) {int tid  = threadIdx.x + blockIdx.x * blockDim.x;unsigned long long ts0=0;asm volatile ("mov.u64 %0, %clock64;" : "=l"(ts0) :: "memory");unsigned int smid;asm volatile("mov.u32 %0, %smid;" : "=r"(smid));if(tid%blockDim.x==0){output_ts[blockIdx.x]=ts0;output_smid[blockIdx.x]=smid;}
}int main(int argc,char *argv[])
{int deviceid=0;cudaSetDevice(deviceid);  cudaDeviceProp deviceProp;cudaGetDeviceProperties(&deviceProp, deviceid);int maxThreadsPerBlock = deviceProp.maxThreadsPerBlock;int sharedMemoryPerBlock = deviceProp.sharedMemPerBlock;int maxBlocksPerMultiprocessor = deviceProp.maxBlocksPerMultiProcessor;int smCount = deviceProp.multiProcessorCount;std::cout << "Device name: " << deviceProp.name << std::endl;std::cout << "Max threads per block: " << maxThreadsPerBlock << std::endl;std::cout << "Shared memory per block: " << sharedMemoryPerBlock << " bytes" << std::endl;std::cout << "Max blocks per SM: " << maxBlocksPerMultiprocessor << std::endl;std::cout << "Number of SMs: " << smCount << std::endl;int block_size=smCount;int thread_block_size=maxThreadsPerBlock;int thread_size=thread_block_size*block_size;int data_size=sizeof(float)*thread_size;int ts_size=sizeof(unsigned long long)*thread_size;int smid_size=sizeof(int)*thread_size;unsigned long long* dev_output_ts=nullptr;unsigned int* dev_smid=nullptr;unsigned long long*host_output_ts=new unsigned long long[thread_size];;unsigned int* host_smid=new unsigned int[thread_size];CHECK_CUDA(cudaMalloc((void**)&dev_output_ts, ts_size));CHECK_CUDA(cudaMalloc((void**)&dev_smid, smid_size));CHECK_CUDA(cudaMemcpy(dev_output_ts,host_output_ts,ts_size,cudaMemcpyHostToDevice));CHECK_CUDA(cudaMemcpy(dev_smid,host_smid,smid_size,cudaMemcpyHostToDevice));printf("dev_output_ts:%p\n",dev_output_ts);printf("dev_smid:%p\n",dev_smid);cudaStream_t stream;cudaStreamCreate(&stream);cudaEvent_t start, stop;cudaEventCreate(&start);cudaEventCreate(&stop);for(int iter=0;iter<3;iter++){cudaEventRecord(start, stream);    kernel<<<block_size, thread_block_size,sharedMemoryPerBlock,stream>>>(dev_output_ts,dev_smid);       cudaEventRecord(stop, stream);CHECK_CUDA(cudaEventSynchronize(stop));float milliseconds = 0;cudaEventElapsedTime(&milliseconds, start, stop);printf("cudaEventElapsedTime:%d %.3f(milliseconds)\n",iter,milliseconds);CHECK_CUDA(cudaMemcpy(host_output_ts,dev_output_ts,ts_size,cudaMemcpyDeviceToHost));CHECK_CUDA(cudaMemcpy(host_smid,dev_smid,smid_size,cudaMemcpyDeviceToHost));unsigned long long _min=0;unsigned long long _max=0;for(int i=0;i<block_size;i++){if(_min==0) _min=host_output_ts[i];if(_max==0) _max=host_output_ts[i];if(host_output_ts[i]<_min){_min=host_output_ts[i];}if(host_output_ts[i]>_max){_max=host_output_ts[i];}printf("blockid:%04d ts:%lld smid:%d\n",i,host_output_ts[i],host_smid[i]);}unsigned long long diff=_max-_min;printf("_max-_min=%lld(cycles) %6.2f(sec)\n",diff,diff/(1.89*1e9));    }CHECK_CUDA(cudaFree(dev_smid));CHECK_CUDA(cudaFree(dev_output_ts));return 0;
}
EOF/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -g -lineinfo -o sm_clock_benchmark sm_clock_benchmark.cu \-I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./sm_clock_benchmark

输出

cudaEventElapsedTime:2 0.006(milliseconds)
blockid:0000 ts:3642438400169 smid:0
blockid:0001 ts:3644393850856 smid:2
blockid:0002 ts:3646612108206 smid:4
blockid:0003 ts:3642438400201 smid:6
blockid:0004 ts:3644393850888 smid:8
blockid:0005 ts:3646612108190 smid:10
blockid:0006 ts:3642438400234 smid:12
blockid:0007 ts:3644393850921 smid:14
blockid:0008 ts:3646612108239 smid:16
blockid:0009 ts:3642438400184 smid:18
blockid:0010 ts:3644393850871 smid:20
blockid:0011 ts:3646612108221 smid:22
blockid:0012 ts:3642438400216 smid:24
blockid:0013 ts:3644393850903 smid:26
blockid:0014 ts:3642438400177 smid:1
blockid:0015 ts:3644393850864 smid:3
blockid:0016 ts:3646612108214 smid:5
blockid:0017 ts:3642438400209 smid:7
blockid:0018 ts:3644393850896 smid:9
blockid:0019 ts:3646612108198 smid:11
blockid:0020 ts:3642438400242 smid:13
blockid:0021 ts:3644393850929 smid:15
blockid:0022 ts:3646612108247 smid:17
blockid:0023 ts:3642438400192 smid:19
blockid:0024 ts:3644393850879 smid:21
blockid:0025 ts:3646612108229 smid:23
blockid:0026 ts:3642438400224 smid:25
blockid:0027 ts:3644393850911 smid:27
_max-_min=4173708078(cycles)   2.21(sec)

这篇关于测试NV GPU SM的时钟是否一致的文章就介绍到这儿,希望我们推荐的文章对编程师们有所帮助!



http://www.chinasem.cn/article/1100434

相关文章

Python中Tensorflow无法调用GPU问题的解决方法

《Python中Tensorflow无法调用GPU问题的解决方法》文章详解如何解决TensorFlow在Windows无法识别GPU的问题,需降级至2.10版本,安装匹配CUDA11.2和cuDNN... 当用以下代码查看GPU数量时,gpuspython返回的是一个空列表,说明tensorflow没有找到

使用Python进行GRPC和Dubbo协议的高级测试

《使用Python进行GRPC和Dubbo协议的高级测试》GRPC(GoogleRemoteProcedureCall)是一种高性能、开源的远程过程调用(RPC)框架,Dubbo是一种高性能的分布式服... 目录01 GRPC测试安装gRPC编写.proto文件实现服务02 Dubbo测试1. 安装Dubb

Python的端到端测试框架SeleniumBase使用解读

《Python的端到端测试框架SeleniumBase使用解读》:本文主要介绍Python的端到端测试框架SeleniumBase使用,具有很好的参考价值,希望对大家有所帮助,如有错误或未考虑完全... 目录SeleniumBase详细介绍及用法指南什么是 SeleniumBase?SeleniumBase

python判断文件是否存在常用的几种方式

《python判断文件是否存在常用的几种方式》在Python中我们在读写文件之前,首先要做的事情就是判断文件是否存在,否则很容易发生错误的情况,:本文主要介绍python判断文件是否存在常用的几种... 目录1. 使用 os.path.exists()2. 使用 os.path.isfile()3. 使用

Java中JSON格式反序列化为Map且保证存取顺序一致的问题

《Java中JSON格式反序列化为Map且保证存取顺序一致的问题》:本文主要介绍Java中JSON格式反序列化为Map且保证存取顺序一致的问题,具有很好的参考价值,希望对大家有所帮助,如有错误或未... 目录背景问题解决方法总结背景做项目涉及两个微服务之间传数据时,需要提供方将Map类型的数据序列化为co

conda安装GPU版pytorch默认却是cpu版本

《conda安装GPU版pytorch默认却是cpu版本》本文主要介绍了遇到Conda安装PyTorchGPU版本却默认安装CPU的问题,文中通过示例代码介绍的非常详细,对大家的学习或者工作具有一定的... 目录一、问题描述二、网上解决方案罗列【此节为反面方案罗列!!!】三、发现的根本原因[独家]3.1 p

python多线程并发测试过程

《python多线程并发测试过程》:本文主要介绍python多线程并发测试过程,具有很好的参考价值,希望对大家有所帮助,如有错误或未考虑完全的地方,望不吝赐教... 目录一、并发与并行?二、同步与异步的概念?三、线程与进程的区别?需求1:多线程执行不同任务需求2:多线程执行相同任务总结一、并发与并行?1、

Python如何判断字符串中是否包含特殊字符并替换

《Python如何判断字符串中是否包含特殊字符并替换》这篇文章主要为大家详细介绍了如何使用Python实现判断字符串中是否包含特殊字符并使用空字符串替换掉,文中的示例代码讲解详细,感兴趣的小伙伴可以了... 目录python判断字符串中是否包含特殊字符方法一:使用正则表达式方法二:手动检查特定字符Pytho

判断PyTorch是GPU版还是CPU版的方法小结

《判断PyTorch是GPU版还是CPU版的方法小结》PyTorch作为当前最流行的深度学习框架之一,支持在CPU和GPU(NVIDIACUDA)上运行,所以对于深度学习开发者来说,正确识别PyTor... 目录前言为什么需要区分GPU和CPU版本?性能差异硬件要求如何检查PyTorch版本?方法1:使用命

Python如何精准判断某个进程是否在运行

《Python如何精准判断某个进程是否在运行》这篇文章主要为大家详细介绍了Python如何精准判断某个进程是否在运行,本文为大家整理了3种方法并进行了对比,有需要的小伙伴可以跟随小编一起学习一下... 目录一、为什么需要判断进程是否存在二、方法1:用psutil库(推荐)三、方法2:用os.system调用