Extras/RigidBodyGpuPipeline/dynamics/basic_demo/Stubs/batchingKernels.cl

   1 /*\r
   2 Copyright (c) 2012 Advanced Micro Devices, Inc.  \r
   3 \r
   4 This software is provided 'as-is', without any express or implied warranty.\r
   5 In no event will the authors be held liable for any damages arising from the use of this software.\r
   6 Permission is granted to anyone to use this software for any purpose, \r
   7 including commercial applications, and to alter it and redistribute it freely, \r
   8 subject to the following restrictions:\r
   9 \r
  10 1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.\r
  11 2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\r
  12 3. This notice may not be removed or altered from any source distribution.\r
  13 */\r
  14 //Originally written by Takahiro Harada\r
  15 \r
  16 \r
  17 #pragma OPENCL EXTENSION cl_amd_printf : enable\r
  18 #pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\r
  19 #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\r
  20 #pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable\r
  21 #pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable\r
  22 \r
  23 #ifdef cl_ext_atomic_counters_32\r
  24 #pragma OPENCL EXTENSION cl_ext_atomic_counters_32 : enable\r
  25 #else\r
  26 #define counter32_t volatile __global int*\r
  27 #endif\r
  28 \r
  29 \r
  30 typedef unsigned int u32;\r
  31 typedef unsigned short u16;\r
  32 typedef unsigned char u8;\r
  33 \r
  34 #define GET_GROUP_IDX get_group_id(0)\r
  35 #define GET_LOCAL_IDX get_local_id(0)\r
  36 #define GET_GLOBAL_IDX get_global_id(0)\r
  37 #define GET_GROUP_SIZE get_local_size(0)\r
  38 #define GET_NUM_GROUPS get_num_groups(0)\r
  39 #define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\r
  40 #define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE)\r
  41 #define AtomInc(x) atom_inc(&(x))\r
  42 #define AtomInc1(x, out) out = atom_inc(&(x))\r
  43 #define AppendInc(x, out) out = atomic_inc(x)\r
  44 #define AtomAdd(x, value) atom_add(&(x), value)\r
  45 #define AtomCmpxhg(x, cmp, value) atom_cmpxchg( &(x), cmp, value )\r
  46 #define AtomXhg(x, value) atom_xchg ( &(x), value )\r
  47 \r
  48 \r
  49 #define SELECT_UINT4( b, a, condition ) select( b,a,condition )\r
  50 \r
  51 #define make_float4 (float4)\r
  52 #define make_float2 (float2)\r
  53 #define make_uint4 (uint4)\r
  54 #define make_int4 (int4)\r
  55 #define make_uint2 (uint2)\r
  56 #define make_int2 (int2)\r
  57 \r
  58 \r
  59 #define max2 max\r
  60 #define min2 min\r
  61 \r
  62 \r
  63 #define WG_SIZE 64\r
  64 \r
  65 \r
  66 \r
  67 typedef struct \r
  68 {\r
  69         float4 m_worldPos[4];\r
  70         float4 m_worldNormal;\r
  71         u32 m_coeffs;\r
  72         int m_batchIdx;\r
  73 \r
  74         u32 m_bodyA;\r
  75         u32 m_bodyB;\r
  76 }Contact4;\r
  77 \r
  78 typedef struct \r
  79 {\r
  80         int m_n;\r
  81         int m_start;\r
  82         int m_staticIdx;\r
  83         int m_paddings[1];\r
  84 } ConstBuffer;\r
  85 \r
  86 typedef struct \r
  87 {\r
  88         u32 m_a;\r
  89         u32 m_b;\r
  90         u32 m_idx;\r
  91 }Elem;\r
  92 \r
  93 #define STACK_SIZE (WG_SIZE*10)\r
  94 //#define STACK_SIZE (WG_SIZE)\r
  95 #define RING_SIZE 1024\r
  96 #define RING_SIZE_MASK (RING_SIZE-1)\r
  97 #define CHECK_SIZE (WG_SIZE)\r
  98 \r
  99 \r
 100 #define GET_RING_CAPACITY (RING_SIZE - ldsRingEnd)\r
 101 #define RING_END ldsTmp\r
 102 \r
 103 u32 readBuf(__local u32* buff, int idx)\r
 104 {\r
 105         idx = idx % (32*CHECK_SIZE);\r
 106         int bitIdx = idx%32;\r
 107         int bufIdx = idx/32;\r
 108         return buff[bufIdx] & (1<<bitIdx);\r
 109 }\r
 110 \r
 111 void writeBuf(__local u32* buff, int idx)\r
 112 {\r
 113         idx = idx % (32*CHECK_SIZE);\r
 114         int bitIdx = idx%32;\r
 115         int bufIdx = idx/32;\r
 116 //      buff[bufIdx] |= (1<<bitIdx);\r
 117         atom_or( &buff[bufIdx], (1<<bitIdx) );\r
 118 }\r
 119 \r
 120 u32 tryWrite(__local u32* buff, int idx)\r
 121 {\r
 122         idx = idx % (32*CHECK_SIZE);\r
 123         int bitIdx = idx%32;\r
 124         int bufIdx = idx/32;\r
 125         u32 ans = (u32)atom_or( &buff[bufIdx], (1<<bitIdx) );\r
 126         return ((ans >> bitIdx)&1) == 0;\r
 127 }\r
 128 \r
 129 //      batching on the GPU\r
 130 __kernel void CreateBatches( __global Contact4* gConstraints, __global Contact4* gConstraintsOut,\r
 131                 __global u32* gN, __global u32* gStart, \r
 132                 ConstBuffer cb )\r
 133 {\r
 134         __local u32 ldsStackIdx[STACK_SIZE];\r
 135         __local u32 ldsStackEnd;\r
 136         __local Elem ldsRingElem[RING_SIZE];\r
 137         __local u32 ldsRingEnd;\r
 138         __local u32 ldsTmp;\r
 139         __local u32 ldsCheckBuffer[CHECK_SIZE];\r
 140         __local u32 ldsFixedBuffer[CHECK_SIZE];\r
 141         __local u32 ldsGEnd;\r
 142         __local u32 ldsDstEnd;\r
 143 \r
 144         int wgIdx = GET_GROUP_IDX;\r
 145         int lIdx = GET_LOCAL_IDX;\r
 146         \r
 147         const int m_n = gN[wgIdx];\r
 148         const int m_start = gStart[wgIdx];\r
 149         const int m_staticIdx = cb.m_staticIdx;\r
 150                 \r
 151         if( lIdx == 0 )\r
 152         {\r
 153                 ldsRingEnd = 0;\r
 154                 ldsGEnd = 0;\r
 155                 ldsStackEnd = 0;\r
 156                 ldsDstEnd = m_start;\r
 157         }\r
 158         \r
 159 //      while(1)\r
 160         for(int ie=0; ie<250; ie++)\r
 161         {\r
 162                 ldsFixedBuffer[lIdx] = 0;\r
 163 \r
 164                 for(int giter=0; giter<4; giter++)\r
 165                 {\r
 166                         int ringCap = GET_RING_CAPACITY;\r
 167                 \r
 168                         //      1. fill ring\r
 169                         if( ldsGEnd < m_n )\r
 170                         {\r
 171                                 while( ringCap > WG_SIZE )\r
 172                                 {\r
 173                                         if( ldsGEnd >= m_n ) break;\r
 174                                         if( lIdx < ringCap - WG_SIZE )\r
 175                                         {\r
 176                                                 int srcIdx;\r
 177                                                 AtomInc1( ldsGEnd, srcIdx );\r
 178                                                 if( srcIdx < m_n )\r
 179                                                 {\r
 180                                                         int dstIdx;\r
 181                                                         AtomInc1( ldsRingEnd, dstIdx );\r
 182                                                         \r
 183                                                         int a = gConstraints[m_start+srcIdx].m_bodyA;\r
 184                                                         int b = gConstraints[m_start+srcIdx].m_bodyB;\r
 185                                                         ldsRingElem[dstIdx].m_a = (a>b)? b:a;\r
 186                                                         ldsRingElem[dstIdx].m_b = (a>b)? a:b;\r
 187                                                         ldsRingElem[dstIdx].m_idx = srcIdx;\r
 188                                                 }\r
 189                                         }\r
 190                                         ringCap = GET_RING_CAPACITY;\r
 191                                 }\r
 192                         }\r
 193 \r
 194                         GROUP_LDS_BARRIER;\r
 195         \r
 196                         //      2. fill stack\r
 197                         __local Elem* dst = ldsRingElem;\r
 198                         if( lIdx == 0 ) RING_END = 0;\r
 199 \r
 200                         int srcIdx=lIdx;\r
 201                         int end = ldsRingEnd;\r
 202 \r
 203                         {\r
 204                                 for(int ii=0; ii<end; ii+=WG_SIZE, srcIdx+=WG_SIZE)\r
 205                                 {\r
 206                                         Elem e;\r
 207                                         if(srcIdx<end) e = ldsRingElem[srcIdx];\r
 208                                         bool done = (srcIdx<end)?false:true;\r
 209 \r
 210                                         for(int i=lIdx; i<CHECK_SIZE; i+=WG_SIZE) ldsCheckBuffer[lIdx] = 0;\r
 211                                         \r
 212                                         if( !done )\r
 213                                         {\r
 214                                                 int aUsed = readBuf( ldsFixedBuffer, e.m_a);\r
 215                                                 int bUsed = readBuf( ldsFixedBuffer, e.m_b);\r
 216 \r
 217                                                 if( aUsed==0 && bUsed==0 )\r
 218                                                 {\r
 219                                                         int aAvailable;\r
 220                                                         int bAvailable;\r
 221 \r
 222                                                         aAvailable = tryWrite( ldsCheckBuffer, e.m_a );\r
 223                                                         bAvailable = tryWrite( ldsCheckBuffer, e.m_b );\r
 224 \r
 225                                                         //aAvailable = (m_staticIdx == e.m_a)? 1: aAvailable;\r
 226                                                         //bAvailable = (m_staticIdx == e.m_b)? 1: bAvailable;\r
 227 \r
 228                                                         bool success = (aAvailable && bAvailable);\r
 229                                                         if(success)\r
 230                                                         {\r
 231                                                                 writeBuf( ldsFixedBuffer, e.m_a );\r
 232                                                                 writeBuf( ldsFixedBuffer, e.m_b );\r
 233                                                         }\r
 234                                                         done = success;\r
 235                                                 }\r
 236                                         }\r
 237 \r
 238                                         //      put it aside\r
 239                                         if(srcIdx<end)\r
 240                                         {\r
 241                                                 if( done )\r
 242                                                 {\r
 243                                                         int dstIdx; AtomInc1( ldsStackEnd, dstIdx );\r
 244                                                         if( dstIdx < STACK_SIZE )\r
 245                                                                 ldsStackIdx[dstIdx] = e.m_idx;\r
 246                                                         else{\r
 247                                                                 done = false;\r
 248                                                                 AtomAdd( ldsStackEnd, -1 );\r
 249                                                         }\r
 250                                                 }\r
 251                                                 if( !done )\r
 252                                                 {\r
 253                                                         int dstIdx; AtomInc1( RING_END, dstIdx );\r
 254                                                         dst[dstIdx] = e;\r
 255                                                 }\r
 256                                         }\r
 257 \r
 258                                         //      if filled, flush\r
 259                                         if( ldsStackEnd == STACK_SIZE )\r
 260                                         {\r
 261                                                 for(int i=lIdx; i<STACK_SIZE; i+=WG_SIZE)\r
 262                                                 {\r
 263                                                         int idx = m_start + ldsStackIdx[i];\r
 264                                                         int dstIdx; AtomInc1( ldsDstEnd, dstIdx );\r
 265                                                         gConstraintsOut[ dstIdx ] = gConstraints[ idx ];\r
 266                                                         gConstraintsOut[ dstIdx ].m_batchIdx = ie;\r
 267                                                 }\r
 268                                                 if( lIdx == 0 ) ldsStackEnd = 0;\r
 269 \r
 270                                                 //for(int i=lIdx; i<CHECK_SIZE; i+=WG_SIZE) \r
 271                                                 ldsFixedBuffer[lIdx] = 0;\r
 272                                         }\r
 273                                 }\r
 274                         }\r
 275 \r
 276                         if( lIdx == 0 ) ldsRingEnd = RING_END;\r
 277                 }\r
 278 \r
 279                 GROUP_LDS_BARRIER;\r
 280 \r
 281                 for(int i=lIdx; i<ldsStackEnd; i+=WG_SIZE)\r
 282                 {\r
 283                         int idx = m_start + ldsStackIdx[i];\r
 284                         int dstIdx; AtomInc1( ldsDstEnd, dstIdx );\r
 285                         gConstraintsOut[ dstIdx ] = gConstraints[ idx ];\r
 286                         gConstraintsOut[ dstIdx ].m_batchIdx = ie;\r
 287                 }\r
 288 \r
 289                 //      in case it couldn't consume any pair. Flush them\r
 290                 //      todo. Serial batch worth while?\r
 291                 if( ldsStackEnd == 0 )\r
 292                 {\r
 293                         for(int i=lIdx; i<ldsRingEnd; i+=WG_SIZE)\r
 294                         {\r
 295                                 int idx = m_start + ldsRingElem[i].m_idx;\r
 296                                 int dstIdx; AtomInc1( ldsDstEnd, dstIdx );\r
 297                                 gConstraintsOut[ dstIdx ] = gConstraints[ idx ];\r
 298                                 gConstraintsOut[ dstIdx ].m_batchIdx = 100+i;\r
 299                         }\r
 300                         GROUP_LDS_BARRIER;\r
 301                         if( lIdx == 0 ) ldsRingEnd = 0;\r
 302                 }\r
 303 \r
 304                 if( lIdx == 0 ) ldsStackEnd = 0;\r
 305 \r
 306                 GROUP_LDS_BARRIER;\r
 307 \r
 308                 //      termination\r
 309                 if( ldsGEnd == m_n && ldsRingEnd == 0 )\r
 310                         break;\r
 311         }\r
 312 \r
 313 \r
 314 }\r
 315 \r
 316 \r
 317 \r
 318 \r
 319 \r
 320 \r
 321 \r
 322 \r
 323 \r
 324 \r
 325 \r
 326 \r
 327 \r
 328 \r
 329 \r
 330 \r
 331 \r
 332 \r
 333 \r
 334 \r
 335 \r
 336 \r
 337 \r
 338 \r