Monthly Archives: January 2015

无人机发展新趋势:自主导航

版权声明:本文系本站作者自己翻译整理,欢迎转载,但转载请以超链接形式注明文章来源(planckscale.info)、作者信息和本声明,否则将追究法律责任。

在技术领域,通常大佬一给力,相关应用就受益,发展迅猛,甚至能改变行业格局,无人机就是受益者之一。

Intel给力,推出其RealSense技术及RealSense 3D摄像头,德国Ascending科技公司看准时机,搭载RealSense 3D摄像头研发出了一套无人机自动驾驶系统。基于此系统的产品Firefly已在CES上展示,被称作无人机乒乓球。它不需要有人专门抱着遥控器去遥控,而是自己看路自己飞,遇到障碍自动避开。

NVIDIA给力,去年推出全球首款嵌入式超级计算机 Jetson TK1 开发组件,搭载拥有192个GPU核心的Tegra K1芯片,又在今年的CES上推出Tegra X1 芯片:1TFLOPS,P < 15W, ARM Cortex A57 * 4 + ARM Cortex A53 * 4 + Maxwell 256 CUDA Cores。话说Tegra X1的运算效能将比Tegra K1提升3倍。真忍不住给老黄点个赞,把我们的多路视频实时全景拼接移植到嵌入式平台就要靠Jetson TK1 开发组件。目前,法国Parrot公司已经利用Jetson TK1平台打造出了新一代的小型无人机,借助Tegra K1强大的视觉计算能力,不仅可以自动避障,还能实时重建周边环境3D模型,帅爆了。看到这个消息,老哥又扼腕一次。幸好现在各方面条件差不多成熟了,我们的3D重建也要纳入进程了。哈。。

Parrotdemovertical-307x500

计算机视觉现在发展的给力,借助此技术的新公司叫板Skydio,由三个来自MIT的人创立(其中两人曾加入过Google X的Project Wing),他们表示用两个最普通的摄像头,就能实现无人机的自主导航及自动避障,而不依赖于激光,声呐或3D相机。他们的目标是用50刀实现5000刀的梦想。他们把两个普通的微型摄像头采集的视频流注入到一个小的Intel媒体中心。利用计算机视觉算法检测障碍物并及时自动避障。他们通过Wi-Fi使手机与无人通讯,用手机代替遥控器,手机对准无人机,开启启动,手臂挥到哪儿无人机就跟到哪儿,比其他的”follow me” ”无人机接口更简单。另外也可在手机上预设好飞行路线,让其自主飞行。

目前的无人机多采用GPS导航,同时需要很多训练过的经验丰富的无人机操作人员花费大量精力来操作,无人机自主导航将会是无人机发展的一大趋势。

Intel RealSense(实感技术)概览

版权声明:本文系本站作者自己翻译整理,欢迎转载,但转载请以超链接形式注明文章来源(planckscale.info)、作者信息和本声明,否则将追究法律责任。

前段时间CES的报道满天飞,网上逛了几圈,感觉料最猛的还数Intel:老树开新花,推出14nm的第五代酷睿处理器;在智能可穿戴设备及物联网上雄起;RealSense实感技术开启未来人机交互模式。之前总听老哥讲,鼠标键盘这些传统交互模式统治了二三十年,目前跟踪识别爆火且技术趋于成熟,新的交互革命快要打响了。看RealSense这阵势,估计不出三五年新的交互便会普及开来。

英特尔早在2012年左右就着重研发实感技术,当时叫Perceptual Computing,即感知计算,并开放英特尔® 感知计算软件开发套件 2013 版(Intel® Perceptual Computing Software Development Kit, SDK 2013),设重奖举办因特尔感知计算挑战赛,吸引众多开发者参与。随着技术完善与成熟,2014年初更名为RealSense,即实感技术,而后发布了新的Intel® RealSense™ SDK 2014 ,同时举行2014英特尔®RealSense™应用挑战赛

基于此技术的应用在IDF2014及今年CES上大放异彩。

Intel® RealSense™ SDK的架构:SDK core,I/O module和Capability modules组成整个SDK堆栈的基础。SDK core管理I/O module和Capability modules同时组织并管理管线的执行。I/O module捕获设备的输入数据而后发送数据到输出设备或Capability modules。Capability modules也叫algorithm modules,主要包括各种模式检测和识别算法(面部跟踪和检测、手部跟踪、手势识别、语音识别及合成等)。

