ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题

  • 15 replies
  • 1826 views
ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 于: 十二月 23, 2021, 12:16:56 pm »
这个metrics的含义是LTS sectors from node GPC accessing system memory (sysmem),根据我目前查找到的资料和我的理解,这个metrics是统计访问system memory时触发的L2 cache sector的数量,其中system memory指的是cpu端的内存,我的测试代码比较简单:
```
   int thread_num = 512;
    int grid_num = 1;
    float *d_src = nullptr;
    float *d_res = nullptr;

    dim3 block(thread_num);
    dim3 grid(grid_num);

    cudaHostAlloc(&d_src, grid_num * thread_num * sizeof(float), cudaHostAllocDefault);
    cudaHostAlloc(&d_res, grid_num * thread_num * sizeof(float), cudaHostAllocDefault);

    CUDA_CHECK(cudaSetDevice(0));
    // kernel
    LtsTSectorsSrcnodeGpcApertureSysmemThreadNum32Kernel<<<grid, block>>>(d_src, d_res);
    CUDA_CHECK(cudaStreamSynchronize(0));
    CUDA_CHECK(cudaGetLastError());

    CUDA_CHECK(cudaFreeHost(d_src));
    CUDA_CHECK(cudaFreeHost(d_res));
```

```
__global__ void LtsTSectorsSrcnodeGpcApertureSysmemThreadNum32Kernel(float *input, float *output) {
    output[threadIdx.x] = input[threadIdx.x];
}
```
当thread_num =512,grid_num = 1时候,测量的结果是lts__t_sectors_srcnode_gpc_aperture_sysmem.sum  sector  128;
我的理解是128是这么得到的:512个线程*4个字节/32字节=64 ,对于load和store都会统计,所以需要64*2=128;
后面陆续增加thread_num和grid_num,发现metrics会一直按照按照上述计算方式线性增加,比如
thread_num =1024,grid_num = 60时,结果为1024*4*60*2/32=15360,和ncu测试的结果一致,但是当继续增加grid_num ,发现如下结果:
thread_num =1024,grid_num = 70时,按照推算的结果是:17920,但是ncu的结果是lts__t_sectors_srcnode_gpc_aperture_sysmem.sum sector 17656,不知道这是为什么?



Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #1 于: 十二月 29, 2021, 06:53:28 pm »
这个metrics的含义是LTS sectors from node GPC accessing system memory (sysmem),根据我目前查找到的资料和我的理解,这个metrics是统计访问system memory时触发的L2 cache sector的数量,其中system memory指的是cpu端的内存,我的测试代码比较简单:
```
   int thread_num = 512;
    int grid_num = 1;
    float *d_src = nullptr;
    float *d_res = nullptr;

    dim3 block(thread_num);
    dim3 grid(grid_num);

    cudaHostAlloc(&d_src, grid_num * thread_num * sizeof(float), cudaHostAllocDefault);
    cudaHostAlloc(&d_res, grid_num * thread_num * sizeof(float), cudaHostAllocDefault);

    CUDA_CHECK(cudaSetDevice(0));
    // kernel
    LtsTSectorsSrcnodeGpcApertureSysmemThreadNum32Kernel<<<grid, block>>>(d_src, d_res);
    CUDA_CHECK(cudaStreamSynchronize(0));
    CUDA_CHECK(cudaGetLastError());

    CUDA_CHECK(cudaFreeHost(d_src));
    CUDA_CHECK(cudaFreeHost(d_res));
```

```
__global__ void LtsTSectorsSrcnodeGpcApertureSysmemThreadNum32Kernel(float *input, float *output) {
    output[threadIdx.x] = input[threadIdx.x];
}
```
当thread_num =512,grid_num = 1时候,测量的结果是lts__t_sectors_srcnode_gpc_aperture_sysmem.sum  sector  128;
我的理解是128是这么得到的:512个线程*4个字节/32字节=64 ,对于load和store都会统计,所以需要64*2=128;
后面陆续增加thread_num和grid_num,发现metrics会一直按照按照上述计算方式线性增加,比如
thread_num =1024,grid_num = 60时,结果为1024*4*60*2/32=15360,和ncu测试的结果一致,但是当继续增加grid_num ,发现如下结果:
thread_num =1024,grid_num = 70时,按照推算的结果是:17920,但是ncu的结果是lts__t_sectors_srcnode_gpc_aperture_sysmem.sum sector 17656,不知道这是为什么?


