DAY23：阅读WDDM和TCC模式
3.5.?Mode Switches
GPUs that have a display output dedicate some DRAM memory to the so-called primary surface, which is used to refresh the display device whose output is viewed by the user. When users initiate a mode switch of the display by changing the resolution or bit depth of the display (using NVIDIA control panel or the Display control panel on Windows), the amount of memory needed for the primary surface changes. For example, if the user changes the display resolution from 1280x1024x32-bit to 1600x1200x32-bit, the system must dedicate 7.68 MB to the primary surface rather than 5.24 MB. (Full-screen graphics applications running with anti-aliasing enabled may require much more display memory for the primary surface.) On Windows, other events that may initiate display mode switches include launching a full-screen DirectX application, hitting Alt+Tab to task switch away from a full-screen DirectX application, or hitting Ctrl+Alt+Del to lock the computer.
If a mode switch increases the amount of memory needed for the primary surface, the system may have to cannibalize memory allocations dedicated to CUDA applications. Therefore, a mode switch results in any call to the CUDA runtime to fail and return an invalid context error.
3.6.?Tesla Compute Cluster Mode for Windows
Using NVIDIA's System Management Interface (nvidia-smi), the Windows device driver can be put in TCC (Tesla Compute Cluster) mode for devices of the Tesla and Quadro Series of compute capability 2.0 and higher.
This mode has the following primary benefits:
It makes it possible to use these GPUs in cluster nodes with non-NVIDIA integrated graphics;
It makes these GPUs available via Remote Desktop, both directly and via cluster management systems that rely on Remote Desktop;
It makes these GPUs available to applications running as a Windows service (i.e., in Session 0).
However, the TCC mode removes support for any graphics functionality.
本文备注/经验分享：
有些卡是有显示器接口的，要负责显示。而有些卡没有接口， 不负责显示（或者有的卡有接口， 但是你没有接显示器）。一般有这三种可能：在第一种可能中， 会遭遇一种情况，有一部分显存要存放最后的显示内容， 然后显卡从显示接口， 将最后的现实内容（例如来自一系列的绘制的最终结果）发送给显示器。这些即将要被显示的内容， 叫主表面。主表面也就是和你的显示器上看到的内容对应的表面。你在显示器上看到的每个点， 对应主表面中的一个元素（请注意这不是精确的定义， 因为还有overlay这种东西存在。但这超出了CUDA手册的内容）。所以你直接将主表面理解成对应显示器内容的那个表面即可 （或者等效的， 理解成一些点构成的矩阵）。 这样的话，主表面既然是一个矩阵（概念上的），它需要占据一定的存储空间（显存的）。例如说，你显示器是1024 * 768 * 32-bit颜色（每个点4字节），需要占据1024 * 768 * 4 / 1024 =大约3MB的大小。 如果这个负责显示的卡，正在运行一个全屏的游戏， 分辨率是1024 * 768，而桌面的分辨率却是1920 * 1080，那么当你不小心按住了WIndows键，或者CTRL-ALT-DELETE之类的，或者ALT-TAB之类的热键，从游戏切换出来。那么显示大小会发生变化（因为分辨率不同， 色深也可能不同。 例如有些游戏只是用16-bit色），原本需要3MB显存，现在可能需要8MB显存（假设），需要多出5MB显存给主表面，因为主表面是即将要被发送给显示器绘制的表面，一般总是固定使用显存分配的，此时如果恰巧显存用光了（例如你有512MB的显存，当前剩余1MB空闲）。 那么WDDM驱动可能会选择“自动的”干掉你的某个CUDA程序，将它的context（和context里面分配的内容）摧毁，释放出足够的空间来。这个时候，你的CUDA程序会挂掉。因为CUDA的Runtime和Driver API都有返回之前的错误代码的现象，你会发现突然从某一刻起，你的所有的CUDA调用都持续的返回错误（无效context，因为已经被干掉了），这个是以前要注意的问题。
现在这个问题基本已经不重要了，因为：（1）现在基本都是液晶显示器了（除了高端美工之类的领域），而为了适应液晶显示器，现在的游戏基本上都是使用固定的最佳分辨率，然后将图片绘制或者拉伸到这个最佳分辨率上。（2）现在的显卡显存正在爆炸式的扩大。以前可能一张512MB显存的显卡就很大了，现在这是垃圾卡。例如现在是16GB（16384MB），你很难正好用到16380MB左右，此时切换一下显示模式，所造成的可能的额外的显存大小使用增加，并不会造成显存耗尽。 从而不会让显卡驱动别无选择，去吃掉你的一个或者多个CUDA应用的显存。所以说，实际上手册这里可以删掉了。但是没有。所以我们应当知道这点，因为无法知道以后的显示器会如何。特别的，如果以后的一个显示器起步都是4K的，甚至8K的， 甚至更多（例如10年后，显示器都变成了888K显示器--假设的）此时一个显示模式切换可能造成非常大的显存分配量变化。此时可能会对CUDA应用（也需要显存）造成干扰。
第二段说了TCC的好处。在Windows Vista+（不算XP），包括Vista，7，10等等，有一种叫WDDM的东西（WDDM=Windows Display Driver Model）
。WDDM导致了很多问题。例如这个问题： your_kernel<<<>>>(....); Your Host Code continues your host side work Now you call cudaDevice/Stream/EventSynchronize() 这个是曾经手册上说过的一个技巧。就是发布了一个kernel后（第一行）Kernel是异步执行的，然后Host此时可以立刻通过调用多种等待方式之一（例如cudaDeviceSynchronize()), 来阻塞住执行，来等待异步的kernel在GPU上慢慢执行完成。也可以中间插入一些Host Code，来同时计算一些其他的Host上的任务（例如准备下一次的kernel启动所需要的数据），这样一旦当前kernel完成，下一个kernel就可以很快的继续被发布了。但是WDDM引入了一个kernel启动延迟的问题。在WDDM驱动下，一次kernel启动（或者一次异步传输之类的）并不能立刻发送给显卡执行的。 WDDM驱动会默默的积攒一些，打包一批，然后集体发送给显卡来执行。具体为何会这样涉及到操作系统，我建议你看一下操作系统的课本。具体说这里会需要进行一次或者多次系统调用(system call), 进入内核态（OS概念里的kernel，不要弄混），才能发送命令给卡。
这样的话，如果是上面那三句： 发布GPU kernel启动命令 CPU继续见缝插针的，干一些活 CPU估计差不多了，等待kernel完成 ——这种在WDDM驱动下是有问题的。 第二句并不能和GPU端的kernel同时在执行。 因为此时很可能任务没有发布出去（WDDM的积攒特性）。所以这种技巧非但不能节省总体时间，减少下次kernel的启动准备时间，反而会浪费时间。 此时的一个解决方案是立刻进行流查询操作(cudaStreamQuery)，此runtime api函数具有不公开的效果， 可以立刻要求进行系统调用，将kernel启动操作要求立刻发布给卡。这样这个技巧才能生效。而TCC驱动，没有WDDM的这个问题。
TCC驱动允许不进行系统调用，不切换进OS的内核态，就能直接从用户态发布命令给显卡，此时不仅仅降低了CPU使用，也减轻了kernel启动延迟（因为CPU上不需要进行昂贵的系统调用了）。因此你会看到在TCC驱动下，有更好的整体性能和kernel启动延迟。这是TCC的一大好处。但是手册不会说这个。因为说了MS会不高兴，但是我们应当知道这个。 也很从容易从profiler的令人意外的时间轴(kernel居然被延迟了这么久才启动）中，发现这点。
而TCC驱动让你无需考虑这些技巧，直接就能享受到福利。这是我的话。然后手册上给出来3大主要好处：
It makes it possible to use these GPUs in cluster nodes with non-NVIDIA integrated graphics;
It makes these GPUs available via Remote Desktop, both directly and via cluster management systems that rely on Remote Desktop;
It makes these GPUs available to applications running as a Windows service (i.e., in Session 0).
而第一点我没有看懂。我说一下后两点。 以前Windows如果想远程使用显卡，WDDM模式的卡是不行的，不能直接通过内置的远程桌面功能连接，而只能用第三方的效果不好的，例如VNC或者TeamViewer之类的连接。后两种软件效果远比远程桌面差。此时必须使用TCC驱动才能无障碍的远程桌面使用。但是此特性已经被放开了，现在普通的显卡（例如Geforce的），不需要TCC就能直接远程桌面了，CUDA可以正常使用的。所以这TCC三大好处中的第二点，已经作废了。 然后第三点是说，以前非TCC的卡只能开发普通CUDA应用，Windows Service（这个是什么，请搜索一下就知道了）中是不可能用的。直接找不到可用的CUDA设备的。必须需要TCC驱动才可以开发Windows Service应用，在里面使用CUDA，然而这TCC的第三点好处也作废了。现在普通卡，无TCC，也能在Windows Service中使用CUDA了。
其实TCC还是有其他好处的。例如说，常见的万恶的WDDM超时。切换成TCC驱动就没事了。 省得以前还得想各种技巧来规避它。
论坛就有例子：http://bbs.gpuworld.cn/thread-58808-1-3.html
例如这里介绍的两种技巧。这两种技巧，如果有TCC驱动，可以直接不用学习。这是TCC的又一大好处。以及，还有很多领域需要P2P Access的，例如深度学习。 而WDDM不支持P2P Access，必须通过内存中转，极大的降低了性能，甚至跑不起来。而此时更换TCC驱动，可以直接享受和Linux下一样的效果。（其实这是为何很多深度学习机器都是配套Linux的原因，不是因为Windows的授权费用，Windows的钱其实相比一台机器并不贵），这是第三点好处。
有不明白的地方，请在本文后留言
或者在我们的技术论坛bbs.gpuworld.cn上发帖