01

各功能:手部和手指跟踪、面部分析、语音识别、背景移除、目标跟踪、增强现实、3D扫描。

1. 手

SDK将手抽象出骨架,并从背景中剥离出来。允许在照相机的0.2–1.2米范围内跟踪手上的22个点的位置和方向,如图。左右手是区分的,因而可以双手进行交互。

02

手势识别:包括静态手势识别和动态手势识别。SDK中内嵌了一系列手势如下图。你可以用内嵌的这些手势组合出新的手势,也可以根据手骨架上那22个点位置创建出新的手势。

 

0304

05

动态手势识别:静态手势可以单独使用,当然也可以组合使用形成一定的动作。组合时要求开始的手势和结束的手势都已在手势识别系统中注册。

06

SDK给出了一系列动作如下:

07

 

2.脸

脸部检测:SDK提供精确的3D脸部检测和跟踪,且可以同时跟踪4张人脸。每张人脸用长方形来标记,你可以获得长方形的XYZ坐标。与2D跟踪相比,3D头部跟踪在头部运动方面更给力。

QQ图片20150118224438

脸部识别:SDK提供识别特定人脸的能力。特定ID对应注册的特定人脸,并将此人脸的信息存储到人脸库的内存中。如果同张人脸被注册多次,那这张人脸被正确识别的机会将会增大。当没被识别的人脸出现时,识别模块将与数据库中的可能数据进行比对,如果找到匹配则返回此人脸对应的ID。

用户不用担心自己的头像被存储,因为存储的只是算法从图像中提取的特征的集合。

头的运动:SDK提供头部运动的3D方向:俯仰、左右转动、左右偏转,如图。

08

所以可以轻松获得用户头部指向哪里。也可以以此做粗略的眼神跟踪,下一版将会推出更精细的眼神跟踪。

标记点跟踪:SDK提供脸部78个标记点的跟踪以提高脸部识别和分析的精确度。在图像和坐标系中给出这78个标记点的位置。脸部标记点跟踪支持头像创建、脸部动画、拟态及简单的表情识别。可以直接用这些点或这些点的相对位置来作分析。不管你有没有刘海、戴不戴眼镜,这种标记点跟踪都支持。但用户头部在屏幕30度内效果最好。

09

面部表情识别:SDK也包括更高级的面部表情识别。这使得你创建卡通头像更简单。每一种表情有1到100的强度等级使得你做的动画更平滑自然。SDK中表情:

001

情感识别:SDK中的情感识别算法用的是2D RGB数据。情感模块是独立的模块,并非脸部模块的一部分。为保障情感识别正常工作,图像中的人脸至少要有48×48个像素。此算法并不局限于RGB数据,灰度数据同样可行。利用SDK,你可以检测并估计以下六种原始情感的强度。

002

头像控制:SDK通过结合面部表情及78个标记点提供简单的头像控制功能。SDK提供动画角色的示例代码,使你的应用可以适应任何脸型并把用户头像动画化。

应用场景:

003

 

3. 语音

主要包括语音命令和控制、听写、从文字转译成语音等功能。只支持英语,语音识别对成年人效果最好。

语音识别:分为命令模式和听写模式。命令模式需提前设定命令列表,特定命令绑定特定动作。听写模式内置了一个通用的词典,包括50k个常用单词。如果你觉得不够用,也可以自己添加词典中没有的单词。听写模式限时30秒。命令和听写模式不可同时开启。

004

语音分析:SDK也可以根据文本动态的生成语音,由女声读出。

005

4. 背景移除

可以实时的移除背景并替换成新的背景。

006

应用场景:

007

5. 目标跟踪

Metaio*3D目标跟踪模块提供基于光学的跟踪技术,可以跟踪视频序列或场景中的目标物。Metaio工具箱可以训练、创建并编辑3D模型,这些模型可以传给各种目标检测及跟踪算法。

跟踪技术支持平面的2D目标跟踪、基于特征的3D跟踪、CAD模型的基于边界的3D跟踪及即时3D跟踪。

009

6. 增强现实

用音频、视频、图像或其他信息来添加、增强或补充英特尔实感技术游戏的内容。

QQ图片20150118210357

7. 三维扫描和打印

可以扫描、编辑、打印和分享三维物体,并与 3D Systems展开合作。

000

