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上发帖

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-05-30

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏腾讯NEXT学位

小游戏内存优化与性能优化

? | 导语 听说你的小游戏内存超标?进来了解一下吧。 本文主要跟大家一起来探讨一下Cocos Creator小游戏开发过程中内存优化、性能优化和包体优化。 ...

4141
来自专栏SAP最佳业务实践

SAP最佳业务实践:SD–带质量管理的销售退货(237)-6退货库存处理

一、MIGO退货库存转移至自有非限制库存 如果可重新使用退回的物料,则需要将其从退货库存移至非限制自有库存,这样退货可任意使用,并会重新包括在 MRP 计算之中...

2765
来自专栏梦里茶室

【Chromium中文文档】进程模型

这个文档描述了Chromium支持的不同线程模型,包括它的渲染器进程,以及现有模型实现的问题。 概述 网页内容已经发展到包含大量在浏览器内运行的活跃代码的地步,...

19910
来自专栏数据和云

Direct IO+asm引起css initialization

作者简介: ? 何剑敏 Oracle ACS华南区售后团队,首席技术工程师 现供职于Oracle ACS华南区售后团队,首席技术工程师。多年从事第一线的数据库...

3395
来自专栏何俊林

VLC框架总结(一)初识庐山真面目

开头说三件事: 1、周末送书已全部寄出。获取者注意查收,考虑到顺丰着实贵了。发的是百世快递,已都帮大家付了邮费。北京市内6块,外省都是10块。觉得这快递价格还是...

2159
来自专栏数据库

爱炫耀的数据库老头儿

作者:刘欣 1 数据库老头儿 我们这个世界很大, 生活着很多人,形形色色,各怀绝技。但是被公认为最拽的一个却是数据库老头儿,年龄挺大,每天都要炫耀几遍他那关系型...

1710
来自专栏向治洪

android开发性能分析

1 背景 其实有点不想写这篇文章的,但是又想写,有些矛盾。不想写的原因是随便上网一搜一堆关于性能的建议,感觉大家你一总结、我一总结的都说到了很多优化注意事项...

2095
来自专栏腾讯移动品质中心TMQ的专栏

GT3.1简化您的App性能测试(1)——全新的性能监测体系

GT君此次给大家送来了强力的安卓GT3.1版本更新。在本次更新中,采用了在被测应用内嵌SDK的方式来获取流畅值,解决了2.x版本测试应用流畅值必须root手机的...

5246
来自专栏代码GG之家

SDL系列讲解(一) 简介

什么是 SDL Simple DirectMedia Layer(SDL)是一个跨平台开发库,主要提供对音频,键盘,鼠标,操纵杆的操作,通过OpenGL和Di...

2205
来自专栏tkokof 的技术,小趣及杂念

随便再聊一点点Coroutine(确实只是一点点~)

  之前写过一点Coroutine相关的东西(这里和这里),大致讲了些自己关于Unity协程的理解,自己在平日的工作中也确实用到了不少相关的知识,遂而引发了一个...

711

扫码关注云+社区