验证4个SMSP是否是串行访问ShareMemory的

2024-08-27 23:12

本文主要是介绍验证4个SMSP是否是串行访问ShareMemory的,希望对大家解决编程问题提供一定的参考价值,需要的开发者们随着小编来一起学习吧!

验证4个SMSP是否是串行访问ShareMemory的

  • 测试过程

原以为4个smsp中的warp在没有bank冲突的情况下,是可以并行访问共享内存的
通过下面的测试发现,其实是串行的,share memory每个cycle只能处理一个请求

测试过程

tee shm_kernel.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
__global__ void shm_kernel(float *input,float *output) {int tid  = threadIdx.x + blockIdx.x * blockDim.x;__shared__ float shm_data[0xc000/4];float vals;clock_t t0=clock64();vals=shm_data[tid];__syncthreads();clock_t t1=clock64();vals*=(tid);output[tid]=vals;if(tid==0){printf("ts:%lld\n",t1-t0);}
}
EOF/usr/local/cuda/bin/nvcc -std=c++17  -dc -lineinfo -arch=sm_86 -ptx shm_kernel.cu -o shm_kernel.ptx
/usr/local/cuda/bin/nvcc -arch=sm_86 shm_kernel.ptx -cubin -o shm_kernel.cubin
/usr/local/cuda/bin/nvcc -arch=sm_86 shm_kernel.cubin -fatbin -o shm_kernel.fatbin
/usr/local/cuda/bin/cuobjdump --dump-sass shm_kernel.fatbintee shm_kernel_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");}CUmodule module;CUfunction function;const char* module_file = "shm_kernel.fatbin";const char* kernel_name = "_Z10shm_kernelPfS_";error = cuModuleLoad(&module, module_file);if(error!=CUDA_SUCCESS){printf("Error happened in load moudle %d!\n",error);}error = cuModuleGetFunction(&function, module, kernel_name);if(error!=CUDA_SUCCESS){printf("get function error!\n");}int thread_size_conf[3]={32,32*4,32*4*4};for(int k=0;k<3;k++){int block_size=1;int thread_size=thread_size_conf[k];int data_size=sizeof(float)*thread_size*block_size;float *output_ptr=nullptr;float *input_ptr=nullptr;int cudaStatus=0;cudaStatus = cudaMalloc((void**)&input_ptr, data_size);if(cudaStatus){printf("cudaMalloc input_ptr Failed\n");}cudaStatus= cudaMalloc((void**)&output_ptr, data_size);if(cudaStatus){printf("cudaMalloc output_ptr Failed\n");        }void *kernelParams[]= {(void*)&output_ptr, (void*)&input_ptr};auto ret=cuLaunchKernel(function,block_size, 1, 1,thread_size, 1, 1,0,0,kernelParams, 0);cudaError_t cudaerr = cudaDeviceSynchronize();if (cudaerr != cudaSuccess)printf("kernel launch failed with error \"%s\".\n",cudaGetErrorString(cudaerr));  cudaFree(output_ptr);cudaFree(input_ptr);        }cuModuleUnload(module);cuCtxDestroy(cuContext);return 0;
}
EOF
g++ shm_kernel_main.cpp -o shm_kernel_main -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcudart -lcuda/usr/local/NVIDIA-Nsight-Compute/ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,\
smsp__sass_l1tex_data_bank_conflicts_pipe_lsu_mem_shared_op_ldgsts.sum,\
smsp__sass_inst_executed_op_shared_ld.sum,\
l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum.peak_sustained,\
l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.avg.peak_sustained,\
l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum ./shm_kernel_main

输出

ts:33
ts:1551
0%....50%....100% - 3 passes
==PROF== Profiling "shm_kernel(float *, float *)" - 1: ts:39
ts:39
ts:1622
0%....50%....100% - 3 passes
==PROF== Profiling "shm_kernel(float *, float *)" - 2: ts:64
ts:57
ts:1706
0%....50%....100% - 3 passes
==PROF== Disconnected from process 657443
[657443] shm_kernel_main@127.0.0.1shm_kernel(float *, float *) (1, 1, 1)x(32, 1, 1), Context 1, Stream 7, Device 0, CC 8.6Section: Command line profiler metrics---------------------------------------------------------------------- ----------- ------------Metric Name                                                            Metric Unit Metric Value---------------------------------------------------------------------- ----------- ------------l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                                      0l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.avg.peak_sustained        1/cycle            1l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                          1l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum.peak_sustained        1/cycle           28smsp__sass_inst_executed_op_shared_ld.sum                                     inst            1smsp__sass_l1tex_data_bank_conflicts_pipe_lsu_mem_shared_op_ldgsts.sum                        0---------------------------------------------------------------------- ----------- ------------shm_kernel(float *, float *) (1, 1, 1)x(128, 1, 1), Context 1, Stream 7, Device 0, CC 8.6Section: Command line profiler metrics---------------------------------------------------------------------- ----------- ------------Metric Name                                                            Metric Unit Metric Value---------------------------------------------------------------------- ----------- ------------l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                                      0l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.avg.peak_sustained        1/cycle            1l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                          4l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum.peak_sustained        1/cycle           28smsp__sass_inst_executed_op_shared_ld.sum                                     inst            4smsp__sass_l1tex_data_bank_conflicts_pipe_lsu_mem_shared_op_ldgsts.sum                        0---------------------------------------------------------------------- ----------- ------------shm_kernel(float *, float *) (1, 1, 1)x(512, 1, 1), Context 1, Stream 7, Device 0, CC 8.6Section: Command line profiler metrics---------------------------------------------------------------------- ----------- ------------Metric Name                                                            Metric Unit Metric Value---------------------------------------------------------------------- ----------- ------------l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                                      0l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.avg.peak_sustained        1/cycle            1l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                         16l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum.peak_sustained        1/cycle           28smsp__sass_inst_executed_op_shared_ld.sum                                     inst           16smsp__sass_l1tex_data_bank_conflicts_pipe_lsu_mem_shared_op_ldgsts.sum                        0---------------------------------------------------------------------- ----------- ------------

