 using std::endl;
 
 template<typename InType, typename OutType>
-__global__ void fir_kernel(int                        ncoeff,
-                           int                        decim, 
-                           int                        ntime, 
-                           int                        nantpol,
+__global__ void fir_kernel(unsigned int               ncoeff,
+                           unsigned int               decim, 
+                           unsigned int               ntime, 
+                           unsigned int               nantpol,
                            const double* __restrict__ coeffs,
                            Complex64*                 state0,
                            Complex64*                 state1,

	// Discover the dimensions of the FIR coefficients.  This uses the 
	// first dimension to set the number of coefficients.  All following
	// dimensions are merged together to get the the number of ant/pols.
	int ncoeff, nantpols;
	ncoeff = coeffs->shape[0];
	nantpols = 1;
	for(int i=1; i<coeffs->ndim; ++i) {

}

class BFfir_impl {
	typedef int          IType;
	typedef unsigned int UType;
	typedef double       FType;
public: // HACK WAR for what looks like a bug in the CUDA 7.0 compiler
	typedef float  DType;
private:
	UType        _ncoeff;
	UType        _decim;
	UType        _nantpol;
	double*      _coeffs = NULL;
	Complex64*   _state0 = NULL;
	Complex64*   _state1 = NULL;
