ACCESSnBodyCS学习笔记之计算着色器

 

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

Nvidia-SDK(1) 

版权申明:本文为博主原创作品,未经博主允许不得转发。

 

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硬件的。

 

ACCESS 1

 

 

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

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

1.3为着色器创设并最先化财富(如缓冲区)

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

1.5取回运算结果

详细的创设进度已经在高斯模糊一篇详细介绍,本文大致简单门船演说下:

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。

只顾:为简易起见以上代码省略了众多变量表明的代码。完整示例代码供给补上那些代码。那里的代码片段仅用来呈现程序中生出的事体。

采用要用的显卡

用IDXGIFactory对象即可枚举系统中设置的显卡,如下边代码所示。首先成立三个IDXGIFactory对象,然后调用EnumAdapters并传到一个意味着正在枚举显卡的整数。倘若不存在,它会回来DXGI_ERROR_NOT_FOUND。

// 获取所有安装的显卡

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;

}

接下去,在调用D3DCreateDevice创设设备的时候从第②个参数传入想用的显卡适配器指针,并且将使得类型设为D3D_DRIVER_TYPE_UNKNOWN。详细音讯请参见D3D11文书档案中D3DCreateDevice函数的帮忙。

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 );

运转计算着色器

着色器(Shader)是在显卡上运维的先后,它并不一样于CPU上实施的主次,能够用HLSL来编排(见后文)。

DirectCompute程序中的总结着色器是由此Dispatch函数执行的:

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

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

上述语句分派了16x十六个线程组。

瞩目,着色器的输入常见考虑成“状态”。就是说你应当在分摊着色器程序在此以前设定情形,而只要分派了,“状态”决定输入变量的值。所以着色器分派代码日常应该像这么:

pd3dImmediateContext->CSSetShader( … );

pd3dImmediateContext->CSSetConstantBuffers( …);

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

// CS 输出

pd3dImmediateContext->CSSetUnorderedAccessViews( …);

// 运行 CS

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

如上全体常量缓冲(constant
buffer),缓冲等足以在着色器程序中阅览东东都以在分摊线程以前经过调用CSSet…()设定的


与CPU进行协同

请小心下边装有的调用都以异步的。CPU方面总是会应声赶回然后才具体实践。如果有必不可少,其后调用的缓冲区“映射”操作(详见下文缓冲区部分)时CPU的调用线程才会停下来等待全数异步操作的完毕。

事件:基本分析和同步操作

DirectCompute提供一种基于“查询”的事件机制API。你能够创立、插入并听候特定情景的查询来判定着色器(或任何异步调用)具体在何时实施。上面的事例创立了三个询问,然后经过等待查询来担保运转到某一点时享有该执行的操作都曾经执行了,再分摊着色器,最终等待另多个查询并认同着色器程序已施行完成。

创办查询对象:

D3D11_QUERY_DESC pQueryDesc;

pQueryDesc.Query = D3D11_QUERY_EVENT;

pQueryDesc.MiscFlags = 0;

ID3D11Query *pEventQuery;

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

然后在一多元调用中插入“篱笆”,再伺机之。要是查询的音信不设有,GetData()将重临S_FALSE。

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 ) {}//自旋等待事件结束

最后用那条语句释放查询对象:

pEventQuery->Release();

请小心成立和自由查询对象以防弄出太多的查询来(尤其是你处理一帧镜头的时候)。

点击那里查看MSDN关于查询的助手.aspx)

 

DirectCompute中的能源

译注:财富是指能够被GPU或CPU访问的数码,是着色器的输入与出口。包含缓冲区和纹理等门类。

DirectX中能源是比照以下步骤创制出来的:

1.率先创立一个财富描述器,用来叙述所要成立的财富。能源描述器是一种内含许多Flag和所需财富新闻的结构体。

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

CPU与GPU之间的通信

