 gpujpeg_huffman_encoder_serialization_kernel(
     struct gpujpeg_segment* d_segment,
     int segment_count, 
-    uint8_t* d_data_compressed
+    const uint8_t* const d_src,
+    uint8_t* const d_dest
 ) {    
 #if __CUDA_ARCH__ >= 200
     // Temp buffer for all threads of the threadblock

        gpujpeg_huffman_encoder_serialization_kernel<<<num_serialization_tblocks, SERIALIZATION_THREADS_PER_TBLOCK>>>(
            coder->d_segment, 
            coder->segment_count, 
            coder->d_temp_huffman
        );
        cudaThreadSynchronize();

            gpujpeg_huffman_encoder_encode_kernel_warp<true><<<grid, thread>>>(
                coder->d_segment, 
                coder->segment_count, 
                coder->d_data_compressed,
                coder->d_block_list,
                coder->d_data_quantized,
                coder->d_component,
