// Autogenerated by WaveGen.py, do not edit! //
#include <algorithm>
#include <cassert>
#include "common.h"
#include "wavelet_common.h"

/// Boundaries (depends on wavelet)
/// This much is reserved at the sides of the signal
/// Must be even!
#define BLEFT 2
#define BRIGHT 2

/// Initial shift (to keep precision in integer wavelets)
#define INITIAL_SHIFT 1
#define INITIAL_OFFSET 1

#define STAGE1_OFFSET 7
#define STAGE1_SHIFT 4

#define STAGE2_OFFSET 16
#define STAGE2_SHIFT 5

/// Vertical pass row management
#define RLEFT 3
#define RRIGHT 2
#define COPYROWS 3

static __global__ void s_transform_h( DATATYPE* data, int width, int stride )
{
    extern __shared__ DATATYPE shared[];  

    const int bid = blockIdx.x;    // row
    const int tid = threadIdx.x;   // thread id within row
    const int tidu16 = ((tid&16)>>4)|((tid&15)<<1)|(tid&~31);
    
    data += __mul24(bid, stride);

    int half = BLEFT+(width>>1)+BRIGHT;

    int ofs;
    if(width&3) // If width is not a multiple of 4, we need to use the slower method
    {
        /// Left part (even coefficients)
        /// Right part (odd coefficients)
        int w2 = (width>>1);
        uint16_t *row = (uint16_t*)data;
        uint16_t *row2 = (uint16_t*)&data[w2];
        uint16_t *dest1 = (uint16_t*)&shared[BLEFT];
        uint16_t *dest2 = (uint16_t*)&shared[half+BLEFT];
        
        ofs = tid;
        while(true)
        {
            if(ofs<w2)
            {
                dest1[ofs] = row[ofs];
                dest2[ofs] = row2[ofs];   
                ofs += BSH;
            } else break;
            if(ofs<w2)
            {
                dest1[ofs] = row[ofs];
                dest2[ofs] = row2[ofs];   
                ofs += BSH;
            } else break;
            if(ofs<w2)
            {
                dest1[ofs] = row[ofs];
                dest2[ofs] = row2[ofs];   
                ofs += BSH;
            } else break;
        }
    } 
    else
    {
        // Shared memory output offset for this thread
        int w2 = (width>>2);
        uint32_t *row = (uint32_t*)data;
        uint32_t *row2 = (uint32_t*)&data[width>>1];
        uint32_t *dest1 = (uint32_t*)&shared[BLEFT];
        uint32_t *dest2 = (uint32_t*)&shared[half+BLEFT];

        /// Left part (even coefficients)
        ofs = tid;
        while(true)
        {
            if(ofs<w2)
            {
                dest1[ofs] = row[ofs];
                dest2[ofs] = row2[ofs];
                ofs += BSH;
            } else break;
            if(ofs<w2)
            {
                dest1[ofs] = row[ofs];
                dest2[ofs] = row2[ofs];
                ofs += BSH;
            } else break;
            if(ofs<w2)
            {
                dest1[ofs] = row[ofs];
                dest2[ofs] = row2[ofs];
                ofs += BSH;
            } else break;
        }
    }

    __syncthreads();

    if(tidu16==0)
    {
        shared[half+BLEFT-2] = shared[half+BLEFT];
        shared[half+BLEFT-1] = shared[half+BLEFT];
        shared[half+BLEFT+(width>>1)] = shared[half+BLEFT+(width>>1)-1];
    }

    __syncthreads();
    
    // Now apply wavelet lifting to entire line at once
    // Process odd
    const int end = BLEFT+(width>>1);
    for(ofs = BLEFT+tidu16; ofs < end; ofs += BSH)
    {
        int acc = STAGE2_OFFSET;

        acc += -1*shared[half+ofs-2];
        acc += __mul24(9,shared[half+ofs-1]);
        acc += __mul24(9,shared[half+ofs+0]);
        acc += -1*shared[half+ofs+1];
        
        shared[ofs] -= acc >> STAGE2_SHIFT;
    }


    __syncthreads();

    if(tidu16==0)
    {    
        shared[BLEFT-1] = shared[BLEFT];
        shared[BLEFT+(width>>1)] = shared[BLEFT+(width>>1)-1];
        shared[BLEFT+(width>>1)+1] = shared[BLEFT+(width>>1)-1];
    }

    __syncthreads();
    
    // Process even
    for(ofs = BLEFT+tidu16; ofs < end; ofs += BSH)
    {
        int acc = STAGE1_OFFSET;

        acc += 1*shared[ofs-1];
        acc += __mul24(-9, shared[ofs+0]);
        acc += __mul24(-9, shared[ofs+1]);
        acc += 1*shared[ofs+2];
        
        shared[ofs + half] -= acc >> STAGE1_SHIFT;
    }

    __syncthreads();

    if(width&3)
    {
        uint32_t *row = (uint32_t*)data;
        int16_t *src1 = (int16_t*)&shared[BLEFT];
        int16_t *src2 = (int16_t*)&shared[half+BLEFT];
        int w2 = (width>>1);
        ofs = tid;
        while(true)
        {
            if(ofs<w2)
            {
                int a = (src1[ofs] + INITIAL_OFFSET)>>INITIAL_SHIFT;
                int b = (src2[ofs] + INITIAL_OFFSET)>>INITIAL_SHIFT;

                row[ofs] = (a&0xFFFF)|((b&0xFFFF)<<16);
                
                ofs += BSH;
            } else break;
            if(ofs<w2)
            {
                int a = (src1[ofs] + INITIAL_OFFSET)>>INITIAL_SHIFT;
                int b = (src2[ofs] + INITIAL_OFFSET)>>INITIAL_SHIFT;

                row[ofs] = (a&0xFFFF)|((b&0xFFFF)<<16);
                
                ofs += BSH;
            } else break;            
            if(ofs<w2)
            {
                int a = (src1[ofs] + INITIAL_OFFSET)>>INITIAL_SHIFT;
                int b = (src2[ofs] + INITIAL_OFFSET)>>INITIAL_SHIFT;

                row[ofs] = (a&0xFFFF)|((b&0xFFFF)<<16);
                
                ofs += BSH;
            } else break;
        }

    }
    else
    {
        i16_4 *row = (i16_4*)data;
        i16_2 *src1 = (i16_2*)&shared[BLEFT];
        i16_2 *src2 = (i16_2*)&shared[half+BLEFT];
        int w2 = (width>>2);
        ofs = tid;
        while(true)
        {
            if(ofs < w2)
            {
                i16_4 x;
                i16_2 a = src1[ofs];
                x.a = (a.a + INITIAL_OFFSET)>>INITIAL_SHIFT;
                x.c = (a.b + INITIAL_OFFSET)>>INITIAL_SHIFT;
                i16_2 b = src2[ofs];
                x.b = (b.a + INITIAL_OFFSET)>>INITIAL_SHIFT;
                x.d = (b.b + INITIAL_OFFSET)>>INITIAL_SHIFT;

                row[ofs] = x;
                
                ofs += BSH;
            } else break;
            if(ofs < w2)
            {
                i16_4 x;
                i16_2 a = src1[ofs];
                x.a = (a.a + INITIAL_OFFSET)>>INITIAL_SHIFT;
                x.c = (a.b + INITIAL_OFFSET)>>INITIAL_SHIFT;
                i16_2 b = src2[ofs];
                x.b = (b.a + INITIAL_OFFSET)>>INITIAL_SHIFT;
                x.d = (b.b + INITIAL_OFFSET)>>INITIAL_SHIFT;

                row[ofs] = x;
                
                ofs += BSH;
            } else break;            
            if(ofs < w2)
            {
                i16_4 x;
                i16_2 a = src1[ofs];
                x.a = (a.a + INITIAL_OFFSET)>>INITIAL_SHIFT;
                x.c = (a.b + INITIAL_OFFSET)>>INITIAL_SHIFT;
                i16_2 b = src2[ofs];
                x.b = (b.a + INITIAL_OFFSET)>>INITIAL_SHIFT;
                x.d = (b.b + INITIAL_OFFSET)>>INITIAL_SHIFT;

                row[ofs] = x;
                
                ofs += BSH;
            } else break;
        }
    }
}