gD3DContext->CopyResouce()函数能够用来读取或复制财富。那里复制是指七个能源之间的复制。如若要在CPU和GPU之间(译注:就是在内部存款和储蓄器和显存之间)实行理并答复制的话,先要创造三个CPU那边的“中间转播”财富。中间转播能源可以映射到CPU的内部存款和储蓄器指针上,那样就足以从中间转播能源中读取数据可能复制数据。之后解除中间转播财富的映照,再用CopyResource()方法举行与GPU之间的复制。

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

CUDA-C语言(CUDA是nVidia的GPU通用总计平台)可以分配定址(pinned)宿主指针和写入联合(write
combined)宿主指针,通过它们能够进行质量最好的GPU数据复制。而在DirectCompute中,缓冲区的“usage”属性决定了内部存款和储蓄器分配的门类和走访时的习性。

·        D3D11_USAGE_STAGING
那种usage的财富是系统内部存款和储蓄器,能够向来由GPU实行读写。然则他们仅能用作复制操作(CopyResource(),CopySubresourceRegion())的源或目的,而不能一贯在着色器中使用。

·        假诺能源创制的时候钦命了D3D11_CPU_ACCESS_W奥德赛ITE
flag那么从CPU到GPU复制的性情最好。

·        假若用了D3D11_CPU_ACCESS_READ该财富将是多少个由CPU缓存的能源,质量较低(可是辅助取回操作)

·        如若同时内定,READ比W安德拉ITE优先。

·        D3D11_USAGE_DYNAMIC
(仅能用来缓冲区型能源,无法用于纹理财富)用于火速的CPU->GPU内部存款和储蓄器传输。这种财富不但能够作为复制和源和指标,还足以看成纹理(用D3D的术语说,叫做着色器能源视图ShaderResourceView)在着色器中读取。不过着色器无法写入那种财富。那个能源的版本由驱动程序来决定,每一趟你用DISCAPRADOD
flag映射内部存款和储蓄器的时候,固然这块内部存款和储蓄器还在被GPU所接纳,驱动程序就会生出一块新的内存来,而不会等GPU的操作停止。它的含义在于提供一种流的法子将数据输送到GPU。

二.示例代码《nBodyCS》 //—————————————————————————–

// Reset the body system to its initial configuration

//—————————————————————————–

HRESULT NBodySystemCS::resetBodies(BodyDataconfigData)

{

    HRESULThr =S_OK;

      m_numBodies= configData.nBodies;

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

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

      D3DXVECTOR4*particleArray =newD3DXVECTOR4[m_numBodies* 3];

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

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

                                                  
   configData.position[i*3 +1],

                                                 
   configData.position[i*3 +2],

                                                     1.0);

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

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

                                                           configData.velocity[i*3

  • 1],

                                                                  
           configData.velocity[i*3 +2],

                                                                  
            1.0);

      }

//——————————————————————————————————

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

ComputeShader的1个很重庆大学的表征是结构化缓冲区和乱序访问视图。结构化缓冲区(structuredbuffer)在估测计算着色器中得以像数组一样访问。任意线程能够读写任意地方(即并行程序的分发scatter和收集gather动作)。乱序访问视图(unordered
access
view,UAV)是一种将调用方创设的财富绑定到着色器中的机制,并且同意……乱序访问。

扬言结构化缓冲区

大家能够用D3D11_RESOURCE_MISC_BUFFER_STRUCTURED来创设结构化缓冲区。下边钦赐的绑定flag表示同意着色器乱序访问。下面选拔的私下认可usage代表它能够被GPU进行读写,但要求复制到中转财富个中才能被CPU读写。

//——————————————————————————————————

      D3D11_SUBRESOURCE_DATAinitData = {particleArray,0, 0 };

 // 创建结构化缓冲区

      D3D11_BUFFER_DESCsbDesc;

      sbDesc.BindFlags           
=D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_SHADER_RESOURCE;

      sbDesc.CPUAccessFlags       = 0;

      sbDesc.MiscFlags           
