RTX3060 FP64测试与猜想

RTX3060 FP64测试与猜想

一.小结

• RTX3060 compute capability为8.6,每个SM有2个FP64 core。每个cycle可输出2个fp64的结果
• RTX3060 有4个subcore,这2个core怎么给4个sub_core分呢
• 执行FP64 DADD指令时,MIO PQ利用率超20%(FADD指令不存在该现象),且fp64 pipe的利用率最多为84%
• 每个smsp 执行一条DADD warp指令 pipe_fp64_cycles_active 增加16个cycle,4个smsp一起运行一条DADD warp指令仍是16个cycle
• 猜测:
  • smsp按 1DADD/cycle 交替发送给2个FP64 core,一个warp需要16个cycle(32inst/16cycle->2inst/cycle)
  • 如果4个sub core同时按这个速度发,则超过了FP64的处理能力(8inst/cycle > 2inst/cycle),但pipe_fp64_cycles_active没有增加
  • 说明,在发射FP64指令之前会检测资源的可用性,如果不足,则不发射,pipe_fp64_cycles_active也就不会增加
  • 也就解释了4个sub core一起执行时,pipe_fp64_cycles_active.max还是16个cycle
  • 执行FP64指令时,4个subcore通过MIO共享FP64实际的执行单元

二.查看FP64的峰值性能

