 #include <stdio.h>
 #include "MatKernel.hpp"
 
-#ifdef __CUDA_ARCH__ 
 #if __CUDA_ARCH__ > 200
 
 #define edcellupdate(RR,RP1,RP2,RPP,WUN,TMP)                                                               \

  printf("__hammingdists() not defined for CUDA Arch < 300\n");
}
#endif
#else
__global__ void __veccmp(int *a, int *b, int *d) {
  printf("__veccmp() not defined for CUDA Arch < 300\n");
}

template<int VECLEN, int NVEC, int TLEN>
__global__ void __hammingdists(int *a, int *b, int *w, int *op, int *ow, int n) {
  printf("__hammingdists() not defined for CUDA Arch < 300\n");
}
#endif

int veccmp(int *a, int *b, int *d) {
  __veccmp<<<1,1>>>(a, b, d);

  asm("vadd4.s32.s32.s32" "%0, %1, %2, %3;": "=r" (TMP) : "r" (MM), "r" (RZ), "r" (RR));                   \
  asm("vmin4.s32.s32.s32" "%0, %1, %2, %3;": "=r" (RR) : "r" (TMP), "r" (RR), "r" (RR));       

#define hammingcell(A0,A1,B0,W0,C,TMP,ZERO)                                                               \
  asm("and.b32" "%0, %1, %2;": "=r" (TMP) : "r" (A0), "r" (B0));                                          \
  asm("vset4.s32.s32.eq" "%0, %1, %2, %3;": "=r" (TMP) : "r" (TMP), "r" (ZERO), "r" (ZERO));              \
  asm("vsub4.s32.s32.s32" "%0, %1, %2, %3;": "=r" (TMP) : "r" (ZERO), "r" (TMP), "r" (ZERO));             \
  asm("vmin4.u32.u32.u32.add" "%0, %1, %2, %3;": "=r" (C) : "r" (W0), "r" (TMP), "r" (C));                \
  asm("vmax4.u32.u32.u32" "%0, %1.b4321, %2.b4321, %3;": "=r" (A0) : "r" (A0), "r" (A1), "r" (ZERO));  

#define hammingend(A0)                               \
  asm("shr.b32" "%0, %1, 8;": "=r" (A0) : "r" (A0)); 