#define BROWS (2*BSVY+COPYROWS) /* Rows to process at once */
#define SKIPTOP COPYROWS

#define PAD_ROWS (WRITEBACK-SKIPTOP+RRIGHT+COPYROWS) /* Rows below which to use s_transform_v_pad */
/// tid is BCOLSxBROWS matrix
/// RLEFT+BROWS+RRIGHT rows
#define TOTALROWS (RLEFT+BROWS+RRIGHT)
#define OVERLAP (RLEFT+RRIGHT+COPYROWS)
#define OVERLAP_OFFSET (TOTALROWS-OVERLAP)
#define WRITEBACK (2*BSVY)

__device__ void doTransform(int xofs)
{
    const int tidx = (threadIdx.x<<1)+xofs;   // column
    const int tidy = threadIdx.y;   // row

    extern __shared__ DATATYPE shared[];
    int ofs;

    ofs = ((RLEFT+(tidy<<1)+3)<<BCOLS_SHIFT) + tidx;
    {
        int acc = STAGE2_OFFSET;

        acc += -1*shared[ofs-3*BCOLS];
        acc += __mul24(9,shared[ofs-BCOLS]);
        acc += __mul24(9,shared[ofs+BCOLS]);
        acc += -1*shared[ofs+3*BCOLS];
        
        shared[ofs] -= acc >> STAGE2_SHIFT;
    }
    
    __syncthreads();
    
    ofs -= BCOLS*3;
    
    {
        int acc = STAGE1_OFFSET;

        acc += 1*shared[ofs-3*BCOLS];
        acc += __mul24(-9, shared[ofs-BCOLS]);
        acc += __mul24(-9, shared[ofs+BCOLS]);
        acc += 1*shared[ofs+3*BCOLS];

        shared[ofs] -= acc >> STAGE1_SHIFT;
    }

}

