 //}
 
 //-----------------------------------------------------------------------------
+__global__ void k_initRofSolver(float* d_u, float* d_u_prev, size_t stride_u,
+                                float2* d_p, size_t stride_p,
+                                imp::cu::Texture2D f_tex,
+                                size_t width, size_t height)
+{
+  int x = blockIdx.x*blockDim.x + threadIdx.x;
+  int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+  if (x<width && y<height)
+  {
+    float val = tex2D<float>(f_tex, x+.5f, y+.5f);
+    d_u[y*stride_u + x] = val;
+    d_u_prev[y*stride_u + x] = val;
+    d_p[y*stride_p + x] = make_float2(0.0f, 0.0f);
+  }
+}
+
+//-----------------------------------------------------------------------------
 __global__ void k_solveRofPrimalIteration(
-    float* d_u, float* d_u_prev, size_t stride_8uC1,
+    float* d_u, float* d_u_prev, size_t stride_u,
     Texture2D f_tex, Texture2D u_tex, Texture2D p_tex,
     float lambda, float tau, float theta, size_t width, size_t height)
 {

}

//-----------------------------------------------------------------------------
__global__ void k_convertResult8uC1(unsigned char* u, size_t stride_u,
                                    imp::cu::Texture2D u_tex,
                                    size_t width, size_t height)
{

                                        cudaFilterModeLinear,
                                        cudaAddressModeClamp,
                                        cudaReadModeElementType);

    // init internal vars
    k_initRofSolver
        <<< fragmentation_->dimGrid, fragmentation_->dimBlock >>> (
        reinterpret_cast<float*>(this->u_->data()),
        reinterpret_cast<float*>(this->u_prev_->data()), this->u_->stride(),
        reinterpret_cast<float2*>(this->p_->data()), this->p_->stride(),
        *this->f_tex_, this->size_.width(), this->size_.height());

    // internal params
    float L = sqrtf(8.0f);
