#NEWLINE# def _execute_reduction( self , dary , size , stream , init = 0 ) : #NEWLINE# #INDENT# from numba import cuda #NEWLINE# ( nbtype , size ) = self . _type_and_size ( dary , size ) #NEWLINE# @ cuda . jit #NEWLINE# def copy_strides( arr , n , stride , tpb ) : #NEWLINE# #INDENT# sm = cuda . shared . array ( 1 , dtype = uint32 ) #NEWLINE# i = cuda . threadIdx . x #NEWLINE# base = 0 #NEWLINE# if ( i == 0 ) : #NEWLINE# #INDENT# sm [ 0 ] = 0 #NEWLINE# #UNINDENT# val = arr [ 0 ] #NEWLINE# while ( base < n ) : #NEWLINE# #INDENT# idx = ( base + i ) #NEWLINE# if ( idx < n ) : #NEWLINE# #INDENT# val = arr [ ( idx * stride ) ] #NEWLINE# #UNINDENT# cuda . syncthreads ( ) #NEWLINE# if ( ( base + i ) < n ) : #NEWLINE# #INDENT# arr [ ( sm [ 0 ] + i ) ] = val #NEWLINE# #UNINDENT# if ( i == 0 ) : #NEWLINE# #INDENT# sm [ 0 ] += tpb #NEWLINE# #UNINDENT# base += tpb #NEWLINE# #UNINDENT# #UNINDENT# while ( size >= 2 ) : #NEWLINE# #INDENT# p2size = ( 2 ** int ( math . log ( size , 2 ) ) ) #NEWLINE# diffsize = ( size - p2size ) #NEWLINE# ( kernel , blocksize ) = self . _get_kernel ( p2size , nbtype ) #NEWLINE# size = p2size #NEWLINE# gridsize = ( size // blocksize ) #NEWLINE# assert ( gridsize <= p2size ) #NEWLINE# if ( gridsize > 0 ) : #NEWLINE# #INDENT# worksize = ( blocksize * gridsize ) #NEWLINE# blocksize = ( blocksize // 2 ) #NEWLINE# assert ( ( size - worksize ) == 0 ) #NEWLINE# kernel [ ( gridsize , blocksize , stream ) ] ( dary , dary , worksize , blocksize ) #NEWLINE# copy_strides [ ( 1 , 512 , stream ) ] ( dary , gridsize , blocksize , 512 ) #NEWLINE# size = gridsize #NEWLINE# #UNINDENT# if ( diffsize > 0 ) : #NEWLINE# #INDENT# itemsize = dary . dtype . itemsize #NEWLINE# dst = dary . gpu_data . view ( ( size * itemsize ) ) #NEWLINE# src = dary . gpu_data . view ( ( p2size * itemsize ) ) #NEWLINE# cuda . driver . device_to_device ( dst , src , ( diffsize * itemsize ) , stream = stream ) #NEWLINE# size += diffsize #NEWLINE# #UNINDENT# #UNINDENT# if ( init != 0 ) : #NEWLINE# #INDENT# ( kernel , _ ) = self . _get_kernel ( 0 , nbtype ) #NEWLINE# kernel ( dary , init )