CUDA新手要首先弄清楚的这些问题

1

问：当下一个新的GPU架构发布时，我必须重写我的CUDA内核吗? 答复：不需要重写的，CUDA具有高层次的描述能力（抽象能力），同时CUDA编译器生成的PTX代码也不是固定于特定硬件的。这样在运行的时候，驱动负责将PTX代码，转换成当前的特定的GPU上的二进制代码。而每当一个新的GPU发布时，驱动程序也随着更新，因此能不断将PTX转换成未来的所有新一代的GPU上的实际代码来运行。当然你可以根据未来的新GPU上增加的数量, 或者变大的共享内存，对代码手工做出进一步优化，但这是可选的。所以，你无需担忧这个，现在就开始写下你的CUDA代码，享受它在未来的所有GPU上运行的能力吧！

2

问：在一个系统里CUDA可以支持多GPU卡么？

答复：应用程序可以跨多个gpu分配工作。但是，这不是自动完成的，而是完全由你，来控制如何使用多卡。请参阅GPU计算SDK中的“multiGPU”示例，以获得编程多个GPU的示例。

注意部分库可以自动多卡，例如cublas里的一些函数， 但是大部分都是需要用户手工写的。

３

问:CPU和GPU可以并行运行吗? 答复：CUDA中的内核调用是异步的，因此驱动程序将在启动内核后立即将控制权返回给应用程序，然后后面的CPU代码将和GPU上的内核并行运行。当进行性能测试的时候，应当通过CudaDeviceSynchronize()这个API调用，来确保所有的GPU上的任务都完成后，然后再停止（CPU上）的计时器。

4

问：我能同时进行CUDA计算和CUDA数据传输么？

答复：CUDA支持通过多流，在GPU计算和数据传输在时间上重叠/同时进行。

5

问：有可能直接通过DMA，从其他PCI-E设备，直接传输数据到显存中吗？

答复：GPUDirect技术允许你直接这样做，具体可以访问这个页面：

https://developer.nvidia.com/gpudirect

6

问：CPU和GPU之间的峰值传输速率是多少?

答复：内存传输的性能取决于许多因素，包括传输的大小和使用的系统主板的类型。您可以使用来自SDK的bandwidthtest样例来测量系统上的带宽。从页面锁定内存传输更快，因为GPU可以直接从这个内存直接DMA。然而，分配过多的页面锁定内存会显著影响系统的整体性能，所以要小心分配。

7

问：为什么我的GPU计算的结果与CPU的结果略有不同?

答复：可能的原因有很多。浮点计算并不能保证在任何一组处理器体系结构上得到相同的结果。在GPU上以数据并行的方式实现算法时，操作的顺序通常是不同的。

这是一个很好的参考浮点算法:

https://developer.nvidia.com/sites/default/files/akamai/cuda/files/NVIDIA-CUDA-Floating-Point.pdf

8

问：我可以从纹理读取双精度浮点数吗?

答复：硬件不支持双精度浮点作为纹理格式，但它可以使用int2强制转换为双精度，只要你不需要纹理硬件对double进行插值。

texture<int2,1> my_texture; static __inline__ __device__ double fetch_double(texture<int2, 1> t, int i) { int2 v = tex1Dfetch(t,i); return __hiloint2double(v.y, v.x); }

（备注，你可以手工通过代码对这样读取到的double数据，进行人工插值，而不是依赖于纹理硬件自动的插值）

9

问：在哪里可以找到有关PTX汇编语言的文档?

答复：这包含在CUDA工具包文档中。

10

问：如何查看程序生成的PTX代码?

答复：VS里面可以直接在CUDA C/C++属性里改。命令行需要用nvcc -keep选项指定（保留中间文件）。 注意是对你的源文件的编译的过程中产生的，而不是你的程序产生的。

11

问：我怎样才能知道我的内核使用了多少寄存器/多少共享/常量内存?

答复：将选项“--ptxas-options=-v”添加到nvcc命令行。编译时，这些信息将输出到控制台。

12

问：CUDA kernel的最大长度是多少?

答复：因为这可能依赖于你的GPU的计算能力——这个问题的最终答案可以在CUDA C编程指南的特性和技术规范部分中找到。

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications

13

问：我如何选择最优的每个block中的线程数量？

答：为了最大化的发挥GPU性能，你应当仔细平衡block中的线程数量，block使用的shared memory大小，以及，每个kernel线程使用的寄存器数量。尽量提升occupancy往往会提升性能，你可以通过CUDA Occupancy Calculator工具来计算特定的kernel在SM上的占用率（或者occupancy不翻译）。改工具随着最新版的CUDA Toolkit发布。没有直接的答案，这个需要反复试验。

实际上，每个kernel的最佳block形状/其中的线程数量，和具体kernel有关。

精确的说，和具体kernel在具体的某个卡上有关。无法直接确定的，得经过实验。

14

问：最大内核执行时间是多少?

答复：在Windows上，单独的GPU程序启动的最大运行时间约为2秒。超过这个时间限制通常会导致通过CUDA驱动程序或CUDA运行时报告的启动失败，但在某些情况下会挂起整个机器，需要硬复位。

这是由Windows的“看门狗”定时器引起的，如果运行时间超过允许的最大时间，则使用主图形适配器的程序超时。

出于这个原因，可以让负责计算的卡不接显示器。这样就可以规避了。但是需要有加一个独立显卡或者集成显卡作为显示输出。以及，还可以用Tesla上TCC驱动。

15

问：什么GPU卡支持CUDA？

答复：目前市面上NVIDIA　RTX、GeForce, Quadro, Tesla 、GRID 以及Ｊｅｔｓｏｎ平台都支持CUDA。

具体关于计算能力，可以访问这里：https://developer.nvidia.com/cuda-gpus#compute