你回复的好快!以及,不用叫您,叫我为你即可。
自从指标变成了nsight风格的这种诡异名字后,很多东西从字面上就看不懂了(我,不包含你可能)。我感觉我们这里需要理解(或者你已经知道,请直接告诉我)至少这些点的东西, 首先是lts__的前缀,这点从NSight手册中,我们知道这个是L2 Slice(多个小L2 cache组成一个整体的大L2)的指标,那么看起来这里的hit和miss是和L2 cache相关的,似乎没错。这是第一点。
然后是aperture_sysmem,这个NSight手册里也有,是interface to system memory, 似乎也比较正常。
然后是sector,这是32B, 无问题,你之前的多次实验也验证了这点了。
那么对于指标:lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit (请一定要原谅nsight的奇葩命名方式),似乎就只剩下srcnode_gpc, lookup_hit了。
这里就需要猜测了。考虑到请求是从GPC发起的(或者说从SM内部),似乎整体可以读作:L2方面,sector计数,从GPC发起的system memory请求,的lookup_hit, 了。其中不知道的就只有最后一项了。
如果说这里的lookup_hit, 等于L2中的cache line hit,似乎说不过去。。我不能知道这里的具体含义了,你帮我继续猜测一下?
不过我建议你继续做一个这样的实验:
分配256MB的显存,连续启动2个kernel,例如K1和K2,其中K1将这256MB填充成0或者threadIdx.x之类的数据,然后K2继续使用你的原本的output[threadIdx.x] = 123;类似这种写入(注意右侧无读取,为了最小化的控制影响因素)。然后使用ncu -k K2的方式进行分析,然后看看结果情况?
我预测会都是lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss之类的。
欢迎实验来证明或者证否,你的测试非常有意思。我非常喜欢你的发帖!
我修改了程序:
__global__ void LtsTSectorsSrcnodeGpcApertureSysmemHitMissThreadNum32Kernel1(float *input, int len) {
int index = threadIdx.x + blockIdx.x * blockDim.x;
if (index < len) {
input[threadIdx.x] = threadIdx.x;
}
}
__global__ void LtsTSectorsSrcnodeGpcApertureSysmemHitMissThreadNum32Kernel2(float *output, int len) {
int index = threadIdx.x + blockIdx.x * blockDim.x;
if (index < len) {
output[threadIdx.x] = 123;
}
}
void LtsTSectorsSrcnodeGpcApertureSysmemLookupHitMiss() {
int thread_num = 512;
int len = 64 * 1024 * 1024;
float *d_src = nullptr;
float *d_res = nullptr;
dim3 block(thread_num);
dim3 grid((len + thread_num) / thread_num);
cudaHostAlloc(&d_src, 64 * 1024 * 1024 * sizeof(float), cudaHostAllocDefault);
cudaHostAlloc(&d_res, 64 * 1024 * 1024 * sizeof(float), cudaHostAllocDefault);
CUDA_CHECK(cudaSetDevice(0));
// kernel
LtsTSectorsSrcnodeGpcApertureSysmemHitMissThreadNum32Kernel1<<<grid, block>>>(d_src, len);
LtsTSectorsSrcnodeGpcApertureSysmemHitMissThreadNum32Kernel2<<<grid, block>>>(d_src, len);
CUDA_CHECK(cudaStreamSynchronize(0));
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaFreeHost(d_src));
CUDA_CHECK(cudaFreeHost(d_res));
}
然后测试的结果为:
sudo ncu --metrics lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit,lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss bin/test_lts_sector_GPC
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit_miss test
==PROF== Connected to process 7629 (/home/yongjian/CUDA_Note/build/bin/test_lts_sector_GPC)
==PROF== Profiling "LtsTSectorsSrcnodeGpcAperture..." - 1: 0%....50%....100% - 1 pass
==PROF== Profiling "LtsTSectorsSrcnodeGpcAperture..." - 2: 0%....50%....100% - 1 pass
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit_miss test finished
==PROF== Disconnected from process 7629
[7629] test_lts_sector_GPC@127.0.0.1
LtsTSectorsSrcnodeGpcApertureSysmemHitMissThreadNum32Kernel1(float *, int), 2022-Jan-04 21:11:37, Context 1, Stream 7
Section: Command line profiler metrics
---------------------------------------------------------------------- --------------- ------------------------------
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit.avg sector 198.631,60
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit.max sector 2.057.156
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit.min sector 0
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit.sum sector
7.945.264 lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss.avg sector 0
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss.max sector 0
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss.min sector 0
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss.sum sector 0
---------------------------------------------------------------------- --------------- ------------------------------
LtsTSectorsSrcnodeGpcApertureSysmemHitMissThreadNum32Kernel2(float *, int), 2022-Jan-04 21:11:37, Context 1, Stream 7
Section: Command line profiler metrics
---------------------------------------------------------------------- --------------- ------------------------------
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit.avg sector 198.242,90
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit.max sector 2.061.616
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit.min sector 0
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit.sum sector
7.929.716 lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss.avg sector 0
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss.max sector 0
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss.min sector 0
lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss.sum sector 0
---------------------------------------------------------------------- --------------- ------------------------------
结果显示miss是0,应该是因为没有load数据 所以没有miss,而hit则是很大的数,而且两个kernel的值还不一样,而且每次测试的结果也不一样,是因为数据量太大导致的计数误差?
当我缩小数据量为1024*4byte时,两个 kernel的hit固定到了128。
我其实也不太理解这个hit的含义,我的猜测是write miss的时候会采取write allocate策略,先把数据load到L2 cache上,然后write hit L2 cache上的数据,但是这样的话为什么miss的sector没有增加呢?