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

                              OutType*     d_out,
                              cudaStream_t stream=0) {
	//cout << "LAUNCH for " << nelement << endl;
	dim3 block(std::min(256u, nantpol), 256u/std::min(256u, nantpol));
	int first = std::min((nantpol-1)/block.x+1, 65535u);
	int secnd = std::min((ntime/decim-1)/block.y+1, 65535u);
	int third = std::min((ntime/decim-secnd*block.y-1)/secnd+2, 65535u);
