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

 

图片 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对象,然后调用Enum艾达pters并传一个意味着着枚举显卡的平头。倘若不存在,它会面重回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的线程,以不同数额资源执行同一之命令。

 

 图片 2

GPU实际上是一个SM的阵列,每个SM包含N个对。一个GPU设备受到带有一个仍旧六个SM。SM内部整合结构图如下所示:

 图片 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,我们可通过是变量来博取硬件协理的线程束的轻重缓急。

 图片 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)个线程,三者的涉而下图:

图片 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的无限可怜数量参看以下表格:

 

 图片 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;

}