目前因特尔已经与腾讯合作创立游戏创新实验室,推进实感技术,并打造了《轩辕传奇》。与京东合作建立“京东因特尔联合创新实验室”,推出虚拟试衣、3D物品展示等,使实感技术落地电商平台。同时因特尔自己也在不断推出自己的实感技术产品。这不前两天在北京推出其全新的第五代酷睿处理器家族,搭载此处理器的多款产品还配备了英特尔实感技术,看来这项技术很快就会遍地开花。

只要你的处理器是第四代及四代以后英特尔® 酷睿™处理器,操作系统满足Microsoft* Windows 8.1(仅限64 bit),再外购一个英特尔® RealSense™ 3D 摄像头就可以下载SDK玩起来。估计不久联想、戴尔、华硕、宏基、惠普等等搭载第五代酷睿处理器并内嵌因特尔3D摄像头的超极本、二合一及一体机设备就会上市,那时玩起来会更爽。

资料来源:RealSenseSDKdesignGuidelinesGold

 

CUDA, 软件抽象的幻影背后 之三

版权声明:原创作品,欢迎转载,但转载请以超链接形式注明文章来源(planckscale.info)、作者信息和本声明,否则将追究法律责任。

上一篇中谈到了编程模型中的Block等概念如何映射到硬件上执行,以及CUDA如何用并行来掩盖延迟。这一篇继续剖析SIMT,谈一谈控制流分叉,指令吞吐和线程间通讯机制。
虽然我们说warp中的线程类似于SIMD,但事实上它是真正的线程。warp中的每一个thread都有自己的指令地址寄存器,允许它们各自执行不同的任务(控制流分叉)。最简单的,比如一个

if(threadIdx < 10)
 {}
 else
 {}

语句,将threadIdx=0…31这一个warp划分成两个分支,各自做不同的事情。这个灵活性以性能为代价,当一个warp中控制流出现分叉时,不同分支的线程会被分组相继执行,直到各分支执行完毕后,控制流重新汇聚成一支(上例中即if语句的结束点)。这种情况下执行单元的利用率较低,因为每个分支执行时都需要关闭其他分支的线程,所以这时一些执行单元是用不到的。
为了尽可能高效的计算,需要约束控制流分叉的出现。除了减少流程控制语句外,还需要注意,并不是只要有流程控制语句就一定会带来控制流分叉。关键是,控制流分叉只是针对同一warp中的线程而言,不同warp的线程原本就是串行化执行的,分叉对其无影响。因此,只有流程控制语句的条件在
同一warp内不一致时,才会有控制流分叉。这样,诸如

if(threadIdx.x / WARPSIZE < n)
{...}
else
{...}

这样的语句是不会有分叉的。当然,更宽松的条件如

if(blockIdx.x < n)
{...}
else
{...}

也不会有分叉。依赖于输入数据的条件如

if(globalArray[threadIdx.x] < n)
{...}
else
{...}

则会带来分叉。

对于简单的指令如32位浮点数的加、乘,32位整数的加减等,通常CUDA Core在一个时钟周期内可以完成一次操作,因而一个周期内完成的操作数目就等于CUDA Cores数目;而对于一些较复杂的指令,如sin/cos等超越函数,执行单元并不能提供这么高的吞吐率。
我们可以用单位周期内进行的操作数目N除以32来计算指令的吞吐率。以GM204为例,它的SM中有32*4 = 128个CUDA Cores,32个SFU(特殊函数单元),在计算32位浮点加法时具有最高吞吐,一个周期内完成128次操作,单位周期内指令吞吐为128/32 = 4;而计算如sin/cos等超越函数时线程不再一一分配到CUDA Cores上,而是要在32个SFU上计算,单位周期内只能完成32次操作,指令吞吐为1条指令每周期.
指令的吞吐率数据可参考CUDA C Programming Guide中 5.4.1. Arithmetic Instructions,该小节以单位时钟周期每SM上能够进行的操作数的形式给出了各指令的吞吐率。
指令吞吐率是我们进行性能优化的有一个重要指标。通常,影响指令吞吐率的因素除了数值计算操作的复杂度、精确度之外,控制流分叉也是一个贡献因子。这里的原因不难理解,控制流分叉时执行单元的利用率下降,使得单位周期内执行的操作数目下降,从而降低了指令吞吐。

