尊敬的各位版主、高手们，新年好！ 本人是CUDA新手，尝试在LINUX下写过几个CUDA程序，用过的卡有560Ti和C2050，发现用设备函数中所使用的寄存器数量是影响GPU所有资源发挥的一大因素，而用纯C语言写程序很难精确控制寄存器的使用数量（用ptx汇编查看每次都比预想的多），我的个人想法是用汇编语言来实现核心设备函数编写来精确使用寄存器的数量，请问这想法是否正确？如果行的话，有无在linux和在windows下的开发教程或文档、实例来参考学习？
刚在官方找到一个文档了：ptx_isa_3.1.pdf ，不知道有没有人使用过
LZ您好，关于您的问题，个人观点如下供您参考：
1：一般推荐使用CUDA runtime API进行开发，这样最为简洁，编码和学习较快，也最为方便。当然您还可以使用CUDA driver API和PTX 汇编来获取更大的操作性，同时付出开发复杂度上升的代价。
2：编译器会尽可能优化寄存器的使用数量，您可以跑一下NVVP（nvidia visual profiler）看看您kernel寄存器数量的使用情况。如果寄存器占用数量依然超过预期，那么您可以，第一：改写您的kernel使之更加轻量化；第二：限制kernel使用寄存器的最大数量，这样会使用local memory，一般会损失速度；第三，可以尝试kepler架构的GPU，kernel内的寄存器最大可用数量要比fermi架构的GPU更高。
个人推荐第一种做法，因为GPU并行本来就比较适合细粒度的并行任务，当然实际情况也许减小kernel规模是有困难的。
3：因为我不用PTX汇编写代码，因而无法给出您具体的建议，也无法估计您用PTX汇编自己写，是否真的比runtime API编译以后的效果好。如您有进一步需求，请其他人补充。另外，您在2#找到的文档是正确的。
祝您春节愉快~
谢谢您细致的回答！
我用Gpu的目的是提高算法实现速度，即追求最快速度，不怕累的。
您说的对，从算法实现方法上优化是优化的最有效手段，但一般很困难，因为我这做的算法都是经过历史检验过的，不好优化了。
CUDA runtime API应该在PTX汇编里也能调用吧，就象ASM里调用WINDOWS API一样，我肯定不会去重写它的。
谢谢您的建议！换个gTX6系列的卡试试，但较早看网上听说6系列的卡在高性能计算上不一定比5系列的快，所以一直没试。
普通的卡，不能在ptx里调用runtime api。
runtime api实际上是给host code调用的。(不考虑GK110的新特性）。
godscu您的例子“asm调用win32 api"，这个例子是对的。但不具有类比性。
此外，的确，有时候换卡可能会提升性能，但也可能会降低性能。要看你的新卡是什么卡，以及，你是否针对它斟酌过你的代码。但作为尝试，还是不错的，热切建议楼主尝试。
目前的6系列GTX显卡，是SM 3.0的，如同评测所说，通用计算能力并不强，更偏向于图形卡，除了sp数量巨大，以及28nm工艺功耗较低。您可以找来试试看。更强大的面向通用计算用途的kepler是SM 3.5的，目前只有telsa K20/K20X这两款。
此外，3.x的优化特性和2.x还是不同的，这也有待您实践摸索。
大致如上，祝您春节愉快~
线程同步函数是不是API？例如syncthread这样的函数
可以内联汇编，象C一样吧
我会尝试一下的，还是比较有意思的
谢谢版主提醒！要提高性能只能换K20了，GTX的6系列就不考虑了，差点要买了。
估计比C2050还要贵吧，我上网看看，如果太贵还是等等吧
楼主您好，您的想法是错误的：不能"在ptx里精准的控制寄存器使用量的“。
PTX使用SSA(Single Static Assignment)来进行方便识别变量生存期，以及进行优化等措施，所以ptx里面的“字面”上的寄存器使用量是巨大的，稍微大的kernel可以看到几百上千的“寄存器”使用。
但这不是最终结果，最终被汇编/jit成SASS(Shader Assembly)后，实际的寄存器使用会大幅度减少，您唯一目前能做到的是，从ptxas的结果中查看寄存器使用，这个才是精确的。
初次之外，您没有其他能“精确”控制寄存器个数的方法。但是您可以尽量通过修改您的算法实现，以及通过__launch_bounds__()来暗示编译器能尽量或多或少的为您的寄存器数目进行优化。
（本文不讨论是否真的没有精准控制寄存器个数的方法，一切回复均只参考公开文档，请勿较真）
此外，上文说的看ptxas的结果，只针对您不需要jit的场合。反之，此结果不适用。（例如您为sm_10编译，但您运行在sm_20上，则会引入jit过程）。
此外，您还可以通过-maxrregcount参数进行上限约束。（但不一定是好事）。
请注意是rreg(Regular Register)，有2个r，不要打错了。
感谢楼主莅临CUDAZone China,
祝您春节愉快！
（1）
嗯嗯。__syncthreads()的确不是runtime api, 而是device中使用的intrinsic functions。
它会编译为bar指令。单条。
（2）我非常赞同您的观点，可以内嵌汇编，但这是一个编译器特性，而不是语言特性。
（3）您说的对，这是很有意思的。
事实上，寄存器数量限制程序性能的案例还是比较少的。
首先您要明确是否真的是此因素导致了您程序性能无法进一步提升。
寄存器影响主要两个方面：active warp 的数量（即occupancy ）和寄存器溢出导致的local memory的传输。
首先看active warp：
什么时候是因为寄存器使用过多导致active warp数量少，导致性能低呢？
第一，程序为延迟密集型程序，也就是说程序大部分时间是在等待高延迟指令返回结果（其他两种为指令密集型和内存密集型，详见GTC 录像中讲座Analysis Driven Optimization with CUDA http://www.gputechconf.com/gtcnew/on-demand-gtc.php ）。因为只有在这种情况下程序才需要更多的active warp来隐藏延迟。
第二，仅仅通过instruction by byte ratio, instruction throughput, memory throughput来确定是否为延迟密集型是不完整的。很多情况下，当我们的程序有较严重的负载不均衡的现象，即某几个block或某些warps运行时间远超过其他时，或者没有分配足够的blocks和threads时。instruction throughput 和 memory throughput 也会很低。这两种情况的特征是实际的occupancy数量远远低于理论的occupancy。实际occupancy可以通过以下公式求出：
actual occupancy = active warps/duration/#sm/frequency_of_gpu/max_active_warps_in_a_sm
其中active warps 可以通过Toolkit 中提供的profiler：nvvp 测得。duration 为kernel运行时间，#sm为此gpu中sm的数量，frequency_of_gpu为gpu的主频，max_active_warps_in_a_sm为一个sm最大可以有多少个resident warps。理论occupance同样可以通过nvvp测得，或者通过toolkit中提供的excel工具测得。
第三，即使理论和实际occupancy接近，也不一定是由于寄存器使用过多导致occupancy少，occupancy还受到shared memory，每个block分配多少线程的影响，可以通过excel工具查看。
只有在程序为延迟密集型程序，且经验证，理论occupancy和实际occupancy接近时，且通过excel工具差得确实寄存器为限制因素时才可的出寄存器使用过导致性能无法进一步提高的结论。
另外，再来看寄存器溢出导致local memory传输制约程序性能的可能性：
如果是内存传输为程序瓶颈，那么kernel一定是内存密集型。其表现为程序sm到L2内存传输速度基本达到gpu峰值。至于判断是寄存器溢出导致的local memory，还是局部数组存储导致的local memory成为瓶颈还是由于一般的global memory的访问或者tex访问。我们可以通过查看Profiler 中的Metrics Local Memory Overhead查看。看是否local memory的吞吐量占主要部分。在local memory的吞吐占主要部分，且代码中用数组存储局部变量的情况极时，才可判断为寄存器溢出导致local memory传输制约程序性能。
如果寄存器使用却是称为程序瓶颈。以下策略可以减少程序中寄存器的使用。
1、 尽量拆分代码为较小的Kernel。
2、 运用maxrregcount编译选项控制寄存器使用。
感谢WZH同学给楼主，以及为我们和其他会员朋友写了一封长文。
里面给出的说法很好很不错。以及写的非常用心，富含技术含量，深入浅出，娓娓道来。值得我们普通版主学习！不愧是原厂支持！展示了NVIDIA的强大丰富的力量！
那个，在敬仰的同时，我给很出一个我们普通人可以使用的算法：
等效的实际occupancy = achieved_occupancy * sm_efficiency %
这个也许简单点。
建议使用，此中的2个参数用profiler可以直接得到。
再次膜拜强大的NVIDIA原厂工程师！
您的到来使得论坛蓬荜生辉。
祝您春节愉快！
（并建议ICE版主将ZHW置顶）
其实需要再说一下的是，我们一般将此2个参数分开看。
例如，后者(sm_efficiency)可以单独的看出为了等待延迟（或者SM无工作而空闲）而导致损失的效率。
当然，合并起来真心不错。值得赞美！