我不清楚profiler报告的该项指标的具体含义,不过你的大胆假设-小心验证的流程很不错!

我们是否可以稍微修改一下kernel的写法,看看这样会如何:
int idx = blockDim.x * blockIdx.x + threadIdx.x;
output[idx] = input[idx];
其中output和input都是来自于cudaMallocHost/cudaHostAlloc的指针,你看这样好不好?

因为我担心该指标会受到某些cache效果的影响?

这样我们将idx都是不重复的,也许可能会好一点?

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #2 于: 十二月 31, 2021, 10:45:17 am »
我不清楚profiler报告的该项指标的具体含义,不过你的大胆假设-小心验证的流程很不错!

我们是否可以稍微修改一下kernel的写法,看看这样会如何:
int idx = blockDim.x * blockIdx.x + threadIdx.x;
output[idx] = input[idx];
其中output和input都是来自于cudaMallocHost/cudaHostAlloc的指针,你看这样好不好?

因为我担心该指标会受到某些cache效果的影响?

这样我们将idx都是不重复的,也许可能会好一点?
的确是收到了cache的影响,另外我想问下ampere架构的L2 cache在遇到write miss的时候会怎么操作?比如说:
__global__ void Kernel(float *output) {
    output[threadIdx.x] = threadIdx.x;
}
核函数直接对output进行赋值操作,L2 cache write miss,这是L2 cache会load数据嘛?还是说直接就写了

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #3 于: 一月 02, 2022, 08:45:05 pm »
的确是收到了cache的影响,另外我想问下ampere架构的L2 cache在遇到write miss的时候会怎么操作?比如说:
__global__ void Kernel(float *output) {
    output[threadIdx.x] = threadIdx.x;
}
核函数直接对output进行赋值操作,L2 cache write miss,这是L2 cache会load数据嘛?还是说直接就写了

我想你询问的是,L2(注意它同时服务于显存,和系统内存)会先读取一段一定粒度的内容,然后改写掉其中写入的部分吗?

从你的代码看,对于float类型的写入,使用类似output[threadIdx.x] = xxx;的风格,很可能不会,因为每个warp总是写入128B,而且如果该output来自于cudaMalloc的话,自带对齐到至少256B的边界的属性。所以这连续写入的128B往往可能覆盖了一个或者多个最小cache line(或者其他的)写入单位。此时可能不会触发任何写入导致的读取放大效应。

不过一般的讨论的话,如果只是随意的往某段地址写入,不能合并成最小的一个或者几个写入单位,则不好说,我们这里不讨论L2服务系统内存的情况(例如分配的内存当作global memory,因为我不知道这种情况是什么样子),则L2下面只是显存的话,得看显存的类型。

对于例如有ECC的显卡来说,写入一段区域内的最小某个单位中的部分内容,往往会触发读取,因为此时要重新ECC计算,重新写入。

对于例如普通显卡来说,不需要有ECC的情况下,因为写入可以带有mask,只标记了要写入了显存DRAM中的最终被更改的那些内容,而越过其他没有改变的,此时没有读取放大。(这点往往和系统内存不同)。

这就是我所知道的一切,更多内容欢迎你测试并发送到论坛和大家分享。

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #4 于: 一月 04, 2022, 02:40:31 pm »
我想你询问的是,L2(注意它同时服务于显存,和系统内存)会先读取一段一定粒度的内容,然后改写掉其中写入的部分吗?

从你的代码看,对于float类型的写入,使用类似output[threadIdx.x] = xxx;的风格,很可能不会,因为每个warp总是写入128B,而且如果该output来自于cudaMalloc的话,自带对齐到至少256B的边界的属性。所以这连续写入的128B往往可能覆盖了一个或者多个最小cache line(或者其他的)写入单位。此时可能不会触发任何写入导致的读取放大效应。

不过一般的讨论的话,如果只是随意的往某段地址写入,不能合并成最小的一个或者几个写入单位,则不好说,我们这里不讨论L2服务系统内存的情况(例如分配的内存当作global memory,因为我不知道这种情况是什么样子),则L2下面只是显存的话,得看显存的类型。