到这里,硬件图景下线程的执行就基本说完了,只剩下一个留到最后的话题:线程间交互。通常,不存在任何相互作用的线程,它们之间才能够以任意的顺序执行,像block。但对于warp这样的线程组,是可能与同一block中其他warp通讯或同步的,这时执行顺序就不能任意。所幸即便在block之内,线程间的交互仍然是较弱的,因而底层可以将block划分成warp来分组串行化执行,遇到交互时再另作处理。我们现在来看看这些交互机制。

线程间交互可以细分为通讯和同步两类。通讯主要由公共存储区域交换数据来实现,但也不排除像shuffle这样的特殊方式存在。
从通讯的粒度来看,可以分为warp内部线程间通讯,block内部线程间通讯,block间通讯,更粗的粒度这里不考虑。block之间的通讯则只能基于global memory,block内部的通讯主要基于shared memory/global memory,warp内部线程间除了可以利用上述所有方式,还有一种特殊的shuffle机制.下面我们以通讯的粒度分类陈述各种通讯的实现方式。

block间通讯通常基于两次kernel发射,一次将通讯数据写入global memory,另一次发射读global memory进行后续处理。这种通讯开销较大,主要来自于global memory访存和kernel发射,所以如果有可能,尽量把任务放在一次kernel发射中完成。
或许有人会问,同一个kernel发射中的两个block具有共同的global memory,是不是也可以利用这个特点来构造同一kernel下block间的通讯呢?通常的答案是no,因为block之间执行顺序不定,很难构造有意义的通讯;但如果要较真,答案是yes,我们真的可以构造一些特殊的block间通讯方式。一个例子如下所示,该实例来自于CUDA C Programming Guide B.5. Memory Fence Functions:

__device__ unsigned int count = 0;
__shared__ bool isLastBlockDone;
__global__ void sum(const float* array, unsigned int N,
	volatile float* result)
{
	// Each block sums a subset of the input array.
	float partialSum = calculatePartialSum(array, N);
	if (threadIdx.x == 0) {
		// Thread 0 of each block stores the partial sum
		// to global memory. The compiler will use
		// a store operation that bypasses the L1 cache
		// since the "result" variable is declared as
		// volatile. This ensures that the threads of
		// the last block will read the correct partial
		// sums computed by all other blocks.
		result[blockIdx.x] = partialSum;
		// Thread 0 makes sure that the incrementation
		// of the "count" variable is only performed after
		// the partial sum has been written to global memory.
		__threadfence();
		// Thread 0 signals that it is done.
		unsigned int value = atomicInc(& count, gridDim.x);
		// Thread 0 determines if its block is the last
		// block to be done.
		isLastBlockDone = (value == (gridDim.x - 1));
	}
	// Synchronize to make sure that each thread reads
	// the correct value of isLastBlockDone.
	__syncthreads();
	if (isLastBlockDone) {
		// The last block sums the partial sums
		// stored in result[0 .. gridDim.x-1]
		float totalSum = calculateTotalSum(result);
		if (threadIdx.x == 0) {
			// Thread 0 of last block stores the total sum
			// to global memory and resets the count
			// varialble, so that the next kernel call
			// works properly.
			result[0] = totalSum;
			count = 0;
		}
	}
}

代码 1. block间通讯实现数组求和
本代码摘录自 CUDA C Programming Guide B.5. Memory Fence Functions

该例实现一个数组的求和,首先各个block计算部分和,然后由最后一个完成部分和计算的block再把所有的部分和加和出最终结果。block间通过一个位于global memory的变量count通讯,它记录了目前已经完成计算的线程数。这样,最后一个完成部分和计算的block就会发现count的数值为最大线
程id,因此可以判定需要由它自己来完成最后从部分和向总和的计算。
不过,为了更好的软件结构,最好还是避免同一kernel的block间产生耦合。同一kernel中block的通讯还涉及到CUDA的weakly-ordered内存模型问题,一个线程中先后两次内存操作在另一个线程看来未必能够保持原有顺序,这产生了相当大的复杂性。我们在下文还会提到这一问题。

block内的线程通讯机制较为丰富,尤其是线程同属一个warp时的shuffle机制。shuffle在Kepler后出现,是一种相当快的线程间通讯方式,它允许同属一个warp的线程间可以互相引用彼此的寄存器,比如下例:

__global__ void bcast(int arg)
{
	int laneId = threadIdx.x & 0x1f;
	int value;
	if (laneId == 0) // Note unused variable for
		value = arg; // all threads except lane 0
	value = __shfl(value, 0); // Get "value" from lane 0
	if (value != arg)
		printf("Thread %d failed.\n", threadIdx.x);
}

代码 2. shuffle机制实现一个值向整个warp的广播
本代码摘录自 CUDA C Programming Guide B.14. Warp Shuffle Functions

laneId是warp中线程的一个index,有threadIdx对32取余得到。__shfl(value, 0)语句使得各线程能够访问laneId==0这一线程中value的值。

更常用的通讯机制自然是shared memory和global memory了。其中shared memory更快速,在大多数时候是构建高性能CUDA程序的必由之路。这些常识不再赘述。基于shared/global memory的线程间数据交换,一定要注意线程的同步。block中线程的同步由__syncthreads()实现。线程会等待同block中其他线程都执行到这一点,并且__syncthreads()语句之前的所有shared/global memory操作都尘埃落定,保证block内所有线程在__syncthreads()之后都能看到这些操作的结果。

最后谈一下CUDA采用的weakly-ordered内存模型。它导致一个线程中相继执行的两个存储器操作在另一个线程看来未必是一样的顺序。例如:

__device__ int X = 1, Y = 2;
//thread 0
__device__ void writeXY()
{
	X = 10;
	Y = 20;
}



//thread 1
__device__ void readXY()
{
	int B = Y;
	int A = X;
}

代码 3. weakly-ordered内存模型示例
本代码摘录自 CUDA C Programming Guide B.5. Memory Fence Functions

这段代码可能产生A=1,B=20这样的结果。原因是有多种可能的,要么thread 1看到的X、Y的写入顺序被颠倒,要么thread 1中读取顺序被颠倒。这种看似相当毁三观的事情确确实实发生在我们的代码背后。在一个线程里两个相继但无依赖的内存操作,其实际完成的顺序可能是不确定的。在这个线程
看来这并没有导致什么不同,因为两个操作无依赖,并不会破坏因果链;但在另一个线程的眼里,它就暴露出来了。
忍不住插句嘴,这简直就是狭义相对论的世界观在计算机世界的翻版:一个参考系的观察者所看到两个类空间隔事件(可以是相继发生但因距离遥远而无因果关联)在另一个参考系中看来是颠倒的,但有因果关联的两事件在所有观察者看来时序都不会改变。好玩吧?

所以,表面的秩序井然背后有着巨大的复杂性怪兽,为了关牢它的笼子,我们需要约束我们的代码,用合适的机制来实现线程间通讯。要保证另一个线程看起来,两组存储器操作具有我们所希望的顺序,需要用 Memory Fence Function. 这里不再涉及,对更多细节感兴趣的同学,请参考CUDA C Programming Guide B.5. Memory Fence Functions等章节。
(未完待续)

CUDA, 软件抽象的幻影背后 之二

先更新到这儿,稍后再回来抛光查错。CUDA比较杂,我一写起来容易满嘴跑火车弄出错误,欢迎拍砖。

**********************************************************************

版权声明:原创作品,欢迎转载,但转载请以超链接形式注明文章来源(planckscale.info)、作者信息和本声明,否则将追究法律责任。

上一篇里说到,有两点对CUDA的计算能力影响甚大:数据并行,以及用多线程掩盖延迟。接下来我们要深入到其硬件实现,看一看这些机制是如何运作的。

通常人们经常说某GPU有几百甚至数千的CUDA核心,这很容易让人联想到多核CPU。不过事实上两种“核心”是不一样的概念,GPU的CUDA核心只相当于处理器中的执行单元,负责执行指令进行运算,并不包含控制单元。可以类比到CPU核心的是流多处理器(Streaming Multiprocessor,简写为SM. Kepler中叫做SMX,Maxwell中叫做SMM),通常一个GPU中有数个SM,而每个SM中包含几十或者上百个CUDA核心,以及数个warp scheduler(相当于控制单元)。如下图GM204中有16个SM,每个SM中有128个CUDA核心,4个warp scheduler。

GeForce_GTX_980_SM_Diagram-545x1024

图 1.  GM204的SM结构图

每个SM中有大量的寄存器资源,在GM204的例子中,有总共64k 32-bit寄存器,可以养活成千上万的线程。SM中另外一个重要资源是Shared Memory,没错,它正是软件抽象中Shared Memory的对应物。在GM204中,每个SM有96KB的Shared Memory.

