 // should be a multiple of 32 (warp size)
 #define BLOCKSIZE 1024
 
-static int blocksize(int N)
-{
-	return BLOCKSIZE;
-}
-
-static long gridsize(long N)
-{
-	return (N + BLOCKSIZE - 1) / BLOCKSIZE;
-}
-
-#ifdef USE_FXDIV
-__device__ static inline uint32_t cuda_fxdiv_quotient_uint32_t(uint32_t n, const struct fxdiv_divisor_uint32_t divisor) {
-
-	const uint32_t t = __umulhi(n, divisor.m);
-	return (t + ((n - t) >> divisor.s1)) >> divisor.s2;
-}
-
-__device__ static inline uint64_t cuda_fxdiv_quotient_uint64_t(uint64_t n, const struct fxdiv_divisor_uint64_t divisor) {
-
-	if (1 == divisor.m && 0 == divisor.s1 && 0 == divisor.s2)
-		return n;
-
-	const uint64_t t = __umul64hi(n, divisor.m);
-	return (t + ((n - t) >> divisor.s1)) >> divisor.s2;
-}
-#endif
-
-
-#define MAX_CONV_DIMS 3
-struct im2col_descriptor_uint64 {
-
-	uint64_t NC;		// number channels
-	long istrs_NC;		// 1
-	long ostrs_NC;		// 1
-
-	int N_conv_dims;	// number convolutions (i.e. 1D, 2D, 3D)
-
-	uint64_t odims[MAX_CONV_DIMS];	// dimensions of the convolution (not including channel)
-	uint64_t kdims[MAX_CONV_DIMS];
-	uint64_t idims[MAX_CONV_DIMS];
-
-	long istrs_odims[MAX_CONV_DIMS];	// input strides of im2col (in elements)
-	long istrs_kdims[MAX_CONV_DIMS];
-	long ostrs_kdims[MAX_CONV_DIMS];	// output strides of im2col (in elements)
-	long ostrs_odims[MAX_CONV_DIMS];
-
-#ifdef USE_FXDIV
-	struct fxdiv_divisor_uint64_t div_NC;			//efficient fixed divisor for dimensions
-	struct fxdiv_divisor_uint64_t div_odims[MAX_CONV_DIMS];
-	struct fxdiv_divisor_uint64_t div_kdims[MAX_CONV_DIMS];
-	struct fxdiv_divisor_uint64_t div_idims[MAX_CONV_DIMS];
-#endif
-
-	long N_in_elements;		// channels * in-dims
-	long N_out_elements;		// channels * out-dims * krn-dims
-	long N_out_elements_o_only;	// channels * out-dims
-	long N_out_elements_k_only;	// channels * krn-dims
+
+static void getBlockSize3_internal(int block[3], const long dims[3], const void* func)
+{
+	cudaFuncAttributes attr;
+	cudaFuncGetAttributes(&attr, func);
+	int threads = attr.maxThreadsPerBlock;
+
+	block[0] = 1;
+	block[1] = 1;
+	block[2] = 1;
+
+	while ((threads >= 2) && (block[0] < dims[0])) {
+
+		block[0] *= 2;
+		threads /= 2;
+	}
+
+	while ((threads >= 2) && (block[1] < dims[1])) {
+
+		block[1] *= 2;
+		threads /= 2;
+	}
+
+	while ((threads >= 2) && (block[2] < dims[2])) {
+
+		block[2] *= 2;
+		threads /= 2;
+	}
+}
+
+static dim3 getBlockSize3(const long dims[3], const void* func)
+{
+	int block[3];
+
+	getBlockSize3_internal(block, dims, func);
+
+	return dim3(block[0], block[1], block[2]);
+}
+
+static long gridsize_int(long N, int blocksize)
+{
+	return (N + blocksize - 1) / blocksize;
+}
+
+static dim3 getGridSize3(const long dims[3], const void* func)
+{
+	int block[3];
+
+	getBlockSize3_internal(block, dims, func);
+
+	return dim3(gridsize_int(dims[0], block[0]), gridsize_int(dims[1], block[1]), gridsize_int(dims[2], block[2]));
+}
+
+
+
+static dim3 blocksize(int N, const void* func)
+{
+	const long dims[3] = { N, 1, 1};
+	return getBlockSize3(dims, func);
+}
+
+static dim3 gridsize(long N, const void* func)
+{
+	const long dims[3] = { N, 1, 1};
+	return getGridSize3(dims, func);
+}
+
+
+template <int DIMS, typename T> 
+struct im2col_descriptor {
+
+	T NC;		// number channels
+	T istrs_NC;		// 1
+	T ostrs_NC;		// 1
+
+	T odims[DIMS];	// dimensions of the convolution (not including channel)
+	T kdims[DIMS];
+	T idims[DIMS];
+
+	T istrs_odims[DIMS];	// input strides of im2col (in elements)
+	T istrs_kdims[DIMS];
+	T ostrs_kdims[DIMS];	// output strides of im2col (in elements)
+	T ostrs_odims[DIMS];
+
+	T N_in_elements;		// channels * in-dims
+	T N_out_elements;		// channels * out-dims * krn-dims
+	T N_out_elements_o_only;	// channels * out-dims
+	T N_out_elements_k_only;	// channels * krn-dims
 
 	bool triv_strides_dilation;	// trivial dilation and strides
 };
 
-// same struct as im2col_descriptor_uint64 but with uint32_t for faster divisons
-struct im2col_descriptor_uint32 {
-
-	uint32_t NC;		// number channels
-	long istrs_NC;		// 1
-	long ostrs_NC;		// 1
-
-	int N_conv_dims;	// number convolutions (i.e. 1D, 2D, 3D)
-
-	uint32_t odims[MAX_CONV_DIMS];	// dimensions of the convolution (not including channel)
-	uint32_t kdims[MAX_CONV_DIMS];
-	uint32_t idims[MAX_CONV_DIMS];
-
-	long istrs_odims[MAX_CONV_DIMS];	// input strides of im2col (in elements)
-	long istrs_kdims[MAX_CONV_DIMS];
-	long ostrs_kdims[MAX_CONV_DIMS];	// output strides of im2col (in elements)
-	long ostrs_odims[MAX_CONV_DIMS];
-
-#ifdef USE_FXDIV
-	struct fxdiv_divisor_uint32_t div_NC;			//efficient fixed divisor for dimensions
-	struct fxdiv_divisor_uint32_t div_odims[MAX_CONV_DIMS];
-	struct fxdiv_divisor_uint32_t div_kdims[MAX_CONV_DIMS];
-	struct fxdiv_divisor_uint32_t div_idims[MAX_CONV_DIMS];
-#endif
-
-	long N_in_elements;		// channels * in-dims
-	long N_out_elements;		// channels * out-dims * krn-dims
-	long N_out_elements_o_only;	// channels * out-dims
-	long N_out_elements_k_only;	// channels * krn-dims
-
-	bool triv_strides_dilation;	// trivial dilation and strides
-};
-
-
-static struct im2col_descriptor_uint64 get_im2col_descriptor_uint64(const long odims[5], const long idims[5], const long kdims[5], const long dilation[5], const long strides[5])
-{
-	struct im2col_descriptor_uint64 config;
+template <int DIMS, typename T> 
+static struct im2col_descriptor<DIMS, T>get_im2col_descriptor(const long odims[5], const long idims[5], const long kdims[5], const long dilation[5], const long strides[5])
+{
+	struct im2col_descriptor<DIMS, T>config;
 
 	config.NC = idims[1];
 	config.istrs_NC = 1;
 	config.ostrs_NC = 1;
 
-#ifdef USE_FXDIV
-	config.div_NC = fxdiv_init_uint64_t(idims[1]);
-#endif
-	config.N_conv_dims = 0;
 	config.N_in_elements = idims[1];
 	config.N_out_elements = idims[1];
 	config.N_out_elements_o_only = idims[1];

	config.ostrs_odims[0] = config.N_out_elements_k_only;
	config.ostrs_kdims[0] = config.NC;

	for (int i = 1; i < config.N_conv_dims; i++) {

		config.ostrs_odims[i] = config.ostrs_odims[i - 1] * config.odims[i - 1];
		config.ostrs_kdims[i] = config.ostrs_kdims[i - 1] * config.kdims[i - 1];

 * */
extern "C" void cuda_im2col_transp(_Complex float* dst, const _Complex float* src, const long odims[5], const long idims[5], const long kdims[5], const long dilation[5], const long strides[5])
{
	cuda_im2col_int2<true>(dst, src, odims, idims, kdims, dilation, strides);
}

