澳门网络娱乐游戏平台-澳门电子游戏娱乐网址-官方直营

澳门官网赌场:介绍一篇不错的CUDA入门博客

澳门游戏平台娱乐登录 1

CUDA是什么

 

引用

CUDA,Compute Unified Device Architecture的简单称谓,是由NVIDIA集团开创的依附他们企业坐褥的图片微型机GPUs(Graphics Processing Units,能够发轫的知道为显卡)的三个并行计算平台和编制程序模型。

nBodyCS<I>学习笔记之总括着色器

出自Bookc的博客,链接在这里

透过CUDA,GPUs能够很方便地被用来张开通用总计(有一点像在CPU中开展的数值总结等等)。在一直不CUDA早先,GPUs日常只用来进行图片渲染(如通过OpenGL,DirectX)。

Nvidia-SDK(1) 

PS:那篇博客出作者的多个校友的手中,写的很好,是CUDA入门的好博客,由此才从她的博客中复制过来,大家也足以查看他博客中任何的随笔,涉猎很广,从语言到开源包。其它,他是做推荐系统的,如若有关于推荐系统的难点也能够在博客给他留言~!——————————————————正文——————————————————————鉴于本人的毕设必要利用GPUCUDA那项工夫,想找一本入门的课本,选拔了杰森Sanders等所著的书《CUDAByExampleanIntroductiontoGeneralPurposeGPUProgramming》。那本书作为入门教材,写的非常不利。自身感觉从知晓与回想的角度的出发,书中好些个故事情节都得以被省略掉,于是就有了那篇博文。此博文记录与计算此书的笔记和透亮。注意本文并不曾遵照书中章节的逐个来写。书中第8章图像互操作性和第11章多GPU系统上的CUDAC,这两章未有看。等不经常光了再看呢,赶紧码字。CUDA是什么CUDA,ComputeUnifiedDeviceArchitecture的简单的称呼,是由NVIDIA集团开创的依赖他们公司生产的图形微机GPUs的一个并行计算平台和编程模型。通过CUDA,GPUs能够很平价地被用来开展通用总括。在并未有CUDA早前,GPUs平常只用来扩充图片渲染。开拓职员能够通过调用CUDA的API,来进行人机联作编制程序,达到高质量总结目的。NVIDIA公司为了抓住更加多的开垦职员,对CUDA进行了编程语言扩张,如CUDAC/C++,CUDAFortran语言。注意CUDAC/C++能够看做二个新的编制程序语言,因为NVIDIA配置了相应的编写翻译器nvcc,CUDAFortran同样。越来越多消息方可参照他事他说加以侦查文献。64位Ubuntu12.04安装CUDA5.5具体步骤请点击这里。[b]对CUDAC的民用懵懂以为[/b]如果强行的认为C语言专门的学问的靶子是CPU和内部存储器条,那么CUDAC工作的的对象正是GPU及GPU上的内部存款和储蓄器,且丰裕利用了GPU多核的优势及收缩了互相编制程序的难度。平常经过C语言把数量从外边读入,再分配数据,给CUDAC,以便在GPU上测算,然后再把总结结果回到给C语言,以便进一层职业,如进一层管理及显示,或另行此进度。要害概念与名称主机将CPU及系统的内存称为主机。设备将GPU及GPU本身的呈现内部存款和储蓄器称为器械。线程(Thread)日常通过GPU的贰个核查行处理。。线程块(Block)1.由五个线程组成。2.各block是并行履行的,block间不恐怕通信,也一贯不举行顺序。3.注意线程块的数量节制为不超过65535。线程格(Grid)由多个线程块组成。线程束在CUDA布局中,线程束是指三个分包三十一个线程的集中,那个线程集合被“编织在一块”并且“志趣相投”的款式进行。在前后相继中的每一行,线程束中的各样线程都将在差异数量上实行同一的下令。核函数1.在GPU上奉行的函数日常称为核函数。2.相仿经过标记符__global__修饰,调用通过参数1,参数2,用于申明内核函数中的线程数量,以至线程是哪些组织的。3.以线程格的格局协会,每一个线程格由若干个线程块组成,而各类线程块又由若干个线程组成。4.是以block为单位进行的。5.叧能在主机端代码中调用。6.调用时必须注明内核函数的实行参数。7.在编制程序时,必得先为kernel函数中用到的数组或变量分配好充裕的半空中,再调用kernel函数,不然在GPU总计时会发生错误,比方越界或报错,以至形成蓝屏和死机。