到这里,SM在软件抽象里的对应也呼之欲出了,没错,正是Block。我们不妨先摆出这个对应:
Block <-> SM
Thread执行 <-> CUDA Cores
Thread数据 <-> Register/Local Memory

同一Grid下的不同Block会被分发到不同的SM上执行。SM上可能同时存在多个Block被执行,它们不一定来自同一个kernel函数。每个Thread中的局域变量被映射到SM的寄存器上,而Thread的执行则由CUDA核心来完成。

SM上可以同时存在多少个Block?这由硬件资源的消耗决定:每个SM会占用一定数量的寄存器和Shared Memory,因此SM上同时存活的Block数目不应当超过这些硬件资源的限制。由于SM上可以同时有来自不同kernel的Block存在,因此有时候即便SM上剩余资源不足以再容纳一个kernel A的Block,但却仍可能容纳下一个kernel B的Block.

接下来一个很重要的问题是Block如何被执行。我们可以看到,SM上的CUDA核心是有限的,它们代表了能够在物理上真正并行的线程数——软件抽象里,Block中所有的线程是并行执行的,这只是个逻辑上无懈可击的抽象,事实上我们不可能对一个任意大小的Block都给出一个同等大小的CUDA核心阵列,来真正并行的执行它们。
因而有了Warp这个概念:物理上,Block被划分成一块块分别映射到CUDA核心阵列上执行,每一块就叫做一个Warp.目前,CUDA中的Warp都是从threadIdx = 0开始,以threadIdx连续的32个线程为一组划分得到,即便最后剩下的线程不足32个,也将其作为一个Warp.CUDA kernel的配置中,我们经常把Block的size设置为32的整数倍,正是为了让它能够精确划分为整数个Warp(更深刻的原因和存储器访问性能有关,但这种情况下仍然和Warp的size脱不了干系)。
在GM204的SM结构图里我们可以看到,SM被划分成四个相同的块,每一块中有单独的Warp Scheduler,以及32个CUDA核心。Warp正是在这里被执行。
Warp的执行非常类似于SIMD. Warp中的活动线程由Warp Scheduler驱动,同步执行。我们可以看到,GM204中32个CUDA核心共享一个Warp Scheduler. 关于Warp执行中可能出现的复杂些的问题,留到下文另外说。

现在可以整理一下这个世界的图景了。SM上存活着几个Block,每个Block中的变量占据着自己的寄存器和Shared Memory,Block被划分成32个线程组成的Warp. 这样,大量的Warp生存在SM上,等待被调度到CUDA核心阵列去执行。

Warp Scheduler正如其名,是这个Warp世界里的调度者。当一个Warp执行中出现等待(存储器读写延迟等)后,Warp Scheduler就迅速切换到下一个可执行的Warp,对其发送指令直到这个Warp又一次出现等待,周而复始。这就是上一篇所说“用多线程掩盖延迟”在硬件图景下的模样。

CPU_GPU_COMPARE

图 2.  GPU用多个Warp掩盖延迟 / 与CPU计算模式的对比

本图引用自PPT “CUDA Overview” from Cliff Woolley, NVIDIA.

如图,GPU用多个Warp快速切换来掩盖延迟,而CPU用快速的寄存器来减小延迟。两者的重要区别是寄存器数目,CPU的寄存器快但少,因此Context Switch代价高;GPU寄存器多而慢,但寄存器数量保证了线程Context Switch非常快。

多少线程才能够掩盖掉常见的延迟呢?对于GPU,最常见的延迟大概要数寄存器写后读依赖,即一个局域变量被赋值后接着不久又被读取,这时候会产生大约24个时钟周期的延迟。为了掩盖掉这个延迟,我们需要至少24个Warp轮流执行,一个Warp遇到延迟后的空闲时间里执行其余23个Warp,从而保持硬件的忙碌。在Compute Capability 2.0,SM中有32个CUDA核心,平均每周期发射一条指令的情况下,我们需要24*32 = 768个线程来掩盖延迟。
保持硬件忙碌,用CUDA的术语来说,就是保持充分的Occupancy,这是CUDA程序优化的一个重要指标。

(未完待续)

CUDA, 软件抽象的幻影背后