__device__ void doTransformTB(int xofs, unsigned int leftover)
{
    const int tidx = (threadIdx.x<<1)+xofs;   // column
    const int tidy = threadIdx.y;   // row
    const int minn = (RLEFT<<BCOLS_SHIFT) + tidx;
    const int maxx = leftover-(2<<BCOLS_SHIFT) + tidx;
    
    extern __shared__ DATATYPE shared[];
    int ofs;

    /// Do the transform
    /// Do procesing on shared mem
    /// Treat all columns the same
    
    /// Process even rows
    ofs = ((RLEFT+(tidy<<1))<<BCOLS_SHIFT) + tidx;
    {
        int acc = STAGE2_OFFSET;

        acc += -1*shared[max(ofs-3*BCOLS, minn+BCOLS)];
        acc += __mul24(9,shared[max(ofs-BCOLS, minn+BCOLS)]);
        acc += __mul24(9,shared[ofs+BCOLS]);
        acc += -1*shared[min(ofs+3*BCOLS, maxx+BCOLS)];

        shared[ofs] -= acc >> STAGE2_SHIFT;
    }
    
    __syncthreads();
    
    /// Process odd rows
    ofs += BCOLS;
    {
        int acc = STAGE1_OFFSET;

        acc += 1*shared[max(ofs-3*BCOLS,minn)];
        acc += __mul24(-9, shared[ofs-BCOLS]);
        acc += __mul24(-9, shared[min(ofs+BCOLS,maxx)]);
        acc += 1*shared[min(ofs+3*BCOLS,maxx)];

        shared[ofs] -= acc >> STAGE1_SHIFT;
    }

}

__device__ void doTransformT(int xofs)
{
    const int tidx = (threadIdx.x<<1)+xofs;   // column
    const int tidy = threadIdx.y;   // row
    const int minn = ((RLEFT+SKIPTOP)<<BCOLS_SHIFT) + tidx;
    
    extern __shared__ DATATYPE shared[];
    int ofs;

    /// Do the transform
    
    /// Do procesing on shared mem
    /// Treat all columns the same
    ofs = ((SKIPTOP+RLEFT+(tidy<<1))<<BCOLS_SHIFT) + tidx;
    
    /// Process even rows
    {
        int acc = STAGE2_OFFSET;

        acc += -1*shared[max(ofs-3*BCOLS, minn+BCOLS)];
        acc += __mul24(9,shared[max(ofs-BCOLS, minn+BCOLS)]);
        acc += __mul24(9,shared[ofs+BCOLS]);
        acc += -1*shared[ofs+3*BCOLS];
        
        shared[ofs] -= acc >> STAGE2_SHIFT;
    }

    __syncthreads();

    /// Process odd rows, except for last
    ofs += BCOLS;
    if(tidy < (BSVY-2))
    {
        int acc = STAGE1_OFFSET;

        acc += 1*shared[max(ofs-3*BCOLS,minn)];
        acc += __mul24(-9, shared[ofs-BCOLS]);
        acc += __mul24(-9, shared[ofs+BCOLS]);
        acc += 1*shared[ofs+3*BCOLS];
        
        shared[ofs] -= acc >> STAGE1_SHIFT;
    }

}

