请问,如何找到性能瓶颈?

我有一个程序,并行10个线程,读取硬盘上的数据,在cpu和gpu里依次做处理之后,然后将输出数据写入硬盘。
硬件设备:Intel i7-2600(4核8线程), 4G内存,nVidia GTX660Ti, 硬盘是威刚SSD AS510S3-120GM-C(读写500MB/s以上)。
Win7 64bit, 程序是按32位方式编译运行的。
用Windows任务管理器查看cpu负载,用GPU-Z这个软件查看GPU的负载。

运行结果:发现CPU和GPU的负载都没有跑满,cpu平均70-80%,gpu平均才30-40%。

检查了输入输出的速度,

  • 硬盘读取大约107M Bytes/s,低于SSD的读取指标
  • 硬盘写入大约9M Bytes/s,远低于SSD的写入指标
  • 从Host到Device的数据,大约107M Bytes/s,低于显卡指标
  • 从Device到Host的数据,大约214M Bytes/s,低于显卡指标

问题:为什么不能跑到cpu/gpu的100%的负载?这个程序的运行瓶颈在哪里呢?该如何找出呢?

(1)
如何在CPU上取得你购买的SSD的峰值速度,这个本论坛不进行技术支持,但建议您咨询您的管理员,甚至您的SSD生产商(为何9MB/S的写入)。

(2)关于PCI-E才100-200MB/S的传输速度,请问您是否每次传输的过小?
一次传输较大的内容,可以提高速度。
以及建议您使用page-locked memory以便取得最大的速度。

(3)为何您的CPU不能100%,抱歉这个本论坛无法提供技术支持。

(4)为何您的GPU load无法100%?
这个最常见的原因就是您的SM经常处于空闲状态,除了罕见的例如只上1个block或者其他block都退出就留一个在慢慢跑之外,最大的可能是您有无法掩盖的延迟。

感谢来访。

  1. 关于SSD的问题,谢谢提醒。我们重新实际测试SSD的性能。
  2. 每次读取都是12M Bytes大小的块,每次写入都是大约1M Bytes大小的,应该不是瓶颈。
    文档里说page-locked memory内存可以在__global__函数里直接访问?为什么我每次操作都失败?
  3. 关于cpu,无.
  4. 一共10个cpu线程,每个cpu线程执行的流程大致如下(每个线程里的gpu计算和拷贝操作都使用cudaStream串行):
    • 使用malloc分配12M内存,从硬盘读取12M Bytes数据
    • 使用cudaMalloc分配24M显存用于输出数据
    • 使用cudaMalloc分配12M显存用于输入数据,使用cudaMemcpyAsync将硬盘读取的12M数据拷贝host->device,等待拷贝完成
    • 1次gpu计算,约15000个block,每个block有128个thread
    • 循环15次gpu计算,约60-2000个block,每个block约有128个thread, gpu计算结果都放置在24M显存里
    • 使用malloc分配24M内存,使用cudaMemcpyAsync拷贝数据device->host,等待拷贝完成
    • 其他cpu计算
    • 写入硬盘约1M Bytes数据
    • 释放所有显存
      问题:我设想,这10个cpu线程之间进行的gpu计算和host ↔ device数据拷贝操作可以重叠并发的。就是说希望有的线程在拷贝host->device,有的在device->host,有的在gpu计算。请问我这样的设计对么?

楼主您好,

page-locked memory的确可以在kernel里直接访问的,您理解的非常正确,但您需要一点点工作:
您需要使用cudaHostGetDevicePointer获取此指针的等价device版,然后使用device版(这不是复制)。
您的cudaHostAlloc()需要指定cudaHostAllocMapped。
在您的程序的最开头使用cudaSetDeviceFlags(cudaDeviceMapHost);。
您进行了这三步后才能直接用。
(或者您立刻买一张tesla卡+TCC驱动,也可以的)

以及,每次您传输12MB么?这个大小可以了,但却只有100-200MB/S的PCI-E传输速度么?您确定您采用的计时方式是准确的?以及,您尝试了我说的用page-locked memory传输的建议了没?

以及,您设想的10个host thread可以同时进行传输和计算是不可能的,任何一种同步,查询,等待,都将阻止这种并行的。如图:

多个流的传输、 计算操作 | 被任何流的同步,查询,等待等操作中断 | 多个流的传输和计算操作

您看,只要某时刻进行了,例如说cudaStreamSynchronize(), 那么任何在这个时刻之前的任何操作和这个时刻之后的任何操作都不能同时进行(包括您说的传输和计算)。

您可以这样修改就能同时进行了:
host线程0启动10个host thread(1-10),
host线程1:复制过去,计算,15次计算,复制回来,线程1结束。(注意不要同步! 下同)
host线程2:复制过去,计算,15次计算,复制回来,线程2结束。
host线程3:复制过去,计算,15次计算,复制回来,线程3结束。