版权声明:原创作品,欢迎转载,但转载请以超链接形式注明文章来源(planckscale.info)、作者信息和本声明,否则将追究法律责任。

今天最酷炫的事情应该就是来自老黄的这条消息:1TFLOPS,P < 15W, ARM Cortex A57 * 4 + ARM Cortex A53 * 4 +  Maxwell 256 CUDA Cores,  Tegra X1.

tegrax1
图1.  Tegra X1

本想挖掘一下写篇博,但目前报道满天飞没太大必要了。于是又想起了这个命途多舛的话题:CUDA. 关于CUDA我写了两次,第一次不满意未发,第二次成文后保存失败灰飞烟灭在热力学第二定律决定的命运里。今天借X1的东风,我们再来聊聊CUDA.

**********************************************************

CUDA是个以性能为第一目标的语言,这也决定了CUDA开发者所要面对的复杂性远远要多于CUDA语言所抽象出来的编程模型本身。这大概会是软件抽象所要面对的永恒话题,我们可以去抽象出一组逻辑上漂亮完备的功能基元,却不能保证从性能的观点看它们同样也是小开销的基本操作。具体在CUDA里,最典型的例子是内存<->显存数据交换,一个简单的拷贝操作在性能上却是让人难以接受的,这背后是PCIE总线;对性能影响稍小些的例子比如Global Memory的读写需要考虑对齐,这是由于硬件层面warp和cache机制的体现;再者如过度臃肿的kernel或block过大导致寄存器耗尽,局域变量被吐到Local Memory导致的性能损失。

所有这些,都要求我们透过CUDA简洁干净的编程模型,看到软件抽象的美丽幻影背后那个不同的世界,它存在于抽象之下我们不熟悉的另一个层次,却透过性能这一个几乎是唯一的方式来影响着我们的软件。这颇类似万有引力与我们世界的关系:引力是唯一能透入额外维度的基本相互作用,如果世界有我们所不知道的维度存在,如何才能感受到那个世界对我们的影响?答案就是用引力。看过《星际穿越》的同学们想必对此有些印象。

在深入GPU的硬件架构之前,我们不妨先探讨一下这个问题:为什么GPU具有这么高的计算能力?我们试着归纳两条最主要的原因。

目前典型的计算模式有两种,CPU式的高速低延迟串行计算,和GPU式的高延迟高吞吐大规模并行计算。CPU是人们熟知的,它具有高速的内部寄存器和Cache,现代CPU又加入了多级流水线,猜测、乱序执行,超线程等技术加速其指令吞吐能力,具有快速的响应能力,但是对于大量数据的处理却相对不够用。这一点3D游戏应用就是典型的例子,当然,这就是GPU崛起的契机。
GPU天生为数据的批量处理而生,它擅长的是在大量数据上同时做同样或几乎一致(这点很重要)的计算。为什么要求一样的计算?这一点可以从很多个角度来回答。
最重要的一个回答是,多个线程同步执行一致的运算,使得我们可以用单路指令流对多个执行单元进行控制,大幅度减少了控制器的个数和系统的复杂度(设想成千上万的线程各自做不同的事情,如果再有线程间通讯/同步,将会是怎样的梦魇)。
另一方面,现实世界中应用在大规模数据上的计算,通常都涵盖在这一计算模式之中,因而考虑更复杂的模式本质上是不必要的。比如计算大气的流动,每一点的风速仅仅取决于该点邻域上的密度和压强分布;再如计算图像的卷积,每一个输出像素都仅是对应源点邻域和一个卷积核的内积。从这些例子中我们可以看到,除了各个数据单元上进行的计算是一样的,计算中数据之间的相互影响也具有某种“局域性”,一个数据单元上的计算最多需要它某个邻域上的数据。这一点意味着线程之间是弱耦合的,邻近线程之间会有一些共享数据(或者是计算结果),远距离的线程间则独立无关。
这个性质反映在CUDA里,就是Block划分的两重天地:Block内部具有Shared Memory,线程间可以共享数据、通讯和同步,Block外部则完全独立,Block间没有通讯机制,相互执行顺序不影响计算结果。这一划分使得我们既可以利用线程间通讯做一些复杂的应用和算法加速,又可以在Block的粒度上自由调度计算任务,在不同计算能力的硬件平台上自适应的调整任务安排。
现在我们把注意力放在“几乎一致”这里。最简单的并行计算方案是多路数据上同时进行完全一致的计算,即SIMD(单指令流多数据流)。这种方案是非常受限的。事实上我们可以看出,“完全一致”是不必要的。只要这些计算在大多数时候完全一致,就可以对它们做SIMD加速,而在计算分叉,各个线程不一致的特殊情况下,只需要分支内并行,分支间串行执行即可,毕竟这些只是很少出现的情况。这样,把“完全一致”这个限制稍微放松,就可以得到更广阔的应用范围和不输于SIMD的计算性能,即SIMT(单指令流多线程)的一个重要环节,这是GPU强大处理能力的第一个原因。