// Process leftover
__device__ void doTransformB(int xofs, unsigned int leftover)
{
    const int tidx = (threadIdx.x<<1)+xofs;   // column
    const int tidy = threadIdx.y;   // row
    const int maxx = leftover-(2<<BCOLS_SHIFT) + tidx;
    
    extern __shared__ DATATYPE shared[];
    int ofs, ofs_t;

    ofs = ((RLEFT+(tidy<<1))<<BCOLS_SHIFT) + tidx;

    for(ofs_t=ofs+3*BCOLS; ofs_t<leftover; ofs_t += BSVY*2*BCOLS)
    {
        int acc = STAGE2_OFFSET;

        acc += -1*shared[ofs_t-3*BCOLS];
        acc += __mul24(9,shared[ofs_t-BCOLS]);
        acc += __mul24(9,shared[ofs_t+BCOLS]);
        acc += -1*shared[min(ofs_t+3*BCOLS, maxx+BCOLS)];
        
        shared[ofs_t] -= acc >> STAGE2_SHIFT;
    }
    
    __syncthreads();
    
    for(ofs_t=ofs; ofs_t<leftover; ofs_t += BSVY*2*BCOLS)
    {
        int acc = STAGE1_OFFSET;

        acc += 1*shared[ofs_t-3*BCOLS];
        acc += __mul24(-9, shared[ofs_t-BCOLS]);
        acc += __mul24(-9, shared[min(ofs_t+BCOLS,maxx)]);
        acc += 1*shared[min(ofs_t+3*BCOLS,maxx)];
        
        shared[ofs_t] -= acc >> STAGE1_SHIFT;
    }
}



#if 0
// Rolled
#define READ_LOOP(rows)         for(; sofs < (rows); sofs += (BCOLS*BSVY), gofs += istride)             *((uint32_t*)&shared[sofs]) = *((uint32_t*)&data[gofs]);

#define WRITE_LOOP(rows)         for(; sofs < (rows); sofs += (BCOLS*BSVY), gofs += istride)             *((uint32_t*)&data[gofs]) = *((uint32_t*)&shared[sofs]);
#endif
#if 1
// Unrolled
#define READ_LOOP_ENTRY(rows)         if(sofs < (rows)) {             *((uint32_t*)&shared[sofs]) = *((uint32_t*)&data[gofs]);             sofs += (BCOLS*BSVY); gofs += istride;         } else break;
#define READ_LOOP(rows)         while(1) {             READ_LOOP_ENTRY(rows);             READ_LOOP_ENTRY(rows);         }

#define WRITE_LOOP_ENTRY(rows)         if(sofs < (rows)) {             *((uint32_t*)&data[gofs]) = *((uint32_t*)&shared[sofs]);             sofs += (BCOLS*BSVY); gofs += istride;         } else break;
#define WRITE_LOOP(rows)         while(1) {             WRITE_LOOP_ENTRY(rows);             WRITE_LOOP_ENTRY(rows);         }

#endif

#if 0
// Unrolled
#define READ_LOOP_ENTRY(rows)         if(sofs < (rows)) {             *((uint32_t*)&shared[sofs]) = *((uint32_t*)&data[gofs]);             sofs += (BCOLS*BSVY); gofs += istride;         }
#define READ_LOOP(rows)         READ_LOOP_ENTRY(rows);         READ_LOOP_ENTRY(rows);         READ_LOOP_ENTRY(rows); 

#define WRITE_LOOP_ENTRY(rows)         if(sofs < (rows)) {             *((uint32_t*)&data[gofs]) = *((uint32_t*)&shared[sofs]);             sofs += (BCOLS*BSVY); gofs += istride;         }

#define WRITE_LOOP(rows)         WRITE_LOOP_ENTRY(rows);         WRITE_LOOP_ENTRY(rows);          WRITE_LOOP_ENTRY(rows); 

#endif