=D3D11_RESOURCE_MISC_BUFFER_STRUCTURED;

      sbDesc.StructureByteStride  =sizeof(D3DXVECTOR4);

      sbDesc.ByteWidth            =sizeof(D3DXVECTOR4) *m_numBodies*
3;

      sbDesc.Usage                =D3D11_USAGE_DEFAULT;

      V_RETURN(m_pd3dDevice->CreateBuffer(&sbDesc,
&initData,&m_pStructuredBuffer) );

 

      //create the Shader Resource View (SRV) for the structured buffer

      D3D11_SHADER_RESOURCE_VIEW_DESCsbSRVDesc;

      sbSRVDesc.Buffer.ElementOffset          = 0;

      sbSRVDesc.Buffer.ElementWidth           =sizeof(D3DXVECTOR4);

      sbSRVDesc.Buffer.FirstElement           = 0;

      sbSRVDesc.Buffer.NumElements            =m_numBodies* 3;

      sbSRVDesc.Format                        =DXGI_FORMAT_UNKNOWN;

      sbSRVDesc.ViewDimension                
=D3D11_SRV_DIMENSION_BUFFER;

      V_RETURN(m_pd3dDevice->CreateShaderResourceView(m_pStructuredBuffer,
&sbSRVDesc,&m_pStructuredBufferSRV) );

宣示乱序访问视图

上面大家注明一(Wissu)个乱序访问视图。注意须求给她三个结构化缓冲区的指针

// 创制二个乱序访问视图,指向结构化缓冲区

      D3D11_UNORDERED_ACCESS_VIEW_DESCsbUAVDesc;

      sbUAVDesc.Buffer.FirstElement       = 0;

      sbUAVDesc.Buffer.Flags              = 0;

      sbUAVDesc.Buffer.NumElements        =m_numBodies* 3;

      sbUAVDesc.Format                    =DXGI_FORMAT_UNKNOWN;

      sbUAVDesc.ViewDimension            
=D3D11_UAV_DIMENSION_BUFFER;

      V_RETURN(m_pd3dDevice->CreateUnorderedAccessView(m_pStructuredBuffer,
&sbUAVDesc,&m_pStructuredBufferUAV) );

      delete[] particleArray;

      returnhr;

}

之后,在分摊着色器线程在此以前,大家须要激活着色器使用的结构化缓冲:

m_pd3dImmediateContext->CSSetUnorderedAccessViews( 0, 1, &g_pStructuredBufferUAV, &initCounts );

分摊线程之后,要是使用CS
4.x硬件,一定要将其免除绑定。因为CS4.x每条渲染流水生产线仅协理绑定一个UAV。

// 运行在 D3D10硬件上的时候: 每条流水线仅能绑定一个UAV

// 设成NULL就可以解除绑定

ID3D11UnorderedAccessView *pNullUAV = NULL;

m_pd3dImmediateContext->CSSetUnorderedAccessViews( 0, 1, &pNullUAV, &initCounts );

 

三.DirectCompute中的常量缓冲

常量缓冲(constantbuffer)是一组总结着色器运维时不能够改变的多寡。用作图形程序是,常量缓冲能够是理念矩阵或颜色常量。在通用总结程序中,常量缓冲能够存放诸如信号过滤的权重和图像处理的印证等数码。

一旦要采用常量缓冲:

·        创造缓冲区财富

·        用内存映射的艺术开始化数据(也得以用成效接口)

·        用CSSetConstantBuffers设定常量缓冲的值

上面代码创制了多个常量缓冲。注意常量缓冲的尺寸,那里我们掌握在HLSL中它是3个四元矢量。

// 创设常量缓冲

      D3D11_BUFFER_DESC cbDesc;

      cbDesc.Usage =D3D11_USAGE_DYNAMIC;

      cbDesc.BindFlags=D3D11_BIND_CONSTANT_BUFFER;