tee fp64_peak_sustained.cu<<-'EOF'
#include <cuda_runtime.h>
#include <cuda.h>
__global__ void fake_kernel(){}
int main(int argc,char *argv[])
{fake_kernel<<<1, 1>>>();cudaDeviceSynchronize();
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -lineinfo  -o fp64_peak_sustained fp64_peak_sustained.cu \-I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics \
sm__sass_thread_inst_executed_op_fp64_pred_on.avg.peak_sustained,\
sm__sass_thread_inst_executed_op_fp64_pred_on.sum.peak_sustained ./fp64_peak_sustained

输出

fake_kernel() (1, 1, 1)x(1, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
---------------------------------------------------------------- ----------- ------------
Metric Name                                                      Metric Unit Metric Value
---------------------------------------------------------------- ----------- ------------
sm__sass_thread_inst_executed_op_fp64_pred_on.avg.peak_sustained  inst/cycle            2  #每个sm的峰值性能
sm__sass_thread_inst_executed_op_fp64_pred_on.sum.peak_sustained  inst/cycle           56  #28个sm
---------------------------------------------------------------- ----------- ------------

• 2 FP64 cores in devices of compute capability 8.6, 8.7 and 8.9
• 问题:这2个core怎么给4个sub_core分呢?

三.打满FP64、FP32的利用率,对比差异

tee fp64_test.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>__global__ void kernel_add_float(volatile float *input,volatile float *output)
{unsigned int tid  = threadIdx.x + blockIdx.x * blockDim.x;float l=input[tid];float r=output[tid];for(int i=0;i<256;i++){l-=r;} input[tid]=l;
}
__global__ void kernel_add_double(volatile double *input,volatile double *output)
{unsigned int tid  = threadIdx.x + blockIdx.x * blockDim.x;double left=input[tid];double right=output[tid];for(int i=0;i<256;i++){left+=right;}       output[tid]=left;
}
EOF/usr/local/cuda/bin/nvcc -std=c++17 -dc -lineinfo -arch=sm_86 -ptx fp64_test.cu -o fp64_test.ptx
/usr/local/cuda/bin/nvcc -arch=sm_86 fp64_test.ptx -cubin -o fp64_test.cubin
/usr/local/cuda/bin/nvcc -arch=sm_86 fp64_test.cubin -fatbin -o fp64_test.fatbin
cat fp64_test.ptx
/usr/local/cuda/bin/cuobjdump --dump-sass fp64_test.fatbin# 删掉除DADD、FADD以外的指令
cuasm.py fp64_test.cubin fp64_test.cuasm
sed '/MOV/d' -i fp64_test.cuasm
sed '/S2R/d' -i fp64_test.cuasm
sed '/ULDC/d' -i fp64_test.cuasm
sed '/IMAD/d' -i fp64_test.cuasm
sed '/LDG/d' -i fp64_test.cuasm
sed '/STG/d' -i fp64_test.cuasm
sed '/F2F/d' -i fp64_test.cuasmcuasm.py fp64_test.cuasm
/usr/local/cuda/bin/nvcc -arch=sm_86 fp64_test.cubin -fatbin -o fp64_test.fatbin
/usr/local/cuda/bin/cuobjdump --dump-sass fp64_test.fatbin
/usr/local/cuda/bin/cuobjdump --dump-resource-usage fp64_test.fatbintee fp64_test_main.cpp<<-'EOF'
#include <stdio.h>
#include <string.h>
#include <cuda_runtime.h>
#include <cuda.h>int main(int argc,char *argv[])
{CUresult error;CUdevice cuDevice;cuInit(0);int deviceCount = 0;error = cuDeviceGetCount(&deviceCount);error = cuDeviceGet(&cuDevice, 0);if(error!=CUDA_SUCCESS){printf("Error happened in get device!\n");}CUcontext cuContext;error = cuCtxCreate(&cuContext, 0, cuDevice);if(error!=CUDA_SUCCESS){printf("Error happened in create context!\n");}int block_count=28*1000;int block_size=32*4*4;int thread_size=block_count*block_size;int data_size=sizeof(double)*thread_size;double *output_ptr=nullptr;double *input_ptr=nullptr;int cudaStatus=0;cudaStatus = cudaMalloc((void**)&input_ptr, data_size);cudaStatus = cudaMalloc((void**)&output_ptr, data_size);void *kernelParams[]= {(void*)&output_ptr, (void*)&input_ptr};CUmodule module;CUfunction double_function;CUfunction float_function;const char* module_file = "fp64_test.fatbin";const char* double_kernel_name = "_Z17kernel_add_doublePVdS0_";const char* float_kernel_name = "_Z16kernel_add_floatPVfS0_";error = cuModuleLoad(&module, module_file);if(error!=CUDA_SUCCESS){printf("Error happened in load moudle %d!\n",error);}error = cuModuleGetFunction(&double_function, module, double_kernel_name);if(error!=CUDA_SUCCESS){printf("get double_function error!\n");}error = cuModuleGetFunction(&float_function, module, float_kernel_name);if(error!=CUDA_SUCCESS){printf("get float_kernel_name error!\n");}    cuLaunchKernel(double_function,block_count, 1, 1,block_size, 1, 1,0,0,kernelParams, 0);cuLaunchKernel(float_function,block_count, 1, 1,block_size, 1, 1,0,0,kernelParams, 0);cudaFree(output_ptr);cudaFree(input_ptr);cuModuleUnload(module);cuCtxDestroy(cuContext);return 0;
}
EOF
g++ fp64_test_main.cpp -o fp64_test_main -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcudart -lcuda
/usr/local/NVIDIA-Nsight-Compute/ncu --metrics \
sm__inst_executed.avg.pct_of_peak_sustained_elapsed,\
smsp__inst_issued.sum,\
sm__issue_active.avg.pct_of_peak_sustained_elapsed,\
sm__pipe_fma_cycles_active.avg.pct_of_peak_sustained_elapsed,\
sm__pipe_fmaheavy_cycles_active.avg.pct_of_peak_sustained_elapsed,\
sm__inst_executed_pipe_cbu_pred_on_any.avg.pct_of_peak_sustained_elapsed,\
sm__mio_inst_issued.avg.pct_of_peak_sustained_elapsed,\
sm__mio_pq_read_cycles_active.avg.pct_of_peak_sustained_elapsed,\
sm__mio_pq_write_cycles_active.avg.pct_of_peak_sustained_elapsed,\
sm__mio_pq_write_cycles_active_pipe_lsu.avg.pct_of_peak_sustained_elapsed,\
sm__mio_pq_write_cycles_active_pipe_tex.avg.pct_of_peak_sustained_elapsed,\
sm__mioc_inst_issued.avg.pct_of_peak_sustained_elapsed,\
sm__mio_inst_issued.avg.pct_of_peak_sustained_elapsed,\
sm__pipe_fp64_cycles_active.avg.pct_of_peak_sustained_elapsed ./fp64_test_main

输出

kernel_add_double(volatile double *, volatile double *) (28000, 1, 1)x(512, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
------------------------------------------------------------------------- ----------- ------------
Metric Name                                                               Metric Unit Metric Value
------------------------------------------------------------------------- ----------- ------------
sm__inst_executed.avg.pct_of_peak_sustained_elapsed                                 %         1.32
sm__inst_executed_pipe_cbu_pred_on_any.avg.pct_of_peak_sustained_elapsed            %         0.02
sm__issue_active.avg.pct_of_peak_sustained_elapsed                                  %         1.32
sm__mio_inst_issued.avg.pct_of_peak_sustained_elapsed                               %         3.51
sm__mio_pq_read_cycles_active.avg.pct_of_peak_sustained_elapsed                     %            0
sm__mio_pq_write_cycles_active.avg.pct_of_peak_sustained_elapsed                    %        21.05
sm__mio_pq_write_cycles_active_pipe_lsu.avg.pct_of_peak_sustained_elapsed           %            0
sm__mio_pq_write_cycles_active_pipe_tex.avg.pct_of_peak_sustained_elapsed           %        21.05 # of cycles where register operands from the register file werewritten to MIO PQ, for the tex pipe
sm__mioc_inst_issued.avg.pct_of_peak_sustained_elapsed                              %         1.32
sm__pipe_fma_cycles_active.avg.pct_of_peak_sustained_elapsed                        %            0
sm__pipe_fmaheavy_cycles_active.avg.pct_of_peak_sustained_elapsed                   %            0
sm__pipe_fp64_cycles_active.avg.pct_of_peak_sustained_elapsed                       %        84.21  #利用率打不满
smsp__inst_issued.sum                                                            inst  115,136,000  #跟fp32相同的指令条数
------------------------------------------------------------------------- ----------- ------------kernel_add_float(volatile float *, volatile float *) (28000, 1, 1)x(512, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
------------------------------------------------------------------------- ----------- ------------
Metric Name                                                               Metric Unit Metric Value
------------------------------------------------------------------------- ----------- ------------
sm__inst_executed.avg.pct_of_peak_sustained_elapsed                                 %        99.76
sm__inst_executed_pipe_cbu_pred_on_any.avg.pct_of_peak_sustained_elapsed            %         1.55
sm__issue_active.avg.pct_of_peak_sustained_elapsed                                  %        99.76
sm__mio_inst_issued.avg.pct_of_peak_sustained_elapsed                               %            0
sm__mio_pq_read_cycles_active.avg.pct_of_peak_sustained_elapsed                     %            0
sm__mio_pq_write_cycles_active.avg.pct_of_peak_sustained_elapsed                    %            0
sm__mio_pq_write_cycles_active_pipe_lsu.avg.pct_of_peak_sustained_elapsed           %            0
sm__mio_pq_write_cycles_active_pipe_tex.avg.pct_of_peak_sustained_elapsed           %            0
sm__mioc_inst_issued.avg.pct_of_peak_sustained_elapsed                              %         0.39
sm__pipe_fma_cycles_active.avg.pct_of_peak_sustained_elapsed                        %        99.37
sm__pipe_fmaheavy_cycles_active.avg.pct_of_peak_sustained_elapsed                   %        99.37
sm__pipe_fp64_cycles_active.avg.pct_of_peak_sustained_elapsed                       %            0
smsp__inst_issued.sum                                                            inst  115,136,000
------------------------------------------------------------------------- ----------- ------------

*猜测,sm__pipe_fp64_cycles_active并不是那2个FP64 core的metrics,而是smsp里通往fp64 core的接口模块的活动cycle数
*4个subcore里的fp64接口模块,连接到2个fp64 core,并且经过了mio模块.因此,无法打满fp64的利用率

四.进一步证明pipe_fp64_cycles_active并不是2个fp64 core的metrics

tee fp64_test.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>__global__ void kernel_add_double(volatile double *input,volatile double *output)
{unsigned int tid  = threadIdx.x + blockIdx.x * blockDim.x;double left=input[tid];double right=output[tid];for(int i=0;i<1;i++){left+=right;}       output[tid]=left;
}
EOF/usr/local/cuda/bin/nvcc -std=c++17 -dc -lineinfo -arch=sm_86 -ptx fp64_test.cu -o fp64_test.ptx
/usr/local/cuda/bin/nvcc -arch=sm_86 fp64_test.ptx -cubin -o fp64_test.cubin
/usr/local/cuda/bin/nvcc -arch=sm_86 fp64_test.cubin -fatbin -o fp64_test.fatbin
cat fp64_test.ptx
/usr/local/cuda/bin/cuobjdump --dump-sass fp64_test.fatbincuasm.py fp64_test.cubin fp64_test.cuasm
sed '/MOV/d' -i fp64_test.cuasm
sed '/S2R/d' -i fp64_test.cuasm
sed '/ULDC/d' -i fp64_test.cuasm
sed '/IMAD/d' -i fp64_test.cuasm
sed '/LDG/d' -i fp64_test.cuasm
sed '/STG/d' -i fp64_test.cuasm
sed '/F2F/d' -i fp64_test.cuasmcuasm.py fp64_test.cuasm
/usr/local/cuda/bin/nvcc -arch=sm_86 fp64_test.cubin -fatbin -o fp64_test.fatbin
/usr/local/cuda/bin/cuobjdump --dump-sass fp64_test.fatbin
/usr/local/cuda/bin/cuobjdump --dump-resource-usage fp64_test.fatbintee fp64_test_main.cpp<<-'EOF'
#include <stdio.h>
#include <string.h>
#include <cuda_runtime.h>
#include <cuda.h>int main(int argc,char *argv[])
{CUresult error;CUdevice cuDevice;cuInit(0);int deviceCount = 0;error = cuDeviceGetCount(&deviceCount);error = cuDeviceGet(&cuDevice, 0);if(error!=CUDA_SUCCESS){printf("Error happened in get device!\n");}CUcontext cuContext;error = cuCtxCreate(&cuContext, 0, cuDevice);if(error!=CUDA_SUCCESS){printf("Error happened in create context!\n");}int block_count=1;int block_size=32*4*4;int thread_size=block_count*block_size;int data_size=sizeof(double)*thread_size;double *output_ptr=nullptr;double *input_ptr=nullptr;int cudaStatus=0;cudaStatus = cudaMalloc((void**)&input_ptr, data_size);cudaStatus = cudaMalloc((void**)&output_ptr, data_size);void *kernelParams[]= {(void*)&output_ptr, (void*)&input_ptr};CUmodule module;CUfunction double_function;const char* module_file = "fp64_test.fatbin";const char* double_kernel_name = "_Z17kernel_add_doublePVdS0_";error = cuModuleLoad(&module, module_file);if(error!=CUDA_SUCCESS){printf("Error happened in load moudle %d!\n",error);}error = cuModuleGetFunction(&double_function, module, double_kernel_name);if(error!=CUDA_SUCCESS){printf("get float_kernel_name error!\n");}    cuLaunchKernel(double_function,block_count, 1, 1,8, 1, 1,0,0,kernelParams, 0);cuLaunchKernel(double_function,block_count, 1, 1,16, 1, 1,0,0,kernelParams, 0);cuLaunchKernel(double_function,block_count, 1, 1,32, 1, 1,0,0,kernelParams, 0);cuLaunchKernel(double_function,block_count, 1, 1,32*2, 1, 1,0,0,kernelParams, 0);cuLaunchKernel(double_function,block_count, 1, 1,32*4, 1, 1,0,0,kernelParams, 0);cuLaunchKernel(double_function,block_count, 1, 1,32*5, 1, 1,0,0,kernelParams, 0);cuLaunchKernel(double_function,block_count, 1, 1,32*4*8, 1, 1,0,0,kernelParams, 0);cudaFree(output_ptr);cudaFree(input_ptr);cuModuleUnload(module);cuCtxDestroy(cuContext);return 0;
}
EOF
g++ fp64_test_main.cpp -o fp64_test_main -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcudart -lcuda/usr/local/NVIDIA-Nsight-Compute/ncu --metrics smsp__pipe_fp64_cycles_active ./fp64_test_main

输出

kernel_add_double(volatile double *, volatile double *) (1, 1, 1)x(8, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
--------------------------------- ----------- ------------
Metric Name                       Metric Unit Metric Value
--------------------------------- ----------- ------------
smsp__pipe_fp64_cycles_active.avg       cycle         0.14
smsp__pipe_fp64_cycles_active.max       cycle           16
smsp__pipe_fp64_cycles_active.min       cycle            0
smsp__pipe_fp64_cycles_active.sum       cycle           16
--------------------------------- ----------- ------------kernel_add_double(volatile double *, volatile double *) (1, 1, 1)x(16, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
--------------------------------- ----------- ------------
Metric Name                       Metric Unit Metric Value
--------------------------------- ----------- ------------
smsp__pipe_fp64_cycles_active.avg       cycle         0.14
smsp__pipe_fp64_cycles_active.max       cycle           16  #不足一个warp跟一个warp 的pipe_fp64_cycles_active一样,说明存在无效计算
smsp__pipe_fp64_cycles_active.min       cycle            0
smsp__pipe_fp64_cycles_active.sum       cycle           16
--------------------------------- ----------- ------------kernel_add_double(volatile double *, volatile double *) (1, 1, 1)x(32, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
--------------------------------- ----------- ------------
Metric Name                       Metric Unit Metric Value
--------------------------------- ----------- ------------
smsp__pipe_fp64_cycles_active.avg       cycle         0.14
smsp__pipe_fp64_cycles_active.max       cycle           16
smsp__pipe_fp64_cycles_active.min       cycle            0
smsp__pipe_fp64_cycles_active.sum       cycle           16
--------------------------------- ----------- ------------kernel_add_double(volatile double *, volatile double *) (1, 1, 1)x(64, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
--------------------------------- ----------- ------------
Metric Name                       Metric Unit Metric Value
--------------------------------- ----------- ------------
smsp__pipe_fp64_cycles_active.avg       cycle         0.29
smsp__pipe_fp64_cycles_active.max       cycle           16
smsp__pipe_fp64_cycles_active.min       cycle            0
smsp__pipe_fp64_cycles_active.sum       cycle           32
--------------------------------- ----------- ------------kernel_add_double(volatile double *, volatile double *) (1, 1, 1)x(128, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
--------------------------------- ----------- ------------
Metric Name                       Metric Unit Metric Value
--------------------------------- ----------- ------------
smsp__pipe_fp64_cycles_active.avg       cycle         0.57
smsp__pipe_fp64_cycles_active.max       cycle           16
smsp__pipe_fp64_cycles_active.min       cycle            0
smsp__pipe_fp64_cycles_active.sum       cycle           64
--------------------------------- ----------- ------------kernel_add_double(volatile double *, volatile double *) (1, 1, 1)x(160, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
--------------------------------- ----------- ------------
Metric Name                       Metric Unit Metric Value
--------------------------------- ----------- ------------
smsp__pipe_fp64_cycles_active.avg       cycle         0.71
smsp__pipe_fp64_cycles_active.max       cycle           32
smsp__pipe_fp64_cycles_active.min       cycle            0
smsp__pipe_fp64_cycles_active.sum       cycle           80
--------------------------------- ----------- ------------kernel_add_double(volatile double *, volatile double *) (1, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
Section: Command line profiler metrics
--------------------------------- ----------- ------------
Metric Name                       Metric Unit Metric Value
--------------------------------- ----------- ------------
smsp__pipe_fp64_cycles_active.avg       cycle         4.57
smsp__pipe_fp64_cycles_active.max       cycle          128
smsp__pipe_fp64_cycles_active.min       cycle            0
smsp__pipe_fp64_cycles_active.sum       cycle          512 #每个smsp 执行一个warp的fp64需要16个pipe_fp64_cycles_active
--------------------------------- ----------- ------------

• 如果是2个fp64 cores的metrics,不会出现这样的现象