static __global__ void s_transform_v( DATATYPE* data, int width, int height, int stride )
{
    extern __shared__ DATATYPE shared[];  

    const unsigned int bid = blockIdx.x;    // slab (BCOLS columns)
    const unsigned int tidx = threadIdx.x<<1;   // column
    const unsigned int tidy = threadIdx.y;   // row    
    const unsigned int swidth = min(width-(bid<<BCOLS_SHIFT), BCOLS); // Width of this slab, usually BCOLS but can be less

    // Element offset in global memory
    //int idata = tidx + (bid<<BCOLS_SHIFT) + __mul24(tidy, stride);
    data += tidx + (bid<<BCOLS_SHIFT) + __mul24(tidy, stride);
    
    const unsigned int istride = __mul24(BSVY, stride);
    const unsigned int sdata = tidx + (tidy<<BCOLS_SHIFT);
    // First read BROWS+RRIGHT
    // After that BROWS
    unsigned int ref = height-(WRITEBACK-SKIPTOP+RRIGHT+COPYROWS);
    unsigned int blocks = ref/WRITEBACK;
    unsigned int leftover = (RLEFT+COPYROWS+RRIGHT+(ref%WRITEBACK))<<BCOLS_SHIFT;
    
    unsigned int gofs,sofs;

    /// More than one block
    /// Read first block of BROWS+RRIGHT rows
    /// Upper RLEFT rows are left unitialized for now, later they should be copied from top
    if(tidx < swidth)
    {
        gofs = 0;
        sofs = sdata + ((RLEFT+SKIPTOP)<<BCOLS_SHIFT);
        READ_LOOP(TOTALROWS<<BCOLS_SHIFT);
    }
// idata_read = idata_write + __mul24(TOTALROWS-RLEFT, stride)

    __syncthreads();
    
    doTransformT(0);
    doTransformT(1);
    
    __syncthreads();

    /// Write back WRITEBACK rows
    if(tidx < swidth)
    {
        gofs = 0;
        sofs = sdata + ((RLEFT+SKIPTOP)<<BCOLS_SHIFT);
        WRITE_LOOP((WRITEBACK+RLEFT)<<BCOLS_SHIFT);
    }
// idata_read = idata_write + __mul24((BROWS+RRIGHT)-WRITEBACK, stride)

// Difference between global mem read and write pointer
#define DATA_READ_DIFF __mul24((BROWS+RRIGHT)-WRITEBACK, stride)
// Advance pointer with this amount after each block
#define DATA_INC __mul24(WRITEBACK, stride)

    data += __mul24(WRITEBACK-SKIPTOP, stride);
    for(unsigned int block=0; block<blocks; ++block)
    {
        __syncthreads();
        /// Move lower rows to top rows
#if OVERLAP <= BSVY 
        if(tidy < OVERLAP)
        {
            unsigned int l = (tidy<<BCOLS_SHIFT)+tidx;
            *((uint32_t*)&shared[l]) = *((uint32_t*)&shared[(WRITEBACK<<BCOLS_SHIFT)+l]);
        }
#else
        for(sofs = (tidy<<BCOLS_SHIFT)+tidx; sofs < (OVERLAP<<BCOLS_SHIFT); sofs += (BSVY<<BCOLS_SHIFT))
            *((uint32_t*)&shared[sofs]) = *((uint32_t*)&shared[(WRITEBACK<<BCOLS_SHIFT)+sofs]);
#endif                
        
        /// Fill shared memory -- read next block of BROWS rows
        /// We can skip RRIGHT rows as we've already copied them for the previous block
        /// and moved them to the top
        if(tidx < swidth)
        {
            gofs = DATA_READ_DIFF;
            sofs = sdata + (OVERLAP<<BCOLS_SHIFT);
            READ_LOOP(TOTALROWS<<BCOLS_SHIFT);
        }

        __syncthreads();

        doTransform(0);
        doTransform(1);

        __syncthreads();

        /// Write back BROWS rows
        if(tidx < swidth)
        {
            gofs = 0;
            sofs = sdata + (RLEFT<<BCOLS_SHIFT);
            WRITE_LOOP((WRITEBACK+RLEFT)<<BCOLS_SHIFT);
        }
        
        data += DATA_INC;
    }
    __syncthreads();

    ///
    /// Handle partial last block
    /// Move lower rows to top rows
#if OVERLAP <= BSVY 
        if(tidy < OVERLAP)
        {
            unsigned int l = (tidy<<BCOLS_SHIFT)+tidx;
            *((uint32_t*)&shared[l]) = *((uint32_t*)&shared[(WRITEBACK<<BCOLS_SHIFT)+l]);
        }
#else
        for(sofs = (tidy<<BCOLS_SHIFT)+tidx; sofs < (OVERLAP<<BCOLS_SHIFT); sofs += (BSVY<<BCOLS_SHIFT))
            *((uint32_t*)&shared[sofs]) = *((uint32_t*)&shared[(WRITEBACK<<BCOLS_SHIFT)+sofs]);
#endif                
    
    /// Fill shared memory -- read next block of BROWS rows
    /// We can skip RRIGHT rows as we've already copied them for the previous block
    /// and moved them to the top
    if(tidx < swidth)
    {
        gofs = DATA_READ_DIFF;
        sofs = sdata + (OVERLAP<<BCOLS_SHIFT);
        READ_LOOP(leftover);
    }

    __syncthreads();

    doTransformB(0, leftover);
    doTransformB(1, leftover);
    
    __syncthreads();
    
    /// Write back leftover
    if(tidx < swidth)
    {
        gofs = 0;
        sofs = sdata + (RLEFT<<BCOLS_SHIFT);
        WRITE_LOOP(leftover);
    }

}

