nBodyCS学习笔记之计着色器

 

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

如上语句分派了16×16独线程组。

顾,着色器的输入常见考虑成“状态”。就是说你该在分摊着色器程序之前设定状态,而若分派了,“状态”决定输入变量的价值。所以在色器分派代码通常应像这样:

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_WRITE
flag那么从CPU到GPU复制的性能最佳。

·        如果就此了D3D11_CPU_ACCESS_READ该资源以是一个是因为CPU缓存的资源,性能比逊色(但是支持取回操作)

·        如果以指定,READ比WRITE优先。

·        D3D11_USAGE_DYNAMIC
(仅能用于缓冲区型资源,不可知用来纹理资源)用于快速的CPU->GPU内存传输。这种资源不但可看作复制和源和目标,还好看做纹理(用D3D的术语说,叫做着色器资源视图ShaderResourceView)在正色器中读取。但是着色器不克写副这种资源。这些资源的版由驱动程序来支配,每次你用DISCARD
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的一个老大重要之特色是结构化缓冲区和乱序访问视图。结构化缓冲区(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) );

扬言乱序访问视图

脚我们声明一个乱序访问视图。注意要让他一个结构化缓冲区的指针

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

      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中它是一个季长矢量。

// 创建常量缓冲

      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加上另外的片段资源做一个sm, streaming multiprocessor.
其他资源为就是是存储资源,共享内存,寄储器等。

WARP:GPU执行顺序时之调度单位,目前cuda的warp的高低为32,同在一个warp的线程,以不同数额资源执行同样的吩咐。

 

 ACCESS 2

GPU实际上是一个SM的阵列,每个SM包含N个审批。一个GPU设备受到寓一个要多独SM。SM内部整合结构图如下所示:

 ACCESS 3

4.1线程网格

一个线程网格是由于若干线程块组成的,每个线程块是二维的,拥有X轴Y轴。此时咱们绝多能开启Y*X*T个线程。thread–>block–>grid:在应用cuda进行编程时,一个grid分为多只block,而一个block分为多只thread。其中任务划分到是不是影响最后之推行效能。划分的基于是职责特点和GPU本身的硬件特性。一个sm只会履行一个block里的warp,当该block里warp执行了才会实施外block里的warp。进行分割时,最好保证每个block里之warp比较客观,那样可以一个sm可以轮流执行中的warp,从而提高效率,此外,在分配block时,要根据GPU的sm个数,分配产生客观之block数,让GPU的sm都下起来,提利用率。分配时,也使考虑到和一个线程block的资源问题,不要出现对应之资源不够。

使我们以羁押同样摆高清的图片,这张图片的分辨率为1920
x1080。通常线程块饱受之线程数量极其是一个线程束大小的平头倍增,即32的平头加倍。本例中我们在线程块上开启192个线程。每个线程块192单线程,很容易计算产生一行图形需要10只线程块(如图4.1)。在此选在192斯是以X轴方向处理的数码大小1920凡它们的平头加倍,192而是线程束大小的整数倍增。GPU上之一个线程束的轻重缓急是32(英伟及公司封存在对斯参数修改的权),他们提供一个原始变量-WRAPSIZE,我们可由此这变量来博硬件支持的线程束的深浅。

 ACCESS 4

图4.1随行分布的线程块

每当X轴方向的顶部我们可以获线程的目,在Y轴方向我们得以抱行号。由于各一行就处理一行像素,每一行来10独线程块,因此我们得1080行来处理整张图片,一共1080*10=10800单线程块。在费米架构的硬件上,一个SM可以处理8只线程块,所以打应用层角度来说无异伙得1350个(总共10800个线程块
/
每个SM能调度的8单线程块)SM来了落实相互之间。但时费米架构的硬件只发生16独SM可供使用(GTx580)即每个SM将于分配675单线程块进行拍卖。

上述例子十分粗略,数据分布整齐容易掌握。但是频繁我们的数额也许无是同样维的,这时我们得行使二维模块或者三维矩阵来囤积数据。

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会为众结合一个block,同一个block中之threads可以协同,也可经过shared
memory通信。

·        grid:多独blocks则会再度成grid。

 

网格(Grid)、线程块(Block)和线程(Thread)的卓绝可怜数额

CUDA中得创造的网格数量和GPU的乘除能力有关,可创的Grid、Block和Thread的最为充分数额参看以下表格:

 

 ACCESS 6

于单纯维度上,程序的履行得由多上3*65535*512=100661760(一亿)个线程并行执行,这对以CPU上创办并行线程来说是不足想像的。

 

线程索引的计算公式

一个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
1只线程组。而各一个线程组中的线程是当在色器代码中用这个语法来指定[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;

}