     // When loading 4:2:2 data of odd width, the data in fact has even width, so round it
     // (at least imagemagick convert tool generates data stream in this way)
     if ( coder->param_image.sampling_factor == GPUJPEG_4_2_2 )
-        image_width = gpujpeg_div_and_round_up(coder->param_image.width, 2) * 2;
+        image_width = (coder->param_image.width + 1) & ~1;
         
     // Prepare unit size
     assert(coder->param_image.sampling_factor == GPUJPEG_4_4_4 || coder->param_image.sampling_factor == GPUJPEG_4_2_2);
-    int unitSize = coder->param_image.sampling_factor == GPUJPEG_4_4_4 ? 3 : 2;
     
     // Prepare kernel
-    int alignedSize = gpujpeg_div_and_round_up(image_width * image_height, RGB_8BIT_THREADS) * RGB_8BIT_THREADS * unitSize;
-    dim3 threads (RGB_8BIT_THREADS);
-    dim3 grid (alignedSize / (RGB_8BIT_THREADS * unitSize));
-    assert(alignedSize % (RGB_8BIT_THREADS * unitSize) == 0);
-    if ( grid.x > GPUJPEG_CUDA_MAXIMUM_GRID_SIZE ) {
-        grid.y = gpujpeg_div_and_round_up(grid.x, GPUJPEG_CUDA_MAXIMUM_GRID_SIZE);
-        grid.x = GPUJPEG_CUDA_MAXIMUM_GRID_SIZE;
-    }
+    const dim3 threads (RGB_8BIT_THREADS);
+    const dim3 grid (gpujpeg_div_and_round_up(image_width, RGB_8BIT_THREADS), image_height);
 
     // Run kernel
     struct gpujpeg_preprocessor_data data;

    // Color transform
    gpujpeg_color_transform<color_space, color_space_internal>::perform(r1, r2, r3);
    
    // Position
    int image_position = gX + x;
    int image_position_x = image_position % image_width;
    int image_position_y = image_position / image_width;
        
    // Store
    if ( image_position < (image_width * image_height) ) {
        gpujpeg_preprocessor_raw_to_comp_store<s_comp1_samp_factor_h, s_comp1_samp_factor_v>::perform((uint8_t)r1, image_position_x, image_position_y, data.comp[0]);
        gpujpeg_preprocessor_raw_to_comp_store<s_comp2_samp_factor_h, s_comp2_samp_factor_v>::perform((uint8_t)r2, image_position_x, image_position_y, data.comp[1]);
        gpujpeg_preprocessor_raw_to_comp_store<s_comp3_samp_factor_h, s_comp3_samp_factor_v>::perform((uint8_t)r3, image_position_x, image_position_y, data.comp[2]);

__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;
    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 * 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];
    }
    __syncthreads();

