<commit_before>#pragma once

namespace dg
{

// general multiply kernel
template<class value_type>
 __global__ void ell_multiply_kernel(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols, const int blocks_per_line,
         const int n, const int size,
         const int right, 
         const value_type* x, value_type *y
         )
{
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rr = row/right, rrn = rr/n;
        int s=rrn/num_rows, 
            i = (rrn)%num_rows, 
            k = (rr)%n, 
            j=row%right;
        int B, J;
        value_type temp=0;
        for( int d=0; d<blocks_per_line; d++)
        {
            B = (data_idx[i*blocks_per_line+d]*n+k)*n;
            J = (s*num_cols+cols_idx[i*blocks_per_line+d])*n;
            for( int q=0; q<n; q++) //multiplication-loop
                temp +=data[ B+q]* x[(J+q)*right+j];
            y[row]=temp;
        }
    }

}

// multiply kernel, n=3, 3 blocks per line
template<class value_type>
 __global__ void ell_multiply_kernel33(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols,
         const int size,
         const int right, 
         const value_type* x, value_type *y
         )
{
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rr = row/right, rrn = rr/3;
        int s=rrn/num_rows, 
            i = (rrn)%num_rows, 
            k = (rr)%3, 
            j=row%right;
        int B, J;
        value_type temp=0;
        {
            B = (data_idx[i*3+0]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+0])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            B = (data_idx[i*3+1]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+1])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            B = (data_idx[i*3+2]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+2])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            y[row]=temp;
        }
    }
}

// multiply kernel, n=3, 2 blocks per line
template<class value_type>
 __global__ void ell_multiply_kernel32(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols, 
         const int size,
         const int right, 
         const value_type* x, value_type *y
         )
{
    //int size = left*num_rows*n*right;
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rr = row/right, rrn = rr/3;
        int s=rrn/num_rows, 
            i = (rrn)%num_rows, 
            k = (rr)%3, 
            j=row%right;
        int B, J;
        value_type temp=0;
        {
            B = (data_idx[i*2+0]*3+k)*3;
            J = (s*num_cols+cols_idx[i*2+0])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            B = (data_idx[i*2+1]*3+k)*3;
            J = (s*num_cols+cols_idx[i*2+1])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            y[row]=temp;
        }
    }

}

// multiply kernel, n=3, 3 blocks per line, right = 1
template<class value_type>
 __global__ void ell_multiply_kernel33x(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols,
         const int size,
         const value_type* x, value_type *y
         )
{
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rr = row/1, rrn = rr/3;
        int s=rrn/num_rows, 
            i = (rrn)%num_rows, 
            k = (rr)%3, 
            j=row%1;
        int B, J;
        value_type temp=0;
        {
            B = (data_idx[i*3+0]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+0])*3;
            temp +=data[ B+0]* x[(J+0)*1+j];
            temp +=data[ B+1]* x[(J+1)*1+j];
            temp +=data[ B+2]* x[(J+2)*1+j];
            B = (data_idx[i*3+1]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+1])*3;
            temp +=data[ B+0]* x[(J+0)*1+j];
            temp +=data[ B+1]* x[(J+1)*1+j];
            temp +=data[ B+2]* x[(J+2)*1+j];
            B = (data_idx[i*3+2]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+2])*3;
            temp +=data[ B+0]* x[(J+0)*1+j];
            temp +=data[ B+1]* x[(J+1)*1+j];
            temp +=data[ B+2]* x[(J+2)*1+j];
            y[row]=temp;
        }
    }
}

// multiply kernel, n=3, 2 blocks per line, right = 1
template<class value_type>
 __global__ void ell_multiply_kernel32x(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols,
         const int size,
         const value_type* x, value_type *y
         )
{
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rrn = row/3, k = (row%3);
        int s=rrn/num_rows, i = (rrn)%num_rows;
        int B0,B1, J0, J1;
        value_type temp=0;
        {
            B0 = (data_idx[i*2+0]*3+k)*3;
            B1 = (data_idx[i*2+1]*3+k)*3;
            J0 = (s*num_cols+cols_idx[i*2+0])*3;
            J1 = (s*num_cols+cols_idx[i*2+1])*3;<commit_msg>performance on gpu is max?<commit_after>#pragma once

namespace dg
{

// general multiply kernel
template<class value_type>
 __global__ void ell_multiply_kernel(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols, const int blocks_per_line,
         const int n, const int size,
         const int right, 
         const value_type* x, value_type *y
         )
{
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rr = row/right, rrn = rr/n;
        int s=rrn/num_rows, 
            i = (rrn)%num_rows, 
            k = (rr)%n, 
            j=row%right;
        int B, J;
        value_type temp=0;
        for( int d=0; d<blocks_per_line; d++)
        {
            B = (data_idx[i*blocks_per_line+d]*n+k)*n;
            J = (s*num_cols+cols_idx[i*blocks_per_line+d])*n;
            for( int q=0; q<n; q++) //multiplication-loop
                temp +=data[ B+q]* x[(J+q)*right+j];
            y[row]=temp;
        }
    }

}

// multiply kernel, n=3, 3 blocks per line
template<class value_type>
 __global__ void ell_multiply_kernel33(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols,
         const int size,
         const int right, 
         const value_type* x, value_type *y
         )
{
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rr = row/right, rrn = rr/3;
        int s=rrn/num_rows, 
            i = (rrn)%num_rows, 
            k = (rr)%3, 
            j=row%right;
        int B, J;
        value_type temp=0;
        {
            B = (data_idx[i*3+0]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+0])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            B = (data_idx[i*3+1]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+1])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            B = (data_idx[i*3+2]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+2])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            y[row]=temp;
        }
    }
}

