Опитвайки се да разберем фрагмента за крайна редукция на cuda sdk по отношение на сложността и условията на изпълнение

Опитвам се да разбера паралелното намаляване в Cuda (много е интересно). В последния ми въпрос относно паралелната редукция Робърт Кровела даде много интуитивно и подробно обяснение, което ми помогна много. Много интуитивен. Сега, гледайки пробите за намаляване на Cuda SDK, има няколко тъмни петна.

1. Защо (в коментара по-долу) сложността на работата е запазена в O(n)? В какъв случай би се случило обратното? Имам същия въпрос относно сложността на стъпките.

   Тази версия добавя множество елементи на нишка последователно. Това намалява общата цена на алгоритъма, като същевременно запазва сложността на работата O(n) и сложността на стъпките
   O(log n).

2. Може ли някой да даде интуитивен пример за коментара относно количеството споделена памет („Забележете, че това ядро ​​се нуждае... разпределете blockSize*sizeof(T) байта“) и как е свързано с кода?

3. Защо nIsPow2 е толкова важно? Някой може ли да обясни или да даде пример?

4. Защо използваме mySum например в следващите задачи? sdata[tid] = mySum = mySum + sdata[tid + 256];, а не само sdata[tid]+=data[tid+256] като в презентацията на Маркс Харис?

   /*This version adds multiple elements per thread sequentially.  This reduces the   
   overall cost of the algorithm while keeping the work complexity O(n) and the       
   step complexity O(log n).
   (Brent's Theorem optimization)
   
   Note, this kernel needs a minimum of 64*sizeof(T) bytes of shared memory.
   In other words if blockSize <= 32, allocate 64*sizeof(T) bytes.
   If blockSize > 32, allocate blockSize*sizeof(T) bytes.*/
   
   template <class T, unsigned int blockSize, bool nIsPow2> 
   __global__ void
   
   reduce6(T *g_idata, T *g_odata, unsigned int n)
   
   {
   T *sdata = SharedMemory<T>();
   
   // perform first level of reduction,
   // reading from global memory, writing to shared memory
   unsigned int tid = threadIdx.x;
   unsigned int i = blockIdx.x*blockSize*2 + threadIdx.x;
   unsigned int gridSize = blockSize*2*gridDim.x;
   
   T mySum = 0;
   
   // we reduce multiple elements per thread.  The number is determined by the
   // number of active thread blocks (via gridDim).  More blocks will result
   // in a larger gridSize and therefore fewer elements per thread
   while (i < n)
   {
       mySum += g_idata[i];
   
       // ensure we don't read out of bounds -- this is optimized away for powerOf2 sized arrays
       if (nIsPow2 || i + blockSize < n)
           mySum += g_idata[i+blockSize];
   
       i += gridSize;
   }
   
   // each thread puts its local sum into shared memory
   sdata[tid] = mySum;
   __syncthreads();
   
   
   // do reduction in shared mem
   if (blockSize >= 512)
   {
       if (tid < 256)
       {
           sdata[tid] = mySum = mySum + sdata[tid + 256];
       }
   
       __syncthreads();
   }
   
   if (blockSize >= 256)
   {
       if (tid < 128)
       {
           sdata[tid] = mySum = mySum + sdata[tid + 128];
       }
   
       __syncthreads();
   }
   
   if (blockSize >= 128)
   {
       if (tid <  64)
       {
           sdata[tid] = mySum = mySum + sdata[tid +  64];
       }
   
       __syncthreads();
   }
   
   if (tid < 32)
   {
       // now that we are using warp-synchronous programming (below)
       // we need to declare our shared memory volatile so that the compiler
       // doesn't reorder stores to it and induce incorrect behavior.
       volatile T *smem = sdata;
   
       if (blockSize >=  64)
       {
           smem[tid] = mySum = mySum + smem[tid + 32];
       }
   
       if (blockSize >=  32)
       {
           smem[tid] = mySum = mySum + smem[tid + 16];
       }
   
       if (blockSize >=  16)
       {
           smem[tid] = mySum = mySum + smem[tid +  8];
       }
   
       if (blockSize >=   8)
       {
           smem[tid] = mySum = mySum + smem[tid +  4];
       }
   
       if (blockSize >=   4)
       {
           smem[tid] = mySum = mySum + smem[tid +  2];
       }
   
       if (blockSize >=   2)
       {
           smem[tid] = mySum = mySum + smem[tid +  1];
       }
   }
   
   // write result for this block to global mem
   if (tid == 0)
       g_odata[blockIdx.x] = sdata[0];
   

   }