自学内容网 自学内容网
自学内容网 > 正文

NVIDIA Nsight Compute Replay测试

🕗 发布于 2024-10-09 11:13 性能优化  性能分析  CUDA  GPU  GPGPU  

NVIDIA Nsight Compute Replay测试

本文测试NVIDIA Nsight Compute 收集同一个Kernel 不同metrics,需要执行replay的次数

一.参考链接

二.背景介绍

在NVIDIA Nsight Compute中，内核重放(Replay)是一种性能数据收集技术，用于获取GPU kernel执行期间的不同度量指标。由于GPU可以同时收集的硬件性能计数器（HW performance counters）数目有限，同时某些软件性能计数器（SW performance counters）可能会严重影响kernel的运行时间（从而影响硬件计数器的准确性），所以在收集所有需要的性能度量时，可能需要对kernel进行一次或多次重放。

内核重放过程：

  1. 数据保存与恢复:
    • 在第一次传递中，内核可以访问的所有GPU内存被保存。
    • 然后确定内核写入的内存子集。
    • 在每次重放之前（除了第一次），这部分内存子集被恢复到原来的位置，这样每次重放内核都会访问相同的内存内容。
  2. 存储位置选取优化:
    • NVIDIA Nsight Compute试图使用最快的可用存储位置来执行保存和恢复策略。例如，如果数据最初分配在设备内存中，并且设备内存还有足够空间，那么数据会直接保存在设备内存中。
    • 如果设备内存不足，数据将转移到CPU主机内存中。
    • 如果数据最初是从CPU主机内存分配的，则工具首先尝试将其保存到同一内存位置（如果可能的话）。
  3. 性能影响:
    • 根据所涉及的内存量，特别是写入操作，保存和恢复过程可能需要更多时间，并因此增加了总体开销。
    • 如果NVIDIA Nsight Compute确定只需要一次重放就足以收集所有请求的度量，则完全不执行保存和恢复操作，以减少开销。

通过这种方法，开发者可以更精确地测量和优化GPU kernel的性能，尽管这可能涉及多次重放和相应的时间增加。

三.实验现象

# 同时获取以下指标 需6 passes
lts__t_bytes.sum                                                  (单独获取该指标,只需1 passes)
lts__t_bytes_equiv_l1sectormiss.sum                               (单独获取该指标,只需1 passes)
lts__t_bytes_equiv_l1sectormiss_pipe_lsu.sum                      (单独获取该指标,只需1 passes)
-------------------------------------------------------------------
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_local_op_ld.sum      |
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_local_op_st.sum      |-->(同时获取这些指标,只需1 passes)
-------------------------------------------------------------------
lts__t_bytes_equiv_l1sectormiss_pipe_tex.sum                      (单独获取该指标,只需1 passes)
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface.sum          (单独获取该指标,只需1 passes)
-------------------------------------------------------------------
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_atom.sum  |
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_ld.sum    |
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_red.sum   |
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_st.sum    |-->(同时获取这些指标,只需1 passes)
-------------------------------------------------------------------
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture.sum          (单独获取该指标,只需1 passes)
-------------------------------------------------------------------
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_ld.sum    |
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_tex.sum   |-->(同时获取这些指标,只需1 passes)
-------------------------------------------------------------------
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_atom.sum   |
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_ld.sum     |
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_red.sum    |
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_st.sum     |-->(同时获取这些指标,需3 passes)
-------------------------------------------------------------------
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_atom.sum   (单独获取该指标,只需1 passes)
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_ld.sum     |
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_red.sum    |
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_st.sum     |-->(同时获取这些指标,只需1 passes)
-------------------------------------------------------------------

单独获取 lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_ld.sum,lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_atom.sum只需1 passes,但同时获取需要3次,原则上最大不应该超过2.这个现象怎么解释呢

可能涉及到更复杂的内部GPU行为和性能监控策略的交互。尽管从直观上认为同时获取两个需要单独1 pass的度量指标最多会需要2 passes，但如果实际需要3 passes，可能有以下几种解释：

  1. 计数器配置和资源重置：
    • 在一个pass中激活一个特定的计数器可能会影响或重置与之关联的其他硬件资源，使得在下一个pass中必须重新配置或恢复这些资源。这种重置可能在单独测试时不明显，但在尝试同时收集多个指标时变得重要。
  2. 计数器互斥与排他性要求：
    • 某些性能计数器之间可能存在互斥性，即开启一个计数器可能会阻止同一时间内开启另一个计数器。这可能是由于硬件设计或内部资源共享导致的。例如，两个指标可能都依赖于一个共享资源，但这个资源在同时请求时无法提供给两个计数器。
  3. 数据依赖和动态优化：
    • GPU的性能分析工具可能根据动态收集到的数据或内核行为调整其数据收集策略。例如，首次pass收集的数据可能指示需要进一步的细粒度分析或额外的数据验证步骤。
  4. 工具和驱动程序的实现细节：
    • 性能监控工具和GPU驱动程序的实现细节也可能影响重播次数。可能存在特定的优化或限制使得理论上的最优passes数在实际操作中不适用。

四.实验步骤

tee replay_test.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(err) << ")" << std::endl; \
            exit(EXIT_FAILURE);                                    \
        }                                                          \
    } while (0)


__global__ void kernel(float *input,float *output)
{
    unsigned int tid  = threadIdx.x + blockIdx.x * blockDim.x;
    output[tid]=input[tid];
    if(tid==0)
    {
        printf("kernel run\n");
    }
}

int main(int argc,char *argv[])
{
    int deviceid=0;
    cudaSetDevice(deviceid);  
    int block_count=1;
    int block_size=32*4;
    int thread_size=block_count*block_size;
    float *input;CHECK_CUDA(cudaMalloc(&input, thread_size*4));
    float *output;CHECK_CUDA(cudaMalloc(&output, thread_size*4));
    kernel<<<block_count, block_size>>>(input,output);
    CHECK_CUDA(cudaDeviceSynchronize());
 
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -lineinfo  -o replay_test replay_test.cu -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda

/usr/local/NVIDIA-Nsight-Compute/ncu --query-metrics | grep "lts__t_bytes" | awk '{ALL=$1".sum,"ALL}END{print ALL}' > metrics.cfg

# 6次replay
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics `cat metrics.cfg` ./replay_test

# 1次replay
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics lts__t_bytes.sum ./replay_test
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics lts__t_bytes_equiv_l1sectormiss.sum ./replay_test
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics lts__t_bytes_equiv_l1sectormiss_pipe_lsu.sum ./replay_test
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics lts__t_bytes_equiv_l1sectormiss_pipe_tex.sum ./replay_test
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics \
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_atom.sum,\
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_ld.sum,\
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_red.sum,\
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_st.sum ./replay_test

# 3次replay
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics \
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_atom.sum,\
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_ld.sum,\
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_red.sum,\
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_st.sum ./replay_test

# 1次replay
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics \
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_ld.sum  ./replay_test
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics \
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_atom.sum  ./replay_test

# 3次replay
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics \
lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_surface_op_atom.sum,\
lts__t_bytes_equiv_l1sectormiss_pipe_lsu_mem_global_op_ld.sum  ./replay_test

原文地址：https://blog.csdn.net/m0_61864577/article/details/142779448

免责声明：本站文章内容转载自网络资源，如本站内容侵犯了原著者的合法权益，可联系本站删除。更多内容请关注自学内容网（zxcms.com）！

相关文章

自学内容网
Copyright © 2015-2024