(dtype_str, nbins, offset)
return kernel
Build a kernel to compute a 64 bin histogram.

    Use templating to generate a customized kernel depending on the input data type.

    Memoized to avoid compiling the same kernel twice.
Build a kernel to compute a 64 bin histogram.
Build a kernel to compute a 64 bin histogram .
def _get_hist_kernel(dtype_str, nbins, offset):
    """
    Build a kernel to compute a 64 bin histogram.

    Use templating to generate a customized kernel depending on the input data type.

    Memoized to avoid compiling the same kernel twice.
    """
    type_str = _ew_types[dtype_str[1:]]
    from string import Template
    code = Template(_common_fp16_to_fp32 + r"""

#define MAX(a,b) (a > b ? a : b)
#define MIN(a,b) (a < b ? a : b)

__global__ void kernel_histo (
    int* d_hist, const $in_type* a1_in,
    int strides, int size)
{
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;

    __shared__ int s[$nbins];
    if(tid < $nbins){
        s[tid] = 0;
    }

    if(bid == 0 && tid < $nbins){
        d_hist[tid] = 0;
    }

    for (int i = tid + blockDim.x*bid; i < size; i += strides)
    {
        float a1 = $convert_to_float(__ldg(a1_in + i));

        float absval = fabs(a1);

        float logabs = round(log2f(absval));

        int bin = MIN($nbins-1, MAX(0, logabs-($offset)));

        atomicAdd(&s[bin], 1);

    }

    __syncthreads();

    if(tid < $nbins){
        atomicAdd(&d_hist[tid], s[tid]);
    }
}
""")

    module = SourceModule(code.substitute(in_type=type_str['type'],
                                          convert_to_float=type_str['cvt'],
                                          nbins=nbins,
                                          offset=offset),
                          options=[])
    kernel = module.get_function("kernel_histo")
    kernel.prepare("PPII")
    return kernel
def _get_hist_kernel ( dtype_str , nbins , offset ) : type_str = _ew_types [ dtype_str [ 1 : ] ] from string import Template code = Template ( _common_fp16_to_fp32 + r"""

#define MAX(a,b) (a > b ? a : b)
#define MIN(a,b) (a < b ? a : b)

__global__ void kernel_histo (
    int* d_hist, const $in_type* a1_in,
    int strides, int size)
{
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;

    __shared__ int s[$nbins];
    if(tid < $nbins){
        s[tid] = 0;
    }

    if(bid == 0 && tid < $nbins){
        d_hist[tid] = 0;
    }

    for (int i = tid + blockDim.x*bid; i < size; i += strides)
    {
        float a1 = $convert_to_float(__ldg(a1_in + i));

        float absval = fabs(a1);

        float logabs = round(log2f(absval));

        int bin = MIN($nbins-1, MAX(0, logabs-($offset)));

        atomicAdd(&s[bin], 1);

    }

    __syncthreads();

    if(tid < $nbins){
        atomicAdd(&d_hist[tid], s[tid]);
    }
}
""" ) module = SourceModule ( code . substitute ( in_type = type_str [ 'type' ] , convert_to_float = type_str [ 'cvt' ] , nbins = nbins , offset = offset ) , options = [ ] ) kernel = module . get_function ( "kernel_histo" ) kernel . prepare ( "PPII" ) return kernel