对于例如有ECC的显卡来说,写入一段区域内的最小某个单位中的部分内容,往往会触发读取,因为此时要重新ECC计算,重新写入。

对于例如普通显卡来说,不需要有ECC的情况下,因为写入可以带有mask,只标记了要写入了显存DRAM中的最终被更改的那些内容,而越过其他没有改变的,此时没有读取放大。(这点往往和系统内存不同)。

这就是我所知道的一切,更多内容欢迎你测试并发送到论坛和大家分享。
L2 cache可以处理自身global memory、peer memory和system memory的访存事务,我在 3080上测试了这两个lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hitlts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss,发现核函数为:
input[threadIdx.x] = input[threadIdx.x]*1.5f;
和核函数为
output[threadIdx.x] = input[threadIdx.x]*1.5f;
时metrics的值是相同的,比如:针对以上的两个核函数,当启动的线程数从32线程变为64线程时,对应的lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss都会增加4个sector,而lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit也都会增加4个sector;
对于lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss这个metrics我感觉比较好理解,因为线程数从32变成64时,会多load 128字节的数据,所以会miss 4个sector;
但是对于lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit我就不太理解了,对于第一个核函数,我认为是write hit,所以这个metrics增加4个sector是能讲得通的,但是对于第二个kernel,应该是对应的write miss,为什么hit还会增加4个sector呢?
global memory和system memory的变化是一样的,都是会增加4个miss和4个hit
我在官方论坛也发了帖子https://forums.developer.nvidia.com/t/ampere-gpu-l2-cache-write-miss-policy/199032,目前还没有人回复,
« 最后编辑时间: 一月 04, 2022, 02:43:39 pm 作者 LibAndLab »

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #5 于: 一月 04, 2022, 06:03:28 pm »
L2 cache可以处理自身global memory、peer memory和system memory的访存事务,我在 3080上测试了这两个lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hitlts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss,发现核函数为:
input[threadIdx.x] = input[threadIdx.x]*1.5f;
和核函数为
output[threadIdx.x] = input[threadIdx.x]*1.5f;
时metrics的值是相同的,比如:针对以上的两个核函数,当启动的线程数从32线程变为64线程时,对应的lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss都会增加4个sector,而lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit也都会增加4个sector;
对于lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_miss这个metrics我感觉比较好理解,因为线程数从32变成64时,会多load 128字节的数据,所以会miss 4个sector;
但是对于lts__t_sectors_srcnode_gpc_aperture_sysmem_lookup_hit我就不太理解了,对于第一个核函数,我认为是write hit,所以这个metrics增加4个sector是能讲得通的,但是对于第二个kernel,应该是对应的write miss,为什么hit还会增加4个sector呢?
global memory和system memory的变化是一样的,都是会增加4个miss和4个hit
我在官方论坛也发了帖子https://forums.developer.nvidia.com/t/ampere-gpu-l2-cache-write-miss-policy/199032,目前还没有人回复,

根据你的测试,我怀疑hit再写入的时候,有特别的含义:完整的覆盖了一个特定大小的最小粒度,中间没有遗漏的原始值。

我建议你将你的测试进行如下修改看看能否验证,或者否定对你的测试实验的猜测:
output[2 * threadIdx.x] = input[threadIdx.x]*1.5f;
(注意output需要分配双倍的空间)

然后重新观察hit和miss的变化。你觉得如何?

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #6 于: 一月 04, 2022, 06:55:51 pm »
根据你的测试,我怀疑hit再写入的时候,有特别的含义:完整的覆盖了一个特定大小的最小粒度,中间没有遗漏的原始值。

我建议你将你的测试进行如下修改看看能否验证,或者否定对你的测试实验的猜测:
output[2 * threadIdx.x] = input[threadIdx.x]*1.5f;
(注意output需要分配双倍的空间)

然后重新观察hit和miss的变化。你觉得如何?

按照您的修改程序后,启动32个线程,此时load的范围是4个sector,write的范围扩大到了8个sector,测试的结果是: miss 4个sector,hit 8个sector,那这里的hit是不是不能理解成cache hit的概念?

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #7 于: 一月 04, 2022, 07:44:51 pm »
按照您的修改程序后,启动32个线程,此时load的范围是4个sector,write的范围扩大到了8个sector,测试的结果是: miss 4个sector,hit 8个sector,那这里的hit是不是不能理解成cache hit的概念?

