   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);

#include <stdio.h>
#include "MatKernel.hpp"

#ifdef __CUDA_ARCH__ 
#if __CUDA_ARCH__ > 200

#define edcellupdate(RR,RP1,RP2,RPP,WUN,TMP)                                                               \

#pragma unroll
        for (k = 0; k < VECLEN; k++) {
          tmp1 = aa[k];
          hammingcell(tmp1, aa[k+1], bb[k], ww[k], c, tmp, zero);
          aa[k] = tmp1;
        }
        tmp1 = aa[VECLEN];
