#NEWLINE# def _precompile_kernels( self , nbtype ) : #NEWLINE# #INDENT# from numba import cuda #NEWLINE# if ( nbtype not in self . _cached_types ) : #NEWLINE# #INDENT# for blocksize in ( 1 , 2 , 4 , 8 , 16 , 32 , 64 , 128 , 256 , 512 ) : #NEWLINE# #INDENT# key = ( nbtype , blocksize ) #NEWLINE# self . _kernels [ key ] = reduction_template ( self . _binop , nbtype , blocksize ) #NEWLINE# #UNINDENT# binop = cuda . jit ( ( nbtype , nbtype ) , device = True ) ( self . _binop ) #NEWLINE# @ cuda . jit ( ( nbtype [ : ] , nbtype ) ) #NEWLINE# def reduction_finish( arr , init ) : #NEWLINE# #INDENT# arr [ 0 ] = binop ( arr [ 0 ] , init ) #NEWLINE# #UNINDENT# self . _kernels [ ( nbtype , 0 ) ] = reduction_finish #NEWLINE# self . _cached_types . add ( nbtype )