你回复的好快!以及,不用叫您,叫我为你即可。

自从指标变成了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之类的。

欢迎实验来证明或者证否,你的测试非常有意思。我非常喜欢你的发帖!


Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #8 于: 一月 04, 2022, 09:18:29 pm »
你回复的好快!以及,不用叫您,叫我为你即可。

自从指标变成了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没有增加呢?

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #9 于: 一月 04, 2022, 11:41:10 pm »
我修改了程序:
__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没有增加呢?

我看到了你的新代码。的确不是我设想的样子。我原本的设想是通过两次写入,这样第二次的kernel写入肯定不可能在cache中的,特别大型的这种大小的缓冲区(例如你的256MB的这个),会有miss计数(假设这是write miss,就是要写入/改写没有在cache中的line/block)。但是的确没有。只有hit计数。

所以这个hit的定义就奇怪了。这是第一点(依然还需要你想,我暂时无法知道它的含义)。


第二点,关于你的结论,fetch on write, 似乎不应该是这样的。否则要写入一个肯定不在cache中的line,如果要load一次的话,那么应该有miss计数的,也没有。

上个帖子的output[2 * idx]的写入,就是想看看这种隔一个元素改写掉一个元素,会不会有load(因为理论上说,这种必须要有load)。结果miss之类的指标都是0.

或者说,只能read的时候有miss,write的时候永远没有。但是无法知道如何能实现隔离1个元素改写一个元素。猜测可能拆分成多次跨越PCI-E的传输了,每次只传输4B,然后间隔一个4B,再4B,如此反复(我不懂PCI-E, 尚不知道这样是否可行)。如果是这样的话,那L2(再对system memory)的写入的时候,就不是fetch on write(allocate)的了。而且用掩码标记了哪些位置被改写了,这个可能接近点目标。从而永远写入没有miss。你看对吗这样想?

以及,上面都没有想页表和页表缓存(TLB)的因素。

更多的,请LibAndLab测试,并教晓我。



Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #10 于: 一月 05, 2022, 10:01:55 am »
我看到了你的新代码。的确不是我设想的样子。我原本的设想是通过两次写入,这样第二次的kernel写入肯定不可能在cache中的,特别大型的这种大小的缓冲区(例如你的256MB的这个),会有miss计数(假设这是write miss,就是要写入/改写没有在cache中的line/block)。但是的确没有。只有hit计数。

所以这个hit的定义就奇怪了。这是第一点(依然还需要你想,我暂时无法知道它的含义)。


第二点,关于你的结论,fetch on write, 似乎不应该是这样的。否则要写入一个肯定不在cache中的line,如果要load一次的话,那么应该有miss计数的,也没有。

上个帖子的output[2 * idx]的写入,就是想看看这种隔一个元素改写掉一个元素,会不会有load(因为理论上说,这种必须要有load)。结果miss之类的指标都是0.

或者说,只能read的时候有miss,write的时候永远没有。但是无法知道如何能实现隔离1个元素改写一个元素。猜测可能拆分成多次跨越PCI-E的传输了,每次只传输4B,然后间隔一个4B,再4B,如此反复(我不懂PCI-E, 尚不知道这样是否可行)。如果是这样的话,那L2(再对system memory)的写入的时候,就不是fetch on write(allocate)的了。而且用掩码标记了哪些位置被改写了,这个可能接近点目标。从而永远写入没有miss。你看对吗这样想?

以及,上面都没有想页表和页表缓存(TLB)的因素。

更多的,请LibAndLab测试,并教晓我。
对于显存 我也进行了相关的测试,也是write的时候全是当作hit 没有当作miss:四个kernel,内部实现是相同的,分别启动,启动的时候对应的显存也是不同的
kernel:
__global__ void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum32Kernel(float *input, float *output) {
    output[threadIdx.x] = threadIdx.x;
}

__global__ void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum64Kernel(float *input, float *output) {
    output[threadIdx.x] = threadIdx.x;
}

__global__ void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum96Kernel(float *input, float *output) {
    output[threadIdx.x] = threadIdx.x;
}

__global__ void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum128Kernel(float *input, float *output) {
    output[threadIdx.x] = threadIdx.x;
}