这篇关于验证4个SMSP是否是串行访问ShareMemory的的文章就介绍到这儿,希望我们推荐的文章对编程师们有所帮助!



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

相关文章

Java通过驱动包(jar包)连接MySQL数据库的步骤总结及验证方式

《Java通过驱动包(jar包)连接MySQL数据库的步骤总结及验证方式》本文详细介绍如何使用Java通过JDBC连接MySQL数据库,包括下载驱动、配置Eclipse环境、检测数据库连接等关键步骤,... 目录一、下载驱动包二、放jar包三、检测数据库连接JavaJava 如何使用 JDBC 连接 mys

Spring Security中用户名和密码的验证完整流程

《SpringSecurity中用户名和密码的验证完整流程》本文给大家介绍SpringSecurity中用户名和密码的验证完整流程,本文结合实例代码给大家介绍的非常详细,对大家的学习或工作具有一定... 首先创建了一个UsernamePasswordAuthenticationTChina编程oken对象,这是S

MySQL中的InnoDB单表访问过程

《MySQL中的InnoDB单表访问过程》:本文主要介绍MySQL中的InnoDB单表访问过程,具有很好的参考价值,希望对大家有所帮助,如有错误或未考虑完全的地方,望不吝赐教... 目录1、背景2、环境3、访问类型【1】const【2】ref【3】ref_or_null【4】range【5】index【6】

前端如何通过nginx访问本地端口

《前端如何通过nginx访问本地端口》:本文主要介绍前端如何通过nginx访问本地端口的问题,具有很好的参考价值,希望对大家有所帮助,如有错误或未考虑完全的地方,望不吝赐教... 目录一、nginx安装1、下载(1)下载地址(2)系统选择(3)版本选择2、安装部署(1)解压(2)配置文件修改(3)启动(4)

如何搭建并配置HTTPD文件服务及访问权限控制

《如何搭建并配置HTTPD文件服务及访问权限控制》:本文主要介绍如何搭建并配置HTTPD文件服务及访问权限控制的问题,具有很好的参考价值,希望对大家有所帮助,如有错误或未考虑完全的地方,望不吝赐教... 目录一、安装HTTPD服务二、HTTPD服务目录结构三、配置修改四、服务启动五、基于用户访问权限控制六、

NGINX 配置内网访问的实现步骤

《NGINX配置内网访问的实现步骤》本文主要介绍了NGINX配置内网访问的实现步骤,Nginx的geo模块限制域名访问权限,仅允许内网/办公室IP访问,具有一定的参考价值,感兴趣的可以了解一下... 目录需求1. geo 模块配置2. 访问控制判断3. 错误页面配置4. 一个完整的配置参考文档需求我们有一

C#实现访问远程硬盘的图文教程

《C#实现访问远程硬盘的图文教程》在现实场景中,我们经常用到远程桌面功能,而在某些场景下,我们需要使用类似的远程硬盘功能,这样能非常方便地操作对方电脑磁盘的目录、以及传送文件,这次我们将给出一个完整的... 目录引言一. 远程硬盘功能展示二. 远程硬盘代码实现1. 底层业务通信实现2. UI 实现三. De

python通过curl实现访问deepseek的API

《python通过curl实现访问deepseek的API》这篇文章主要为大家详细介绍了python如何通过curl实现访问deepseek的API,文中的示例代码讲解详细,感兴趣的小伙伴可以跟随小编... API申请和充值下面是deepeek的API网站https://platform.deepsee

Nginx 访问 /root/下 403 Forbidden问题解决

《Nginx访问/root/下403Forbidden问题解决》在使用Nginx作为Web服务器时,可能会遇到403Forbidden错误,文中通过示例代码介绍的非常详细,对大家的学习或者工作... 目录解决 Nginx 访问 /root/test/1.html 403 Forbidden 问题问题复现Ng

Linux内核参数配置与验证详细指南

《Linux内核参数配置与验证详细指南》在Linux系统运维和性能优化中,内核参数(sysctl)的配置至关重要,本文主要来聊聊如何配置与验证这些Linux内核参数,希望对大家有一定的帮助... 目录1. 引言2. 内核参数的作用3. 如何设置内核参数3.1 临时设置(重启失效)3.2 永久设置(重启仍生效