host线程10:复制过去,计算,15次计算,复制回来,线程10结束。
host线程0执行一次WaitForMultipleObjects(), 等待次10个host线程结束,然后在线程0里执行一次cudaDeviceSynchronize()。这样才有可能并行,而不是被这10个线程的各自的同步打乱他们之前的可能的并行性。

只有这样,您才可能实现您上一段的设想(有的在计算,有的在传输到卡,有的从卡传输回来)。

  1. 我目前使用的是非page-locked的显存,明天试试page-locked的结果。
  2. 我现在的做法是,在每个cpu线程里都create一个stream,10个线程共有10个stream,线程里有用到cudaStreamSynchronize(stream)来同步,以确保device->host数据拷贝完成后才进行后续的cpu计算。
    我理解这个cudaStreamSynchronize只同步同一个stream的操作有效吧,不会同步其他线程里的操作吧?
    10个线程各自不断的循环前述的流程,我以为应该是可以达成我的设计目标吧。为什么你说cudaStreamSynchronize会造成这种线程间的操作并发无效呢?

请您重新阅读4#, 谢谢!

为何一个stream的同步会干扰其他stream以及多个stream间的同步是因为事实如此,非要我说个理由。。。。里可以找找老黄问为何这么设计。我只知道事实,但不知道内因。

这是我查到的文档。你说的老黄是哪位?
5.4.3.6 cudaError_t cudaStreamSynchronize (cudaStream_t stream)
Blocks until stream has completed all operations. If the cudaDeviceScheduleBlockingSync flag was set for this
device, the host thread will block until the stream is finished with all of its tasks.

老黄是NV的老大。。。。

话说您查到的这个并不能否定我的帖子,您这个是cudaStreamSynchronize()能对一个流进行同步,这点没错。但这点无法排除不能干扰其他流中的工作。您觉得呢?

能举个简单的例子来说明一个stream的同步,导致对其他stream的影响吗?

这个随意跑个profiler就能看到的。一个cudaStreamSynchronize()将明显的将之前和之后的工作阻隔成2块的。

以及,这个是事实,您不相信无需强迫。简单的不相信我的话即可。

以及我之前给出的如何避免同步干扰的示范代码也请无视。您可以继续坚持自我。谢谢。

继续昨天的测试.

  1. 先测试SSD性能,发现SSD读写是个瓶颈!!!
    因此改写了测试代码, 不把结果写入硬盘, 而且预先读入一组12M数据, 以后每次处理只从内存拷贝这12M数据.
    屏蔽了SSD读写之后,使用8-20个线程分别测试,每次得到的结果几乎一样:
    • 整体速度大幅度提升了70%。cpu可以跑到80%。
    • 从Host到Device的数据,大约200M Bytes/s,低于显卡指标
    • 从Device到Host的数据,大约400M Bytes/s,低于显卡指标

经检查发现,gpu负载还是远远没有跑满!
继续分析.

  1. host ↔ device涉及到的内存改为用page-locked内存。测试结果是对整体性能提升几乎没有影响。
    代码:
    cudaHostAlloc(&p, size, cudaHostAllocDefault);
    还是用cudaMemcpyAsync来拷贝。不知道如果不做拷贝,而是用map过去的device指针来操作,速度是不是有提升?

  2. 我在一个kernel函数里是用到了4K shared memory,profile显示shared memory限制了gpu占用率。我要改为每次执行kernel函数时,动态改变sharedmemory大小(不是每次都要用4K的)。怎么做到呢?

还有个问题
我的程序是32位的,用Visual profile(64bit), 每次run profile都失败。
改为64位后,可以被Profile.
一定要是64bit的菜可以吗?

LZ您好:

1:SSD的问题我不清楚,暂无法评价。亦尚不清楚您为何divice和host间的copy是这样的速度,或许和您的代码实现有关。

2:page-locked memory才能真正使异步copy的runtime函数发挥作用,若不使用page-locked memory,那么您的异步copy的函数依然可以运作,只是无法产生计算和传输相掩盖的作用。关于计算和传输相掩盖,传输和传输相掩盖的内容,不久前的帖子有详细讨论,您不妨翻看一下。

以及,zero-copy一般建议少量的只用一次的数据访问的时候使用。如果量比较大或者需要多次使用,请主动copy。

3:可以使用动态的shared memory的,请参阅programming guide,具体章节已经在另外的帖子为您指出。

12#答复:
我这里可以使用64bit的profiler(CUDA 5.0 64bit版自带的)跑32bit的工程编译出来的exe文件,所以未能重现您的问题。估计您还有其他问题造成此现象。

祝您编码顺利~

非常感谢!!

不客气的,欢迎您常来论坛~

祝您好运~