测试主程序:
void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMiss() {
    int thread_num[4]{32, 64, 96, 128};
    int grid_num = 1;
    float *d_src[4] = {nullptr};
    float *d_res[4] = {nullptr};

    for (int i = 0; i < 4; i++) {
        dim3 block(thread_num);
        dim3 grid(grid_num);

        CUDA_CHECK(cudaMalloc(&d_src, grid_num * thread_num * sizeof(float)));
        CUDA_CHECK(cudaMalloc(&d_res, grid_num * thread_num * sizeof(float)));
        CUDA_CHECK(cudaSetDevice(0));
        // kernel
        switch (i) {
            case 0: {
                LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum32Kernel<<<grid, block>>>(d_src, d_res);
                break;
            }

            case 1: {
                LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum64Kernel<<<grid, block>>>(d_src, d_res);
                break;
            }

            case 2: {
                LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum96Kernel<<<grid, block>>>(d_src, d_res);
                break;
            }

            case 3: {
                LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum128Kernel<<<grid, block>>>(d_src, d_res);
                break;
            }
            default:
                break;
        }

        CUDA_CHECK(cudaStreamSynchronize(0));
        CUDA_CHECK(cudaGetLastError());

        CUDA_FREE(d_src);
        CUDA_FREE(d_res);
    }
}

