 __global__ void 
 gpujpeg_preprocessor_raw_to_comp_kernel_4_4_4(struct gpujpeg_preprocessor_data data, const uint8_t* d_data_raw, int image_width, int image_height)
 {
-    const unsigned int image_position_x = threadIdx.x + blockIdx.x * blockDim.x;
-    const unsigned int image_position_y = blockIdx.y;
-    const unsigned int x = threadIdx.x;
-    const unsigned int pix_offset = image_position_y * image_width + blockIdx.x * blockDim.x;
+    int x  = threadIdx.x;
+    int gX = (blockIdx.y * gridDim.x + blockIdx.x) * blockDim.x;
             
     // Load to shared
     __shared__ unsigned char s_data[RGB_8BIT_THREADS * 3];
     if ( (x * 4) < RGB_8BIT_THREADS * 3 ) {
         int* s = (int*)d_data_raw;
         int* d = (int*)s_data;
-        d[x] = s[((pix_offset * 3) >> 2) + x];
+        d[x] = s[((gX * 3) >> 2) + x];
     }
     __syncthreads();
 

__global__ void 
gpujpeg_preprocessor_raw_to_comp_kernel_4_2_2(struct gpujpeg_preprocessor_data data, const uint8_t* d_data_raw, int image_width, int image_height)
{
    const unsigned int image_position_x = threadIdx.x + blockIdx.x * blockDim.x;
    const unsigned int image_position_y = blockIdx.y;
    const unsigned int x = threadIdx.x;
    const unsigned int pix_offset = image_position_y * image_width + blockIdx.x * blockDim.x;
        
    // Load to shared
    __shared__ unsigned char s_data[RGB_8BIT_THREADS * 2];
    if ( (x * 4) < RGB_8BIT_THREADS * 2 ) {
        int* s = (int*)d_data_raw;
        int* d = (int*)s_data;
        d[x] = s[((pix_offset * 2) >> 2) + x];
    }
    __syncthreads();


    assert(coder->param_image.comp_count == 3);

    cudaMemset(coder->d_data, 0, coder->data_size * sizeof(uint8_t));
    
    // Select kernel
    gpujpeg_preprocessor_encode_kernel kernel = (gpujpeg_preprocessor_encode_kernel)coder->preprocessor;
    assert(kernel != NULL);
        
    int image_width = coder->param_image.width;
    int image_height = coder->param_image.height;
    
