@cuda.jit
def gpu_single_block_sum(arr, out):
 temp = cuda.shared.array(gpu_block_sum_max_blockdim, dtype=float32)
 tid = cuda.threadIdx.x
 blksz = cuda.blockDim.x
 temp[tid] = 0
 for i in range(tid, arr.size, blksz):
  temp[tid] += arr[i]
 cuda.syncthreads()
 if (tid == 0):
  for i in range(1, blksz):
   temp[0] += temp[i]
  out[0] = temp[0]

@ cuda . jit def gpu_single_block_sum ( arr , out ) : temp = cuda . shared . array ( gpu_block_sum_max_blockdim , dtype = float32 ) tid = cuda . threadIdx . x blksz = cuda . blockDim . x temp [ tid ] = 0 for i in range ( tid , arr . size , blksz ) : temp [ tid ] += arr [ i ] cuda . syncthreads ( ) if ( tid == 0 ) : for i in range ( 1 , blksz ) : temp [ 0 ] += temp [ i ] out [ 0 ] = temp [ 0 ]
a naive single threadblock sum reduction .