// CPU 可写, 这样我们得以每帧更新数据

      cbDesc.CPUAccessFlags=D3D11_CPU_ACCESS_WRITE;

      cbDesc.MiscFlags= 0;

      cbDesc.ByteWidth=sizeof(CB_DRAW);

      V_RETURN( pd3dDevice->CreateBuffer( &cbDesc,NULL, &m_pcbDraw)
);

 

      cbDesc.ByteWidth=sizeof(CB_UPDATE);

      V_RETURN( pd3dDevice->CreateBuffer( &cbDesc,NULL,
&m_pcbUpdate) );

     

      cbDesc.Usage =D3D11_USAGE_IMMUTABLE;

      cbDesc.BindFlags=D3D11_BIND_CONSTANT_BUFFER;

      cbDesc.CPUAccessFlags= 0;

      cbDesc.ByteWidth=sizeof(CB_IMMUTABLE);

      D3D11_SUBRESOURCE_DATA initData= { &cbImmutable, 0, 0 };

     V_RETURN( pd3dDevice->CreateBuffer(&cbDesc,&initData,
&m_pcbImmutable) );

接下去用内部存款和储蓄器映射的主意将数据发送到常量缓冲。常常程序员会在CPU程序里定义和HLSL一样的结构体,由此会用sizeof取得尺寸,然后将缓冲区的指针映射到组织体来填充数据

// 必须用 D3D11_MAP_WRITE_DISCARD

      D3D11_MAPPED_SUBRESOURCEmappedResource;

      V( m_pd3dImmediateContext->Map( m_pcbUpdate,0, D3D11_MAP_WRITE_DISCARD,
0, &mappedResource ) );

      CB_UPDATE* pcbUpdate= (CB_UPDATE*)mappedResource.pData;

      pcbUpdate->g_timestep=dt;

     pcbUpdate->g_softeningSquared = 0.01f;

      pcbUpdate->g_numParticles=m_numBodies;

     pcbUpdate->g_readOffset =m_readBuffer*m_numBodies;

     pcbUpdate->g_writeOffset = (1 -m_readBuffer)*m_numBodies;

m_pd3dImmediateContext->Unmap(m_pcbUpdate,0);

专注持筹握算着色器的输入变量(在那里正是常量缓冲)是真是“状态”变量的,因而在分摊总括着色器在此之前需求用CSSetShader()函数设置意况。那样测算着色器执行的时候就能访问到那些变量

// 在盘算着色器中激活

m_pd3dImmediateContext->CSSetShader(m_pCSUpdatePositionAndVelocity,NULL,
0 );

m_pd3dImmediateContext->CSSetConstantBuffers( 0, 1, &m_pcbUpdate );

// Run the CS

m_pd3dImmediateContext->Dispatch(m_numBodies/ 256, 1, 1 );

最后当总括着色器运转的时候,m_pCSUpdatePositionAndVelocity所针对的着色器就可以访问这些m_pcbUpdate常量缓冲

4.总括着色器(CS)HLSL编制程序

运转在显卡上的臆度着色器是用HLSL(High Level Shader Language
高级着色器语言)写成的。在我们的例证中它是以文件格局存在,并且在运行时动态编写翻译的。总计着色器是一种单一程序被众多线程并行执行的顺序。这一个线程分成四个“线程组”,在线程组内的线程之间能够共享数据或相互同步。

GPU硬件架构

GPU硬件结构首要由以下多少个首要模块组成:内部存款和储蓄器(全局的,常量的,共享的);流处理器簇(SM);流处理器(SP)。如图所示:

SP: 最宗旨的处理单元,streamingprocessor
最后实际的通令和职责都以在sp上处理的。GPU实行并行总括,也正是累累个sp同时做拍卖

SM:多少个sp加上其余的部分能源结合1个sm, streaming multiprocessor.
其余能源也正是存款和储蓄财富,共享内部存款和储蓄器,寄储器等。

WAENVISIONP:GPU执行顺序时的调度单位,如今cuda的warp的轻重缓急为32,同在二个warp的线程,以分裂数额财富执行同样的通令。

 

 ACCESS 2

