我有一个程序，并行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%。
检查了输入输出的速度，
问题：为什么不能跑到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都退出就留一个在慢慢跑之外，最大的可能是您有无法掩盖的延迟。
感谢来访。
楼主您好，
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个线程的各自的同步打乱他们之前的可能的并行性。
只有这样，您才可能实现您上一段的设想（有的在计算，有的在传输到卡，有的从卡传输回来）。
请您重新阅读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块的。
以及，这个是事实，您不相信无需强迫。简单的不相信我的话即可。
以及我之前给出的如何避免同步干扰的示范代码也请无视。您可以继续坚持自我。谢谢。
继续昨天的测试.
经检查发现，gpu负载还是远远没有跑满！
继续分析.
host ↔ device涉及到的内存改为用page-locked内存。测试结果是对整体性能提升几乎没有影响。
代码：
cudaHostAlloc(&p, size, cudaHostAllocDefault);
还是用cudaMemcpyAsync来拷贝。不知道如果不做拷贝，而是用map过去的device指针来操作，速度是不是有提升？
我在一个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文件，所以未能重现您的问题。估计您还有其他问题造成此现象。
祝您编码顺利~
非常感谢！!
不客气的，欢迎您常来论坛~
祝您好运~