 #include <stdint.h>
 #include <string.h>
 #include <limits.h>
-#include <cuda.h>
+#include <unistd.h>
 
 #include "common.h"
-#include "lzss_gpu_help.h"
+#include "lzss_help.h"
 
 #define DISP_BITS 11
 #define LEN_BITS 4

#define IDX_BY_BIT(arr,idx) ((arr)[(idx)/8])
#define GET_BIT(arr,idx) ((IDX_BY_BIT(arr,idx) >> (7-((idx)%8))) & 0x1)

#define END (min((int64_t)(input_len-partition_off),(int64_t)partition_size))
/* Make sure flags is zeroed out before passed in */
__global__ void gpu_compress(uint8_t* input, uint64_t input_len, uint8_t* dst, uint8_t* flags, comp_size_t *size)
{
  /* threadIdx.x; */
  uint64_t global_idx = blockIdx.x * blockDim.x + threadIdx.x;

  uint64_t partition_size = (global_idx * MAX_MATCH * 2 >= input_len) ?
    0 : min(input_len - global_idx * MAX_MATCH * 2, MAX_MATCH * 2);
  int64_t partition_off = (global_idx * MAX_MATCH * 2);

  input = input + partition_off;
  dst = dst + partition_off;
  flags = flags + partition_off;
  size = size + global_idx;

  match_expanded_t match;
  match_t m;
  uint64_t i=0;
  uint64_t w=0;
  uint64_t b=0;
  uint8_t *curr;
  /*
  printf("thread ID: %ld, last idx: %ld\n", global_idx + 100000, min((int64_t)(input_len-partition_off),(int64_t)partition_size));
  return;
  */
  for(;i<END && partition_size > 0;)
  {
    curr = input + i;

    uint64_t window_offset = min(partition_off + i,WINDOW);
    window_match(curr, window_offset, min((uint16_t) input_len-partition_off-i, (uint16_t) partition_size-i), &match);
    if( match.l < MIN_MATCH )
    {
      /* add 0 bit and the byte */
        /* IDX_BY_BIT(flags,b+j) |= PUT_BIT(0,b+j); */
      flags[b] = 0;
      dst[w] = curr[0];
      b += max(1,match.l);
      w += max(1,match.l);
      i += max(1,match.l);

    {
      /* match.d is displacement */
      /* match.l is length of match */
      IDX_BY_BIT(flags,b) |= PUT_BIT(1,b);
      pack_match(&m,&match);
      memcpy(dst + w,&m,sizeof(match_t));
      i += match.l;
      w += sizeof(match_t);
      b++;
