Simple convenience function so that the a trous kernels can be easily accessed by any function.

def gpu_a_trous():
    """
    Simple convenience function so that the a trous kernels can be easily accessed by any function.
    """

    ker1 = SourceModule("""
                        __global__ void gpu_a_trous_row_kernel(float *in1, float *in2, float *wfil, int *scale)
                        {
                            const int len = gridDim.x*blockDim.x;
                            const int col = (blockDim.x * blockIdx.x + threadIdx.x);
                            const int i = col;
                            const int row = (blockDim.y * blockIdx.y + threadIdx.y);
                            const int j = row*len;
                            const int tid2 = i + j;
                            const int lstp = exp2(float(scale[0] + 1));
                            const int sstp = exp2(float(scale[0]));

                            in2[tid2] = wfil[2]*in1[tid2];

                            if (row < lstp)
                                { in2[tid2] += wfil[0]*in1[col + len*(lstp - row - 1)]; }
                            else
                                { in2[tid2] += wfil[0]*in1[tid2 - lstp*len]; }

                            if (row < sstp)
                                { in2[tid2] += wfil[1]*in1[col + len*(sstp - row - 1)]; }
                            else
                                { in2[tid2] += wfil[1]*in1[tid2 - sstp*len]; }

                            if (row >= (len - sstp))
                                { in2[tid2] += wfil[3]*in1[col + len*(2*len - row - sstp - 1)]; }
                            else
                                { in2[tid2] += wfil[3]*in1[tid2 + sstp*len]; }

                            if (row >= (len - lstp))
                                { in2[tid2] += wfil[4]*in1[col + len*(2*len - row - lstp - 1)]; }
                            else
                                { in2[tid2] += wfil[4]*in1[tid2 + lstp*len]; }
                        }
                        """, keep=True)

    ker2 = SourceModule("""
                        __global__ void gpu_a_trous_col_kernel(float *in1, float *in2, float *wfil, int *scale)
                        {
                            const int len = gridDim.x*blockDim.x;
                            const int col = (blockDim.x * blockIdx.x + threadIdx.x);
                            const int i = col;
                            const int row = (blockDim.y * blockIdx.y + threadIdx.y);
                            const int j = row*len;
                            const int tid2 = i + j;
                            const int lstp = exp2(float(scale[0] + 1));
                            const int sstp = exp2(float(scale[0]));

                            in2[tid2] = wfil[2]*in1[tid2];

                            if (col < lstp)
                                { in2[tid2] += wfil[0]*in1[j - col + lstp - 1]; }
                            else
                                { in2[tid2] += wfil[0]*in1[tid2 - lstp]; }

                            if (col < sstp)
                                { in2[tid2] += wfil[1]*in1[j - col + sstp - 1]; }
                            else
                                { in2[tid2] += wfil[1]*in1[tid2 - sstp]; }

                            if (col >= (len - sstp))
                                { in2[tid2] += wfil[3]*in1[j + 2*len - sstp - col - 1]; }
                            else
                                { in2[tid2] += wfil[3]*in1[tid2 + sstp]; }

                            if (col >= (len - lstp))
                                { in2[tid2] += wfil[4]*in1[j + 2*len - lstp - col - 1]; }
                            else
                                { in2[tid2] += wfil[4]*in1[tid2 + lstp]; }

                        }
                       """, keep=True)

    return ker1.get_function("gpu_a_trous_row_kernel"), ker2.get_function("gpu_a_trous_col_kernel")