GPU实际上是二个SM的阵列,每种SM包蕴N个核。二个GPU设备中富含1个或多个SM。SM内部整合结构图如下所示:

 ACCESS 3

4.1线程网格

七个线程网格是由若干线程块结成的,各类线程块是二维的,拥有X轴Y轴。此时我们最多能开启Y*X*T个线程。thread–>block–>grid:在使用cuda进行编程时,三个grid分为五个block,而3个block分为七个thread。在那之中职分划分到是还是不是影响最后的执行职能。划分的基于是职责特点和GPU本人的硬件天性。1个sm只会执行二个block里的warp,当该block里warp执行完才会实施其它block里的warp。进行划分时,最棒有限支持每种block里的warp比较客观,那样能够2个sm能够轮换执行里面的warp,从而进步功能,别的,在分配block时,要基于GPU的sm个数,分配出合理的block数,让GPU的sm都应用起来,提利用率。分配时,也要考虑到同二个线程block的能源难点,不要现身对应的能源不够。

倘诺我们在看一张高清的图形,那张图片的分辨率为一九一八x1080。日常线程块中的线程数量极其是3个线程束大小的整数倍,即32的平头倍。本例中我们在线程块上打开19一个线程。每一个线程块19二个线程,很简单总结出一行图形必要11个线程块(如图4.1)。在那里选在192那么些是因为X轴方向处理的数码大小一九一九是它的整数倍,192又是线程束大小的平头倍。GPU上的四个线程束的大大小小是32(Intel公司封存着对这些参数修改的义务),他们提供三个土生土长变量-WRAPSIZE,我们能够透过那些变量来取得硬件支撑的线程束的轻重。

 ACCESS 4

图4.1按行分布的线程块

在X轴方向的顶部我们得以博得线程的目录,在Y轴方向大家能够赢得行号。由于每一行只处理一行像素,每一行有12个线程块,因而我们必要1080行来处理整张图片,一共1080*10=10800个线程块。在费米架构的硬件上,贰个SM能够处理8个线程块,所以从应用层角度来说一共索要1355个(总共10800个线程块
/
每一种SM能调度的几个线程块)SM来完全完结相互之间。但当下费米架构的硬件唯有拾几个SM可供使用(GTx580)即种种SM将被分配6七十个线程块举办处理。

上述例子非常粗大略,数据分布整齐不难精晓。可是频仍大家的多寡只怕不是一维的,那时我们能够行使二维模块恐怕三维矩阵来存款和储蓄数据。

4.2网格(Grid)、线程块(Block)和线程(Thread)的组织关系

CUDA的软件架构由网格(Grid)、线程块(Block)和线程(Thread)组成,也就是把GPU上的乘除单元分为若干(2~3)个网格,各个网格内含有若干(65535)个线程块,各类线程块包含若干(512)个线程,三者的关系如下图:

ACCESS 5

Thread,block,grid是CUDA编制程序上的定义,为了有利于程序员软件设计,组织线程。

·        thread:一个CUDA的并行程序会被以众多少个threads来执行。

·        block:数个threads会被群结合3个block,同二个block中的threads能够联手,也足以因而shared
memory通讯。

·        grid:多个blocks则会再结合grid。

 

网格(Grid)、线程块(Block)和线程(Thread)的最大数额

CUDA中得以创设的网格数量跟GPU的持筹握算能力有关,可创立的Grid、Block和Thread的最大数据参看以下表格:

 

 ACCESS 6

在单纯维度上,程序的执行能够由多达3*65535*512=100661760(一亿)个线程并行执行,那对在CPU上创设并行线程来说是不行想像的。

 

线程索引的总结公式

3个Grid能够分包四个Blocks,Blocks的团伙章程得以是一维的,二维恐怕三维的。block包涵多个Threads,那个Threads的团体章程也能够是一维,二维大概三维的。

CUDA中每二个线程都有三个唯一的标识ID—ThreadIdx,那一个ID随着Grid和Block的分割格局的两样而转变,那里给出Grid和Block差异划分方式下线程索引ID的总括公式。

 