// multiply kernel, n=3, 2 blocks per line
template<class value_type>
 __global__ void ell_multiply_kernel32(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols, 
         const int size,
         const int right, 
         const value_type* x, value_type *y
         )
{
    //int size = left*num_rows*n*right;
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rr = row/right, rrn = rr/3;
        int s=rrn/num_rows, 
            i = (rrn)%num_rows, 
            k = (rr)%3, 
            j=row%right;
        int B, J;
        value_type temp=0;
        {
            B = (data_idx[i*2+0]*3+k)*3;
            J = (s*num_cols+cols_idx[i*2+0])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            B = (data_idx[i*2+1]*3+k)*3;
            J = (s*num_cols+cols_idx[i*2+1])*3;
            temp +=data[ B+0]* x[(J+0)*right+j];
            temp +=data[ B+1]* x[(J+1)*right+j];
            temp +=data[ B+2]* x[(J+2)*right+j];
            y[row]=temp;
        }
    }

}

// multiply kernel, n=3, 3 blocks per line, right = 1
template<class value_type>
 __global__ void ell_multiply_kernel33x(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols,
         const int size,
         const value_type* x, value_type *y
         )
{
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rr = row/1, rrn = rr/3;
        int s=rrn/num_rows, 
            i = (rrn)%num_rows, 
            k = (rr)%3;
        int B, J;
        value_type temp=0;
        {
            B = (data_idx[i*3+0]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+0])*3;
            temp +=data[ B+0]* x[(J+0)];
            temp +=data[ B+1]* x[(J+1)];
            temp +=data[ B+2]* x[(J+2)];
            B = (data_idx[i*3+1]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+1])*3;
            temp +=data[ B+0]* x[(J+0)];
            temp +=data[ B+1]* x[(J+1)];
            temp +=data[ B+2]* x[(J+2)];
            B = (data_idx[i*3+2]*3+k)*3;
            J = (s*num_cols+cols_idx[i*3+2])*3;
            temp +=data[ B+0]* x[(J+0)];
            temp +=data[ B+1]* x[(J+1)];
            temp +=data[ B+2]* x[(J+2)];
            y[row]=temp;
        }
    }
}

// multiply kernel, n=3, 2 blocks per line, right = 1
template<class value_type>
 __global__ void ell_multiply_kernel32x(
         const value_type* data, const int* cols_idx, const int* data_idx, 
         const int num_rows, const int num_cols,
         const int size,
         const value_type* x, value_type *y
         )
{
    const int thread_id = blockDim.x * blockIdx.x + threadIdx.x;
    const int grid_size = gridDim.x*blockDim.x;
    //every thread takes num_rows/grid_size rows
    for( int row = thread_id; row<size; row += grid_size)
    {
        int rrn = row/3, k = (row%3);
        int s=rrn/num_rows, i = (rrn)%num_rows;
        int B0,B1, J0, J1;
        value_type temp=0;
        {
            B0 = (data_idx[i*2+0]*3+k)*3;
            B1 = (data_idx[i*2+1]*3+k)*3;
            J0 = (s*num_cols+cols_idx[i*2+0])*3;
            J1 = (s*num_cols+cols_idx[i*2+1])*3;
{'fraction_non_alphanumeric': 0.15817967830521773, 'fraction_numerical': 0.025892506865437426, 'mean_word_length': 1.9745624270711786, 'pattern_counts': {'":': 0, '<': 27, '<?xml version=': 0, '>': 13, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '6336204', 'n_tokens_mistral': 4944, 'n_tokens_neox': 4859, 'n_words': 1100}