测试结果
sudo ncu --metrics lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit,lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss bin/test_lts_sector_GPC
lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit_miss test
==PROF== Connected to process 20693 (/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
==PROF== Profiling "LtsTSectorsSrcnodeGpcAperture..." - 3: 0%....50%....100% - 1 pass
==PROF== Profiling "LtsTSectorsSrcnodeGpcAperture..." - 4: 0%....50%....100% - 1 pass
==PROF== Disconnected from process 20693
[20693] test_lts_sector_GPC@127.0.0.1
  LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum32Kernel(float *, float *), 2022-Jan-05 09:57:41, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.avg                       sector                          12,40
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.max                       sector                             80
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.min                       sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.sum                       sector                            496
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.avg                      sector                           2,52
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.max                      sector                             16
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.min                      sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.sum                      sector                            101
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum64Kernel(float *, float *), 2022-Jan-05 09:57:41, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.avg                       sector                          12,50
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.max                       sector                             80
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.min                       sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.sum                       sector                            500
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.avg                      sector                           2,52
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.max                      sector                             16
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.min                      sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.sum                      sector                            101
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum96Kernel(float *, float *), 2022-Jan-05 09:57:41, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.avg                       sector                          12,62
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.max                       sector                             88
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.min                       sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.sum                       sector                            505
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.avg                      sector                           2,52
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.max                      sector                             21
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.min                      sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.sum                      sector                            101
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum128Kernel(float *, float *), 2022-Jan-05 09:57:41, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.avg                       sector                          12,72
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.max                       sector                             80
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.min                       sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.sum                       sector                            509
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.avg                      sector                           2,52
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.max                      sector                             16
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.min                      sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.sum                      sector                            101

结果分析
从结果来看,随着线程数的增加miss的数量并没有增加,反而hit的数量在依次增加4个sector

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #11 于: 一月 05, 2022, 12:34:11 pm »
对于显存 我也进行了相关的测试,也是write的时候全是当作hit 没有当作miss:四个kernel,内部实现是相同的,分别启动,启动的时候对应的显存也是不同的
kernel:
__global__ void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum32Kernel(float *input, float *output) {
    output[threadIdx.x] = threadIdx.x;
}

__global__ void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum64Kernel(float *input, float *output) {
    output[threadIdx.x] = threadIdx.x;
}

__global__ void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum96Kernel(float *input, float *output) {
    output[threadIdx.x] = threadIdx.x;
}

__global__ void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum128Kernel(float *input, float *output) {
    output[threadIdx.x] = threadIdx.x;
}

测试主程序:
void LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMiss() {
    int thread_num[4]{32, 64, 96, 128};
    int grid_num = 1;
    float *d_src[4] = {nullptr};
    float *d_res[4] = {nullptr};

    for (int i = 0; i < 4; i++) {
        dim3 block(thread_num);
        dim3 grid(grid_num);

        CUDA_CHECK(cudaMalloc(&d_src, grid_num * thread_num * sizeof(float)));
        CUDA_CHECK(cudaMalloc(&d_res, grid_num * thread_num * sizeof(float)));
        CUDA_CHECK(cudaSetDevice(0));
        // kernel
        switch (i) {
            case 0: {
                LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum32Kernel<<<grid, block>>>(d_src, d_res);
                break;
            }

            case 1: {
                LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum64Kernel<<<grid, block>>>(d_src, d_res);
                break;
            }

            case 2: {
                LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum96Kernel<<<grid, block>>>(d_src, d_res);
                break;
            }

            case 3: {
                LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum128Kernel<<<grid, block>>>(d_src, d_res);
                break;
            }
            default:
                break;
        }

        CUDA_CHECK(cudaStreamSynchronize(0));
        CUDA_CHECK(cudaGetLastError());

        CUDA_FREE(d_src);
        CUDA_FREE(d_res);
    }
}

测试结果
sudo ncu --metrics lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit,lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss bin/test_lts_sector_GPC
lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit_miss test
==PROF== Connected to process 20693 (/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
==PROF== Profiling "LtsTSectorsSrcnodeGpcAperture..." - 3: 0%....50%....100% - 1 pass
==PROF== Profiling "LtsTSectorsSrcnodeGpcAperture..." - 4: 0%....50%....100% - 1 pass
==PROF== Disconnected from process 20693
[20693] test_lts_sector_GPC@127.0.0.1
  LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum32Kernel(float *, float *), 2022-Jan-05 09:57:41, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.avg                       sector                          12,40
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.max                       sector                             80
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.min                       sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.sum                       sector                            496
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.avg                      sector                           2,52
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.max                      sector                             16
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.min                      sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.sum                      sector                            101
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum64Kernel(float *, float *), 2022-Jan-05 09:57:41, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.avg                       sector                          12,50
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.max                       sector                             80
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.min                       sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.sum                       sector                            500
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.avg                      sector                           2,52
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.max                      sector                             16
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.min                      sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.sum                      sector                            101
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum96Kernel(float *, float *), 2022-Jan-05 09:57:41, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.avg                       sector                          12,62
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.max                       sector                             88
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.min                       sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.sum                       sector                            505
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.avg                      sector                           2,52
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.max                      sector                             21
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.min                      sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.sum                      sector                            101
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTSectorsSrcnodeGpcApertureDeviceLookupHitMissThreadNum128Kernel(float *, float *), 2022-Jan-05 09:57:41, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.avg                       sector                          12,72
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.max                       sector                             80
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.min                       sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_hit.sum                       sector                            509
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.avg                      sector                           2,52
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.max                      sector                             16
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.min                      sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_lookup_miss.sum                      sector                            101

结果分析
从结果来看,随着线程数的增加miss的数量并没有增加,反而hit的数量在依次增加4个sector

你没有回答我的猜测。

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #12 于: 一月 05, 2022, 02:33:24 pm »
你没有回答我的猜测。
我觉得你的猜测是很有道理的,write miss policy除了 Write-allocate还有No-write-allocate:This is just what it sounds like! If you have a write miss in a no-write-allocate cache, you simply notify the next level down (similar to a write-through operation). You don't kick anything out of your own cache. 但是具体是如何写入的就不太清楚了,我觉得掩码标记可能性更大,真实的情况可能需要NVIIDA官方人员来解释下了

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #13 于: 一月 05, 2022, 08:07:34 pm »
我觉得你的猜测是很有道理的,write miss policy除了 Write-allocate还有No-write-allocate:This is just what it sounds like! If you have a write miss in a no-write-allocate cache, you simply notify the next level down (similar to a write-through operation). You don't kick anything out of your own cache. 但是具体是如何写入的就不太清楚了,我觉得掩码标记可能性更大,真实的情况可能需要NVIIDA官方人员来解释下了

所以对于显存(你的上上贴)和内存都是一样的了。。。差不多可以结帖了。很有意思的发现。你是一个很好的妹子。

Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 回复 #14 于: 一月 06, 2022, 10:36:50 am »
所以对于显存(你的上上贴)和内存都是一样的了。。。差不多可以结帖了。很有意思的发现。你是一个很好的妹子。
嗯嗯,可以结贴了。PS:我不是妹子 :)