一个或许让每个初学者都惊讶的事实是这样一组数据:Global Memory访存延迟可以达到数百个时钟周期,即便是最快的Shared Memory和寄存器在有写后读依赖时也需要数十个时钟周期。这似乎和CUDA强大的处理能力完全相悖——如果连寄存器都这么慢,怎么会有高性能呢?难道这不会成为最大的瓶颈吗?
答案恰恰就出乎意料:不,这不是瓶颈,这个高延迟的开销被掩盖了,掩盖在大量线程之下。更清楚的说,当一组线程(同步执行,类似于SIMD的一个线程组,在CUDA里叫做warp)因为访存或其他原因出现等待时,就将其挂起,转而执行另一组线程,GPU的硬件体系允许同时有大量线程存活于GPU的SM(流多处理器)之中,控制单元在多组线程之间快速切换,从而保证资源的最大利用率——控制单元始终有指令可以发放,执行单元始终有任务可以执行,仍然可以保持最高的指令吞吐,每个单元基本都能保持充分的忙碌。
这就是GPU硬件设计中非常有特色的基本思想:用多线程掩盖延迟。这一设计区别于CPU的特点是,大量高延迟寄存器取代了少量低延迟寄存器,寄存器的数量保证了可以有大量线程同时存活,且可以在各组线程间快速切换。尽管每个线程是慢的,但庞大的线程数成就了GPU的数据吞吐能力。此为高性能的第二个原因。

这文又要写成未完待续了。接下来的日子,不填完旧坑不再开新话题。

Dronecode Project:无人机的开源平台

dc_slide_introdrone_video_slide

两个月前,Linux Foundation 联合 12 家技术公司推出了Dronecode Project,意在整合分散的无人机项目资源,从而创造一个共同共享的无人机开源平台,为无人机开发者们提供操作系统、导航工具、飞行控制等方面的大量开源代码。

Dronecode的创始成员:3D Robotics, Baidu, Box, DroneDeploy, Intel, jDrones, Laser Navigation, Qualcomm, SkyWard, Squadrone System, Walkera 和 Yuneec。

目前有超过1200名开发者参与Dronecode工作,某些项目每天有超过150个code commit。项目实例包括 APM/ArduPilot, Mission Planner, MAVLink 和 DroidPlanner。具有无人机前沿技术的很多组织包括Skycatch, DroneDeploy, HobbyKing, Horizon Ag, PrecisionHawk, Agribotix, and Walkera等都已采用了这个平台。

无人机这两年不可谓不热闹,大佬们似乎会在任何新兴领域里都插一脚:Amazon无人机Prime Air已迭代至第九代,网购随买随到的时代指日可待;Facebook和Google争相收购无人机制造公司:Facebook收购Ascenta;Google收购itan Aerospace,且在14年8月曝光了其研发了两年的无人机Project Wing;似乎GoPro也准备进军无人机市场。国内无人机老大DJI 大疆一边融资一边在美国发布航拍神器Inspire 1“悟”,性价比极高,秒杀N多航拍无人机;极飞科技融资2000万美元并与顺丰合作推动无人机送货;14年刚刚成立的Ehang推出了Ghost,在各众筹网站火了一把,话说其估值在6 个月内翻了 24 倍,创业者崭露头角的机会来了。

Dronecode 对无人机开发者来说是一大福音。无人机是我家老哥童年的梦想啊,等有精力了肯定要玩起来。把我们现在的多路视频实时全景拼接移植到Jetson TK1上做成嵌入式产品,并结合无人机拍摄出超炫全景视频。3D重建也是我们一直热衷准备做的,用无人机航拍城市,重建出城市的3D模型。虽有公司在做这些且已出产品,每次看着自己想做的事情被别人抢先做了,都恨自己分身无术啊,只能慢慢来。