开荒职员能够通过调用CUDA的API,来扩充人机联作编制程序,达到高品质总括指标。NVIDIA公司为了吸引越来越多的开拓人士,对CUDA进行了编制程序语言增加,如CUDA C/C++,CUDA Fortran语言。注意CUDA C/C++能够当作三个新的编程语言,因为NVIDIA配置了对应的编写翻译器nvcc,CUDA Fortran同样。越来越多音信能够仿效文献。

版权评释:本文为博主原创小说,未经博主允许不得转发。

/**@file_nameHelloWorld.cu后缀名称.cu*/#includestdio.h#includecuda_runtime.h//头文件//核函数声明,前面的关键字__global____global__voidkernel(void){}intmain(void){//核函数的调用,注意1,1,第一个1,代表线程格里只有一个线程块;第二个1,代表一个线程块里只有一个线程。kernel1,1();printf("Hello,World!n");return0;}

非常重要概念与名称

 

dim3构造类型1.dim3是基亍uint3定义的矢量类型,卓殊亍由3个unsignedint型组成的构造体。uint3连串有多个数据成员unsignedintx;unsignedinty;unsignedintz;2.可采纳亍一维、二维或三维的目录来标记线程,构成一维、二维或三个维度线程块。3.dim3构造类型变量用在核函数调用的,中。4.有关的多少个放置变量4.1.threadIdx,以文害辞取得线程thread的ID索引;假使线程是一维的那么就取threadIdx.x,二维的还足以多取到八个值threadIdx.y,由此及彼到三个维度threadIdx.z。4.2.blockIdx,线程块的ID索引;同样有blockIdx.x,blockIdx.y,blockIdx.z。4.3.blockDim,线程块的维度,相似有blockDim.x,blockDim.y,blockDim.z。4.4.gridDim,线程格的维度,相同有gridDim.x,gridDim.y,gridDim.z。5.对于一维的block,线程的threadID=threadIdx.x。6.对此大小为的二维block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x。7.对此大小为的三个维度block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。8.对此计算线程索引偏移增量为已开发银行线程的总额。如stride=blockDim.x*gridDim.x;threadId+=stride。函数修饰符1.__global__,注明被修饰的函数在设施上举行,但在主机上调用。2.__device__,表明被修饰的函数在道具上施行,但一定要在任何__device__函数可能__global__函数中调用。常用的GPU内部存款和储蓄器函数cudaMalloc(卡塔尔(قطر‎1.函数原型:cudaError_澳门游戏平台娱乐登录,tcudaMalloc(void**devPtr,size_澳门官网赌场,tsize卡塔尔。2.函数用项:与C语言中的malloc函数同样,只是此函数在GPU的内部存款和储蓄器你分配内部存款和储蓄器。3.注意事项:3.1.能够将cudaMalloc(卡塔尔国分配的指针传递给在设施上推行的函数;3.2.足以在设备代码中利用cudaMalloc(卡塔尔(قطر‎分配的指针进行配备内部存款和储蓄器读写操作;3.3.能够将cudaMalloc(卡塔尔分配的指针传递给在主机上试行的函数;3.4.不能在主机代码中接受cudaMalloc(State of Qatar分配的指针实行主机内部存款和储蓄器读写操作。cudaMemcpy()1.函数原型:cudaError_tcudaMemcpy(void*dst,constvoid*src,size_tcount,cudaMemcpyKindkindState of Qatar。2.函数效用:与c语言中的memcpy函数相同,只是此函数能够在主机内部存款和储蓄器和GPU内存之间人机联作拷贝数据。3.函数参数:cudaMemcpyKindkind表示数据拷贝方向,假设kind赋值为cudaMemcpyDeviceToHost表示数据从设备内部存款和储蓄器拷贝到主机内部存款和储蓄器。4.与C中的memcpy(卡塔尔国相似,以贰只方式实践,即当函数再次来到时,复制操作就曾经成功了,并且在输出缓冲区中隐含了复制进去的内容。5.对应的有个异步方式推行的函数cudaMemcpyAsync(State of Qatar,这么些函数详细解释请看下边包车型大巴流一节有关内容。cudaFree()1.函数原型:cudaError_tcudaFree(void*devPtr卡塔尔(قطر‎。2.函数职能:与c语言中的free(卡塔尔(قطر‎函数相像,只是此函数释放的是cudaMalloc(卡塔尔分配的内部存储器。上边实例用于解释上面八个函数

主机

DirectX向来是Windows上航海用体育场所片和游玩开辟的大旨技艺。DirectX提供了一种在显卡上运营的前后相继——着色器(Shader)。从DirectX11最初,DirectX扩展了一种计算着色器(Compute Shader),它是特意为与图片非亲非故的通用总结设计的。因而DirectX就改成了一个通用GPU计算的平台。鉴于GPU具备特别强盛的交互作用运算技术,学习运用DirectCompute是很有含义的。基本上,DirectCompute供给通过测算着色器5.0(Compute Shader)编制程序模型(即CS 5.0)技艺完全完毕。但是CS 5.0急需DirectX 11硬件本事支撑,本文私下认可机器辅助DirectX11硬件的。

#includestdio.h#includecuda_runtime.h__global__voidadd(inta,intb,int*c){*c=a+b;}intmain(void){intc;int*dev_c;//cudaMalloc()cudaMalloc((void**)dev_c,sizeof(int));//核函数执行add1,1(2,7,dev_c);//cudaMemcpy()cudaMemcpy(c,dev_c,sizeof(int),cudaMemcpyDeviceToHost);printf("2+7=%dn",c);//cudaFree()cudaFree(dev_c);return0;}

将CPU及系统的内部存款和储蓄器(内部存款和储蓄器条)称为主机。

 

GPU内部存款和储蓄器分类全局内部存款和储蓄器领头意义上的配备内部存款和储蓄器。共享内存1.职责:设备内部存款和储蓄器。2.花样:关键字__shared__拉长到变量表明中。如__shared__floatcache[10]。3.指标:对于GPU上运维的每种线程块,CUDAC编译器都将制造该分享变量的叁个副本。线程块中的每一个线程都分享那块内部存款和储蓄器,但线程却敬敏不谢看出也无法改善其余线程块的变量别本。那样使得叁个线程块中的四个线程可以在总计上通讯和合营。常量内部存款和储蓄器1.地方:设备内部存款和储蓄器2.款式:关键字__constant__丰裕到变量证明中。如__constant__floats[10];。3.指标:为了提高品质。常量内部存款和储蓄器接纳了分歧于规范全局内部存款和储蓄器的管理情势。在少数情状下,用常量内部存款和储蓄器替换全局内部存储器能有效地减小内部存款和储蓄器带宽。4.特征:常量内部存款和储蓄器用于保存在核函数施行时期不会爆发变化的数码。变量的拜谒限定为只读。NVIDIA硬件提供了64KB的常量内部存款和储蓄器。不再须求cudaMalloc(卡塔尔也许cudaFree(卡塔尔(قطر‎,而是在编写翻译时,静态地分配空间。5.渴求:当我们供给拷贝数据到常量内存中应该利用cudaMemcpyToSymbol(卡塔尔(قطر‎,而cudaMemcpy(卡塔尔国会复制到全局内部存储器。6.品质晋升的自始至终的经过:6.1.对常量内部存款和储蓄器的单次读操作能够播放到别的的“相近”线程。那将节约十七回读取操作。6.2.常量内部存款和储蓄器的数据将缓存起来,由此对同样地方的接连读操作将不会时有发生额外的内部存款和储蓄器通讯量。纹理内部存款和储蓄器1.职位:设备内部存款和储蓄器2.目标:能够裁减对内部存款和储蓄器的央浼并提供高速的内部存储器带宽。是极度为那些在内存访谈格局中存在大气空中局地性的图形应用程序设计,意味着一个线程读取之处只怕与将近线程读取的岗位“极度周围”。如下图:3.纹理变量必须注脚为文件功用域内的全局变量。4.款式:分为一维纹理内部存款和储蓄器和二维纹理内部存款和储蓄器。4.1.一维纹理内部存款和储蓄器4.1.1.用texture类型类型注脚,如texturefloattexIn。4.1.2.由此cudaBindTexture(卡塔尔(قطر‎绑定到纹理内部存款和储蓄器中。4.1.3.通过tex1Dfetch(卡塔尔(قطر‎来读取纹理内部存款和储蓄器中的数码。4.1.4.经过cudaUnbindTexture(卡塔尔(قطر‎撤销绑定纹理内部存储器。4.2.二维纹理内部存款和储蓄器4.2.1.用texture类型,数字类型申明,如texturefloat,2texIn。4.2.2.通过cudaBindTexture2D(卡塔尔绑定到纹理内存中。4.2.3.经过tex2D(卡塔尔(قطر‎来读取纹理内部存款和储蓄器中的数量。4.2.4.透过cudaUnbindTexture(卡塔尔国撤废绑定纹理内部存款和储蓄器。原则性内部存储器1.地方:主机内部存储器。2.定义:也堪称页锁定内部存款和储蓄器还是不可分页内部存款和储蓄器,操作系统将不会对这块内部存款和储蓄器分页并交流到磁盘上,进而确定保障了该内部存款和储蓄器始终驻留在物理内部存款和储蓄器中。由此操作系统能够平安地使有些应用程序访问该内部存款和储蓄器的概略地址,因为那块内部存款和储蓄器将不会毁掉只怕再度定位。3.目的:升高访问速度。由于GPU知道主机内部存款和储蓄器的概略地址,由此得以因此“间接内部存款和储蓄器访谈DMA7.流同步:通过cudaStreamSynchronize(卡塔尔来和睦。8.流销毁:在退出应用程序此前,必要销毁对GPU操作举行排队的流,调用cudaStreamDestroy(卡塔尔国。9.针对八个流:9.1.记得对流进行同步操作。9.2.将操作放入流的系列时,应运用肥瘦优先进轨范式,而非深度优先的办法,换句话说,不是率先添加第0个流的有所操作,再相继增进前面包车型客车第1,2,…个流。而是轮换实行增加,举个例子将a的复制操作增添到第0个流中,接着把a的复制操作增加到第三个流中,再持续别的的近乎更改增加的作为。9.3.要结实记住操作放入流中的类别中的顺序影响到CUDA驱动程序调治那个操作和流以至推行的法子。技巧1.当线程块的数目为GPU中拍卖多少的2倍时,将达到最优质量。2.核函数实行的第叁个总计就是简政放权输入数据的舞狮。各种线程的开端偏移都是0到线程数量减1之间的某些值。然后,对偏移的增量为已运行线程的总量。实例程序感兴趣的读者能够下载本书附带的示范代码点击这里下载_by_example.zip。

设备

澳门游戏平台娱乐登录 2

将GPU及GPU自己的展现内部存款和储蓄器称为设备。

 

线程(Thread)

 

诚如经过GPU的二个核举办拍卖。(能够象征成一维,二维,三个维度,具体上面再细说)。

1.测算着色器创造步骤 1.1初步化设备和上下文

线程块(Block)

1.2从HLSL文件加载着色器程序并编写翻译

1. 由三个线程组成(可以代表成一维,二维,三个维度,具体下边再细说)。

1.3为着色器创立并带头化财富(如缓冲区)

2. 各block是并行实行的,block间不恐怕通讯,也一贯不实行各种。

1.4设定着色器状态,并实行

澳门官网赌场:介绍一篇不错的CUDA入门博客。3. 在乎线程块的数额限定为不超过65535(硬件限定)。

1.5取回运算结果

线程格(Grid)

详细的制造进程已经在高斯模糊一篇详细介绍,本文大约简单门船演说下:

由多少个线程块组成(可以象征成一维,二维,三个维度,具体下边再细说)。

D3D_FEATURE_LEVEL levelsWanted[] =

{

    D3D_FEATURE_LEVEL_11_0,

    D3D_FEATURE_LEVEL_10_1,

    D3D_FEATURE_LEVEL_10_0

};

UINT numLevelsWanted = sizeof( levelsWanted ) /sizeof( levelsWanted[0] );

 

D3D_DRIVER_TYPE driverTypes[] =

{

    D3D_DRIVER_TYPE_REFERENCE,

    D3D_DRIVER_TYPE_HARDWARE,

};

UINT numDriverTypes = sizeof( driverTypes ) /sizeof( driverTypes[0] );

 

// 遍历每一种驱动类型,先尝试参考驱动,然后是硬件驱动

// 成功创建一种之后就退出循环。

// 你可以更改以上顺序来尝试各种配置

// 这里我们只需要参考设备来演示API调用

for( UINT driverTypeIndex = 0; driverTypeIndex < numDriverTypes; driverTypeIndex++ )

{

    D3D_DRIVER_TYPE g_driverType = driverTypes[driverTypeIndex];

    UINT createDeviceFlags = NULL;

    hr = D3D11CreateDevice( NULL, g_driverType, NULL, createDeviceFlags,

        levelsWanted, numLevelsWanted, D3D11_SDK_VERSION,

        &g_pD3DDevice, &g_D3DFeatureLevel, &g_pD3DContext );

}

线程束

立业成家运营后,这段代码将时有产生三个设施指针,多少个上下文指针还或者有贰个性情级其余Flag。

在CUDA布局中,线程束是指多少个包罗31个线程的聚合,那么些线程集合被“编织在一同”並且“同心同德”的款型进行。在程序中的每一行,线程束中的各类线程都将要分化数量上举行相通的一声令下。

在乎:为轻易起见以上代码省略了大多变量评释的代码。完整示例代码须求补上那么些代码。这里的代码片段仅用来体现程序中发出的作业。

核函数(Kernel)

接纳要用的显卡

1. 在GPU上实行的函数经常称为核函数。

用IDXGIFactory对象就可以枚举系统中安装的显卡,如上面代码所示。首先创立多少个IDXGIFactory对象,然后调用EnumAdapters并传到多个意味着正在枚举显卡的整数。若是荒诞不经,它会回来DXGI_ERROR_NOT_FOUND。

2. 相通经过标志符__global__修饰,调用通过<<<参数1,参数2>>>,用于评释内核函数中的线程数量,甚至线程是如何组织的。

// 获取所有安装的显卡

std::vector<IDXGIAdapter1*> vAdapters;

IDXGIFactory1* factory;

CreateDXGIFactory1(__uuidof(IDXGIFactory1), (void**)&factory);

IDXGIAdapter1 * pAdapter = 0;

UINT i=0;

while(factory->EnumAdapters1(i, &pAdapter) != DXGI_ERROR_NOT_FOUND)

{

    vAdapters.push_back(pAdapter);

    ++i;

}

3. 以线程格(Grid)的样式组织,种种线程格由若干个线程块(block)组成,而各种线程块又由若干个线程(thread)组成。

接下去,在调用D3DCreateDevice创立设备的时候从第二个参数字传送入想用的显卡适配器指针,而且将使得类型设为D3D_DRIVER_TYPE_UNKNOWN。详细音信请参见D3D11文书档案中D3DCreateDevice函数的援救。

4. 是以block为单位施行的。

g_driverType = D3D_DRIVER_TYPE_UNKNOWN;

hr = D3D11CreateDevice( vAdapters[devNum], g_driverType, NULL, createDeviceFlags, levelsWanted,

            numLevelsWanted, D3D11_SDK_VERSION, &g_pD3DDevice, &g_D3DFeatureLevel, &g_pD3DContext );

5. 叧能在主机端代码中调用。

运转总计着色器

6. 调用时必得注解内核函数的实践参数。

着色器(Shader)是在显卡上运营的程序,它并差异于CPU上施行的顺序,能够用HLSL来编排(见后文)。

7. 在编制程序时,必需先为kernel函数中用到的数组或变量分配好丰硕的上空,再调用kernel函数,不然在GPU计算时会产生错误,比方越界或报错,以至招致蓝屏和死机。

DirectCompute程序中的总计着色器是透过Dispatch函数推行的:

dim3布局类型

// 现在分派(运行) 计算着色器, 分成16x16个线程组。

g_pD3DContext->Dispatch( 16, 16, 1 );

1. dim3是基亍uint3定义的矢量类型,万分亍由3个unsigned int型组成的构造体。uint3品种有七个数据成员unsigned int x; unsigned int y; unsigned int z;

如上语句分派了16x拾陆个线程组。

2. 可接受亍一维、二维或三个维度的目录来标记线程,构成一维、二维或三个维度线程块。

留意,着色器的输入视而不见思索成“状态”。便是说你应当在分摊着色器程序早前设定意况,而只要分派了,“状态”决定输入变量的值。所以着色器分派代码日常应该像这么:

3. dim3结构类型变量用在核函数调用的<<<,>>>中。

pd3dImmediateContext->CSSetShader( ... );

pd3dImmediateContext->CSSetConstantBuffers( ...);

pd3dImmediateContext->CSSetShaderResources( ...);  // CS 输入

// CS 输出

pd3dImmediateContext->CSSetUnorderedAccessViews( ...);

// 运行 CS

pd3dImmediateContext->Dispatch( dimx, dimy, 1 );

4. 有关的多少个放置变量

以上全体常量缓冲(constant buffer),缓冲等得以在着色器程序中来看东东都是在分摊线程在此之前经过调用CSSet…(State of Qatar设定的

4.1. threadIdx,一面之识得到线程thread的ID索引;借使线程是一维的那么就取threadIdx.x,二维的还足以多取到多少个值threadIdx.y,依此类推到三个维度threadIdx.z。


与CPU进行联合

4.2. blockIdx,线程块的ID索引;相符有blockIdx.x,blockIdx.y,blockIdx.z。

请在意上边装有的调用都以异步的。CPU方面总是会立时回到然后工夫体实施。假设有须求,其后调用的缓冲区“映射”操作(详见下文缓冲区部分)时CPU的调用线程才会停下来等待全数异步操作的产生。

4.3. blockDim,线程块的维度,相同有blockDim.x,blockDim.y,blockDim.z。

事件:基本剖判和同步操作

4.4. gridDim,线程格的维度,同样有gridDim.x,gridDim.y,gridDim.z。

DirectCompute提供一种基于“查询”的事件机制API。你能够创设、插入并等候特定情景的询问来剖断着色器(或任何异步调用)具体在曾几何时实践。上边包车型大巴例子创造了三个询问,然后通过等待查询来确定保障运维到某一点时持有该施行的操作都已经实行了,再分摊着色器,最终等待另叁个询问并确认着色器程序已进行完结。

5. 对于一维的block,线程的threadID=threadIdx.x。

创立查询对象:

6. 对于大小为(blockDim.x, blockDim.y)的 二维 block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x。

D3D11_QUERY_DESC pQueryDesc;

pQueryDesc.Query = D3D11_QUERY_EVENT;

pQueryDesc.MiscFlags = 0;

ID3D11Query *pEventQuery;

g_pD3DDevice->CreateQuery( &pQueryDesc, &pEventQuery );

7. 对于大小为(blockDim.x, blockDim.y, blockDim.z)的 三维 block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。

接下来在一多元调用中插入“篱笆”,再伺机之。假使查询的新闻不设有,GetData(卡塔尔(قطر‎将再次来到S_FALSE。

8. 对此总括线程索引偏移增量为已运转线程的总和。如stride = blockDim.x * gridDim.x; threadId += stride。

g_pD3DContext->End(pEventQuery); // 在 pushbuffer 中插入一个篱笆

 while( g_pD3DContext->GetData( pEventQuery, NULL, 0, 0 ) == S_FALSE ) {}//自旋等待事件结束

 

 g_pD3DContext->Dispatch(,x,y,1);//启动着色器

 

 g_pD3DContext->End(pEventQuery);//在 pushbuffer 中插入一个篱笆

 while( g_pD3DContext->GetData( pEventQuery, NULL, 0, 0 ) == S_FALSE ) {}//自旋等待事件结束

函数修饰符

末段用那条语句释放查询对象:

1. __global__,评释被修饰的函数在设备上施行,但在主机上调用。

pEventQuery->Release();

2. __device__,声明被修饰的函数在设施上实施,但只好在其他__device__函数也许__global__函数中调用。

请小心创立和刑释查询对象避防弄出太多的询问来(非常是您管理一帧镜头的时候)。

常用的GPU内部存款和储蓄器函数

点击这里查看MSDN关于查询的帮手.aspx)

cudaMalloc()

 

1. 函数原型: cudaError_t cudaMalloc (void **devPtr, size_t size)。

DirectCompute中的能源

2. 函数用场:与C语言中的malloc函数雷同,只是此函数在GPU的内部存储器你分配内部存款和储蓄器。

译注:能源是指能够被GPU或CPU访谈的数额,是着色器的输入与输出。包括缓冲区和纹理等体系。

3. 注意事项:

DirectX中能源是遵守以下步骤创制出来的:

3.1. 足以将cudaMalloc(卡塔尔国分配的指针传递给在装置上实践的函数;

1.先是创设二个财富描述器,用来描述所要创造的能源。财富描述器是一种内含非常多Flag和所需财富音讯的布局体。

3.2. 得以在装置代码中动用cudaMalloc(卡塔尔(قطر‎分配的指针进行设备内部存款和储蓄器读写操作;

2.调用某种Create系方法,传入描述器作为参数并创制能源。

3.3. 足以将cudaMalloc(State of Qatar分配的指针传递给在主机上进行的函数;

CPU与GPU之间的报道

3.4. 不可能在主机代码中采纳cudaMalloc(卡塔尔国分配的指针进行主机内部存款和储蓄器读写操作(即不能开展解引用)。

gD3DContext->CopyResouce(卡塔尔函数能够用来读取或复制财富。这里复制是指五个财富之间的复制。假使要在CPU和GPU之间(译注:便是在内部存款和储蓄器和显存之间)举行理并答复制的话,先要创建多少个CPU那边的“中间转播”能源。中间转播财富得以映射到CPU的内部存款和储蓄器指针上,那样就足以从当中间转播能源中读取数据或许复制数据。之后废除中间转播财富的照射,再用CopyResource(卡塔尔方法开展与GPU之间的复制。

cudaMemcpy()

CPU与GPU之间缓冲区复制的属性

1. 函数原型:cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)。

CUDA-C语言(CUDA是nVidia的GPU通用总计平台)能够分配定址(pinned)宿主指针和写入联合(write combined)宿主指针,通过它们能够开展质量最棒的GPU数据复制。而在DirectCompute中,缓冲区的“usage”属性决定了内存分配的花色和做客时的习性。

2. 函数成效:与c语言中的memcpy函数雷同,只是此函数能够在主机内部存款和储蓄器和GPU内部存款和储蓄器之间人机联作拷贝数据。

·        D3D11_USAGE_STAGING 这种usage的财富是系统内存,能够直接由GPU进行读写。可是她们仅能用作复制操作(CopyResource(卡塔尔,CopySubresourceRegion(State of Qatar)的源或指标,而不可能一向在着色器中选择。

3. 函数参数:cudaMemcpyKind kind表示数据拷贝方向,若是kind赋值为cudaMemcpyDeviceToHost表示数据从设备内部存款和储蓄器拷贝到主机内部存款和储蓄器。

·        借使财富创制的时候钦赐了D3D11_CPU_ACCESS_WPAJEROITE flag那么从CPU到GPU复制的质量最棒。

4. 与C中的memcpy(卡塔尔相似,以合作方式实行,即当函数重回时,复制操作就曾经达成了,何况在出口缓冲区中包涵了复制进去的原委。

·        假诺用了D3D11_CPU_ACCESS_READ该能源将是一个由CPU缓存的能源,质量超级低(可是帮衬取回操作)

5. 相应的有个异步情势进行的函数cudaMemcpyAsync(卡塔尔(قطر‎,那些函数详整请看上面包车型大巴流一节有关内容。

·        如若相同的时间钦定,READ比WEnclaveITE优先。

cudaFree()

·        D3D11_USAGE_DYNAMIC (仅能用来缓冲区型财富,不能够用于纹理财富)用于火速的CPU->GPU内部存款和储蓄器传输。这种财富不但能够看作复制和源和目的,还足以看做纹理(用D3D的术语说,叫做着色器能源视图ShaderResourceView)在着色器中读取。然则着色器不能够写入这种能源。那几个财富的版本由驱动程序来调控,每一遍你用DISCARubiconD flag映射内部存款和储蓄器的时候,假诺那块内部存储器还在被GPU所采取,驱动程序就能够发出一块新的内部存款和储蓄器来,而不会等GPU的操作截止。它的意义在于提供一种流的点子将数据输送到GPU。

1. 函数原型:cudaError_t cudaFree ( void* devPtr )。

二.示例代码《nBodyCS》 //-----------------------------------------------------------------------------

2. 函数功效:与c语言中的free(卡塔尔(قطر‎函数同样,只是此函数释放的是cudaMalloc(卡塔尔(قطر‎分配的内部存储器。

// Reset the body system to its initial configuration

上面实例用于解释上面三个函数

//-----------------------------------------------------------------------------

GPU内部存款和储蓄器分类

HRESULT NBodySystemCS::resetBodies(BodyDataconfigData)

大局内部存款和储蓄器

{

通俗意义上的装备内存。

    HRESULThr =S_OK;

分享内部存款和储蓄器

      m_numBodies= configData.nBodies;

1. 地点:设备内部存款和储蓄器。

      //for compute shader on CS_4_0, we can only have a single UAV per shader, so wehave to store particle

2. 形式:关键字__shared__增进到变量注解中。如__shared__ float cache[10]。

      //position and velocity in the same array: all positions followed by all velocities

3. 目标:对于GPU上运维的各样线程块,CUDA C编写翻译器都将创造该分享变量的二个别本。线程块中的每种线程都分享那块内部存款和储蓄器,但线程却爱莫能助看出也不能够改改别的线程块的变量别本。这样使得叁个线程块中的七个线程能够在计算上通讯和协作。

      D3DXVECTOR4*particleArray =newD3DXVECTOR4[m_numBodies* 3];

常量内部存款和储蓄器

      for(unsignedinti=0; i < m_numBodies; i++){

1. 岗位:设备内部存款和储蓄器

           particleArray[i] =D3DXVECTOR4(configData.position[i*3 + 0],

2. 形式:关键字__constant__拉长到变量注明中。如__constant__ float s[10];。

                                                      configData.position[i*3 +1],

3. 目标:为了进步品质。常量内部存款和储蓄器接收了区别于标准全局内存的处理方式。在好几情状下,用常量内存替换全局内部存款和储蓄器能有效地回退内部存款和储蓄器带宽。

                                                     configData.position[i*3 +2],

4. 特点:常量内部存款和储蓄器用于保存在核函数实施时期不会产生变化的数目。变量的访谈约束为只读。NVIDIA硬件提供了64KB的常量内部存款和储蓄器。不再需求cudaMalloc(卡塔尔(قطر‎只怕cudaFree(State of Qatar,而是在编写翻译时,静态地分配空间。

                                                     1.0);

5. 渴求:当我们要求拷贝数据到常量内部存款和储蓄器中应该使用cudaMemcpyToSymbol(State of Qatar,而cudaMemcpy(卡塔尔国会复制到全局内部存款和储蓄器。

           particleArray[i +m_numBodies]=particleArray[i];

6. 性质进步的原由:

        particleArray[i + 2 *m_numBodies]=D3DXVECTOR4(configData.velocity[i*3 +0],

6.1. 对常量内部存款和储蓄器的单次读操作能够播放到别的的“附近”线程。这将节约16次读取操作。(为何是15,因为“周边”指半个线程束,二个线程束包括三十几个线程的会合。)

                                                           configData.velocity[i*3

6.2. 常量内部存储器的数量将缓存起来,因此对相同地方的三番若干次读操作将不会生出额外的内部存储器通信量。

  • 1],

纹理内部存款和储蓄器

                                                                              configData.velocity[i*3 +2],

1. 地点:设备内存

                                                                               1.0);

2. 目标:能够降低对内部存储器的伸手并提供快捷的内部存款和储蓄器带宽。是特意为那个在内部存款和储蓄器访谈情势中设有大量上空局部性的图样应用程序设计,意味着二个线程读取的岗位可能与周围线程读取的岗位“非常相仿”。如下图:

      }

  1. 纹理变量(引用)必得证明为文件成效域内的全局变量。

//------------------------------------------------------------------------------------------------------

4. 花样:分为一维纹理内部存款和储蓄器 和 二维纹理内部存储器。

构造化缓冲区和乱序访问视图

4.1. 一维纹理内部存款和储蓄器

ComputeShader的四个很关键的特色是构造化缓冲区和乱序访问视图。布局化缓冲区(structuredbuffer)在总结着色器中得以像数组同样访谈。大肆线程能够读写猖狂地方(即并路程序的分发scatter和访问gather动作)。乱序访谈视图(unordered access view,UAV)是一种将调用方创设的财富绑定到着色器中的机制,并且同意……乱序访问。

4.1.1. 用texture<类型>类型申明,如texture texIn。

扬言布局化缓冲区

4.1.2. 通过cudaBindTexture(State of Qatar绑定到纹理内部存款和储蓄器中。

咱俩得以用D3D11_RESOURCE_MISC_BUFFER_STRUCTURED来成立构造化缓冲区。上面钦赐的绑定flag表示同意着色器乱序访谈。下面接纳的私下认可usage代表它能够被GPU进行读写,但必要复制到中间转播能源个中手艺被CPU读写。

4.1.3. 透过tex1Dfetch(State of Qatar来读取纹理内部存款和储蓄器中的数量。

//------------------------------------------------------------------------------------------------------

4.1.4. 透过cudaUnbindTexture(卡塔尔国撤消绑定纹理内部存款和储蓄器。

      D3D11_SUBRESOURCE_DATAinitData = {particleArray,0, 0 };

本文由澳门网络娱乐游戏平台发布于操作系统,转载请注明出处:澳门官网赌场:介绍一篇不错的CUDA入门博客

相关阅读