         // first thread uses DC table for its even value
         d_table_even = d_table_dc;
         
-        // update last DC coefficient
+        // update last DC coefficient (saved at the special place at the end of the shared bufer)
         const int original_in_even = in_even;
-        in_even -= *last_dc;
-        *last_dc = original_in_even;
+        in_even -= ((int*)s_out)[last_dc_idx];
+        ((int*)s_out)[last_dc_idx] = original_in_even;
     }
     
     // last thread handles special block-termination symbol

    int warpidx = threadIdx.x >> 5;
    int tid = threadIdx.x & 31;

    __shared__ uint4 s_out_all[64 * WARPS_NUM];
    unsigned int * s_out = (unsigned int*)(s_out_all + warpidx * 64);
    
    // Number of remaining codewords in shared buffer
    int remaining_codewords = 0;

        return;
    struct gpujpeg_segment* segment = &d_segment[segment_index];
    
    // Initialize last DC coefficients
    if(tid < 3) {
        s_out[256 + tid] = 0;
    }
    
    // Prepare data pointers
    unsigned int * data_compressed = (unsigned int*)(d_data_compressed + segment->data_temp_index);