1、 grid划分成1维,block划分为1维

    int threadId = blockIdx.x *blockDim.x + threadIdx.x;  
      
2、 grid划分成1维,block划分为2维  

    int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y*
blockDim.x + threadIdx.x;  
    
3、 grid划分成1维,block划分为3维  

    int threadId = blockIdx.x * blockDim.x * blockDim.y *
blockDim.z 
                      + threadIdx.z * blockDim.y * blockDim.x  
                      + threadIdx.y * blockDim.x + threadIdx.x;  

4、 grid划分成2维,block划分为1维  

    int blockId = blockIdx.y * gridDim.x + blockIdx.x;  
    int threadId = blockId * blockDim.x + threadIdx.x;  
  5、 grid划分成2维,block划分为2维 

    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
                      + (threadIdx.y * blockDim.x) + threadIdx.x;  
6、 grid划分成2维,block划分为3维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y *
blockDim.z) 
                      + (threadIdx.z * (blockDim.x * blockDim.y))  
                      + (threadIdx.y * blockDim.x) + threadIdx.x;    
7、 grid划分成3维,block划分为1维 

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     +gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * blockDim.x + threadIdx.x;  
   
  
8、 grid划分成3维,block划分为2维  

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     +gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
                      + (threadIdx.y * blockDim.x) + threadIdx.x;  
   
  
9、 grid划分成3维,block划分为3维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     +gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y *
blockDim.z) 
                      + (threadIdx.z * (blockDim.x * blockDim.y))  
                      + (threadIdx.y * blockDim.x) + threadIdx.x;   
 
4.3 nBodyCS.hlsl

 

m_pd3dImmediateContext->Dispatch( m_numBodies / 256,1, 1 );

Dispatch(m_numBodies/ 256, 1, 1)创建了(m_numBodies/ 256)x 1 x
一个线程组。而每二个线程组中的线程是在着色器代码中用这些语法来钦赐[numthreads(BLOCK_SIZE,1,1)]

/* 那象征线程组中的线程数,本例中是BLOCK_SIZE x1x1 = 256个线程

[numthreads(BLOCK_SIZE,1,1)]

void NBodyUpdate(uint threadId        : SV_GroupIndex,

                 uint3 groupId        : SV_GroupID,

                 uint3 globalThreadId :SV_DispatchThreadID)

{

float4pos = particles[g_readOffset + globalThreadId.x];

   float4 vel = particles[2 * g_numParticles + globalThreadId.x];

            //compute acceleration

      float3 accel = computeBodyAccel(pos, threadId, groupId);

      //Leapfrog-Verlet integration of velocity and position

      vel.xyz+= accel * g_timestep;

      pos.xyz+= vel   * g_timestep;

 

      particles[g_writeOffset+ globalThreadId.x]      = pos;

      particles[2* g_numParticles + globalThreadId.x] = vel;

}

 

// Computes the total acceleration on the body with position myPos 

// caused by the gravitational attraction of all other bodies in 

// the simulation

float3 computeBodyAccel(float4 bodyPos, uint threadId, uint blockId)

{

    float3 acceleration = {0.0f, 0.0f, 0.0f};

    uint p = BLOCK_SIZE;

    uint n = g_numParticles;

    uint numTiles = n / p;

 

    for (uint tile = 0; tile < numTiles; tile++) 

    {

        // 取全体线程组中千篇一律索引的threadid 索引threadid 【0,255】范围

        sharedPos[threadId] = particles[g_readOffset + tile * p +
threadId];

     //
 阻止组中享有线程的施行,直到全数组共享访问成功,组中的具备线程都

      //已达到规定的标准此调用。

        GroupMemoryBarrierWithGroupSync();

     
 //总计 bodyPos同相同threakid分歧的线程组的粒子坐标的重力引力影响值

        acceleration = gravitation(bodyPos, acceleration);

        GroupMemoryBarrierWithGroupSync();

    }

    return acceleration;

}