/// Use this if the image is lower than PAD_ROWS
static __global__ void s_transform_v_pad( DATATYPE* data, int width, int height, int stride )
{
    extern __shared__ DATATYPE shared[];  

    const unsigned int bid = blockIdx.x;    // slab (BCOLS columns)
    const unsigned int tidx = threadIdx.x<<1;   // column
    const unsigned int tidy = threadIdx.y;   // row
    const unsigned int swidth = min(width-(bid<<BCOLS_SHIFT), BCOLS); // Width of this slab, usually BCOLS but can be less

    // Element offset in global memory
    //int idata = tidx + (bid<<BCOLS_SHIFT) + __mul24(tidy, stride);
    data +=  tidx + (bid<<BCOLS_SHIFT) + __mul24(tidy, stride);
    const unsigned int istride = __mul24(BSVY, stride);
    const unsigned int sdata = tidx + ((tidy+RLEFT)<<BCOLS_SHIFT); // Does this get converted into a shift?
    // First read BROWS+RRIGHT
    // After that BROWS
    unsigned int leftover = (RLEFT+height) << BCOLS_SHIFT; /// How far to fill buffer on last read
    //unsigned int blocks = (height-RRIGHT)/BROWS;
    
    unsigned int gofs, sofs;
    
    /// Fill shared memory -- read next block of BROWS rows
    /// We can skip RRIGHT rows as we've already copied them for the previous block
    /// and moved them to the top
    if(tidx < swidth)
    {
        gofs = 0; // Read from row (cur+RRIGHT)
        sofs = sdata;
        READ_LOOP(leftover);
    }

    __syncthreads();
    
    doTransformTB(0, leftover);
    doTransformTB(1, leftover);
    
    __syncthreads();
    
    /// Write back leftover
    if(tidx < swidth)
    {
        gofs = 0;
        sofs = sdata;
        WRITE_LOOP(leftover);
    }
}

void cuda_iiwt_13_5(int16_t *d_data, int lwidth, int lheight, int stride, cudaStream_t stream)
{
    /** Invoke kernel */
    dim3 block_size;
    dim3 grid_size;
    int shared_size;

#ifdef VERTICAL
    block_size.x = BSVX;
    block_size.y = BSVY;
    block_size.z = 1;
    grid_size.x = (lwidth+BCOLS-1)/BCOLS;
    grid_size.y = 1;
    grid_size.z = 1;
    shared_size = BCOLS*(BROWS+RLEFT+RRIGHT)*2; 
	
    if(lheight < PAD_ROWS)
        s_transform_v_pad<<<grid_size, block_size, shared_size, stream>>>(d_data, lwidth, lheight, stride);
    else
        s_transform_v<<<grid_size, block_size, shared_size, stream>>>(d_data, lwidth, lheight, stride);
#endif
#ifdef HORIZONTAL
    block_size.x = BSH;
    block_size.y = 1;
    block_size.z = 1;
    grid_size.x = lheight;
    grid_size.y = 1;
    grid_size.z = 1;
    shared_size = (lwidth+BLEFT*2+BRIGHT*2) * sizeof(DATATYPE); 
    s_transform_h<<<grid_size, block_size, shared_size, stream>>>(d_data, lwidth, stride);
#endif
}