最新帖子

页: 1 2 [3] 4 5 ... 10
21
CUDA / Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 最后发表 作者 屠戮人神 一月 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的变化。你觉得如何?
22
CUDA / Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 最后发表 作者 LibAndLab 一月 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,目前还没有人回复,
23
CUDA / Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 最后发表 作者 屠戮人神 一月 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中的最终被更改的那些内容,而越过其他没有改变的,此时没有读取放大。(这点往往和系统内存不同)。

这就是我所知道的一切,更多内容欢迎你测试并发送到论坛和大家分享。
24
CUDA / Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 最后发表 作者 LibAndLab 十二月 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数据嘛?还是说直接就写了
25
CUDA / Re: Linux系统:cuFFT使用错误 undefined reference to XXX
« 最后发表 作者 屠戮人神 十二月 29, 2021, 07:12:01 pm »
linux系统在nsight edition中使用cufft库函数,已经#include <cufft.h>,还是出现以下错误:
     undefined reference to `cufftPlan1d'
查阅后发现windows下解决方法如下:
    项目——属性——平台选x64——链接器——输入——附加依赖项 在这个里面加入cufft.lib这个名字。
但是linux中的IDE没有这个选项,请教一下要怎么改正?

你需要链接cufft的库文件,但是我不是Eclipse(nsight edition)的用户,暂时不知道如何设定这个,你要不要询问一下你周围的使用eclipse的人?
26
CUDA / Re: ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 最后发表 作者 屠戮人神 十二月 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都是不重复的,也许可能会好一点?
27
CUDA / ncu测试lts__t_sectors_srcnode_gpc_aperture_sysmem发现的问题
« 最后发表 作者 LibAndLab 十二月 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,不知道这是为什么?


28
CUDA / Linux系统:cuFFT使用错误 undefined reference to XXX
« 最后发表 作者 jinyer 十二月 23, 2021, 10:23:26 am »
linux系统在nsight edition中使用cufft库函数,已经#include <cufft.h>,还是出现以下错误:
     undefined reference to `cufftPlan1d'
查阅后发现windows下解决方法如下:
    项目——属性——平台选x64——链接器——输入——附加依赖项 在这个里面加入cufft.lib这个名字。
但是linux中的IDE没有这个选项,请教一下要怎么改正?
29
CUDA / Re: CUDA多个核函数kernel之间的数据如何传递
« 最后发表 作者 屠戮人神 十二月 15, 2021, 12:39:46 pm »
假设程序中有两个核函数:kernel1和kernel2

        kernel2中想要用kernel1 的结果,可kernel2使用的时候值为空,即一个kernel运行的结果没有传递给另外一个kernel使用,那么多个kernel之间的数据应该怎么做才能传递使用??

         下边程序中开辟的内存也没有释放,为什么kernel1:initV 时对u0、u1成功赋值之后kernel2:convPy 使用时值却为Nan呢?
程序代码: [选择]

// 先在device开辟内存
cudaMalloc((void**)&u0, mea * wr * sizeof(float));
cudaMalloc((void**)&u1, mea * wr * sizeof(float));

// 通过initV对u0、u1 赋值
initV <<<grid1, block1 >>> (u0, u1, k);

// 同步,确保上边u0、u1成功初始化
cudaDeviceSynchronize();

// convPy函数需要用到u0、u1的值来计算temp
convPy <<<grid22, block22 >>> (temp, u0, u1);

// 最后才释放内存
cudaFree(u0);
cudaFree(u1);


你这是2个独立的问题:
(1)2个kernel调用之间的缓冲区(指针)的传输问题。
(2)为何第二个kernel没有取得预想中的正确结果。

关于第一个问题,如你所写,直接将第一个kernel的结果缓冲区(指针)输出,作为第二个kernel的输入参数即可,这没什么特别的,和普通的C函数调用一个道理。注意你可能并不需要2个kernel调用之间的同步,当然写了也不错。

关于第二个问题,为何某kernel没有运行得到成功结果。这个其实和第一个问题无关,只要你给的参数是正确的,同时你确定(刚才你说的)了第一个kernel的结果正常,那么只可能你第二个kernel写的有问题,跪了。这个得你好好检查一下(而不是这种连续调用2个kernel不行。连续调用2个kernel是常见的实践,逻辑上无问题的)。

30
CUDA / CUDA多个核函数kernel之间的数据如何传递
« 最后发表 作者 jinyer 十二月 15, 2021, 10:22:52 am »
假设程序中有两个核函数:kernel1和kernel2

        kernel2中想要用kernel1 的结果,可kernel2使用的时候值为空,即一个kernel运行的结果没有传递给另外一个kernel使用,那么多个kernel之间的数据应该怎么做才能传递使用??

         下边程序中开辟的内存也没有释放,为什么kernel1:initV 时对u0、u1成功赋值之后kernel2:convPy 使用时值却为Nan呢?
程序代码: [选择]

// 先在device开辟内存
cudaMalloc((void**)&u0, mea * wr * sizeof(float));
cudaMalloc((void**)&u1, mea * wr * sizeof(float));

// 通过initV对u0、u1 赋值
initV <<<grid1, block1 >>> (u0, u1, k);

// 同步,确保上边u0、u1成功初始化
cudaDeviceSynchronize();

// convPy函数需要用到u0、u1的值来计算temp
convPy <<<grid22, block22 >>> (temp, u0, u1);

// 最后才释放内存
cudaFree(u0);
cudaFree(u1);

页: 1 2 [3] 4 5 ... 10