 	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++) {
+	for (int i = 1; i < 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(_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])
{
	struct im2col_descriptor_uint64 config64 = get_im2col_descriptor_uint64(odims, idims, kdims, dilation, strides);
	struct im2col_descriptor_uint32 config32 = get_im2col_descriptor_uint32(odims, idims, kdims, dilation, strides);

	if ((1 < config32.NC) && (config64.triv_strides_dilation)) {

		if (config32.N_out_elements < INT32_MAX)
			kern_im2col_valid_no_dil_str<<<gridsize(config32.N_in_elements), blocksize(config32.N_in_elements)>>>(config32, (cuFloatComplex*) dst, (cuFloatComplex*) src);
		else
	 		kern_im2col_valid_no_dil_str<<<gridsize(config64.N_in_elements), blocksize(config32.N_in_elements)>>>(config64, (cuFloatComplex*) dst, (cuFloatComplex*) src);

	} else {

		if (config32.N_out_elements < INT32_MAX)
			kern_im2col_valid<<<gridsize(config32.N_out_elements), blocksize(config32.N_out_elements)>>>(config32, (cuFloatComplex*) dst, (cuFloatComplex*) src);
		else
			kern_im2col_valid<<<gridsize(config64.N_out_elements), blocksize(config64.N_out_elements)>>>(config64, (cuFloatComplex*) dst, (cuFloatComplex*) src);
	}
}

__global__ static void kern_im2col_valid_no_dil_str_transp(struct im2col_descriptor_uint64 config, cuFloatComplex* dst, const cuFloatComplex* src)
{
	int start = threadIdx.x + blockDim.x * blockIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (long i = start; i < config.N_in_elements; i += stride) {

		uint64_t i_cur = i;
		uint64_t i_new;

#ifdef USE_FXDIV
		i_new = cuda_fxdiv_quotient_uint64_t(i_cur, config.div_NC);
#else
		i_new = i_cur / config.NC;
#endif
		uint64_t c = i_cur - i_new * config.NC;

		i_cur = i_new;
#ifdef USE_FXDIV
		i_new = cuda_fxdiv_quotient_uint64_t(i_cur, config.div_idims[0]);
#else
		i_new = i_cur / config.idims[0];
#endif
		uint64_t ix = i_cur - i_new * config.idims[0];

		i_cur = i_new;
#ifdef USE_FXDIV
		i_new = cuda_fxdiv_quotient_uint64_t(i_cur, config.div_idims[1]);
#else
		i_new = i_cur / config.idims[1];
#endif
		uint64_t iy = i_cur - i_new * config.idims[1];
		uint64_t iz = i_new;

		cuFloatComplex result = dst[i];

		for (long kz = 0; kz < (long)config.kdims[2]; kz++) {

			long oz = (long)iz - kz;

			if ((0 > oz) || ((int)config.odims[2] <= oz))
				continue;

			long offset_z = config.N_out_elements_k_only * config.odims[0] * config.odims[1] * oz + config.NC * config.kdims[0] * config.kdims[1] * kz;

			for (long ky = 0; ky < (long)config.kdims[1]; ky++) {

				long oy = (long)iy - ky;

				if ((0 > oy) || ((int)config.odims[1] <= oy))
					continue;

				long offset_y = config.N_out_elements_k_only * config.odims[0] * oy + config.NC * config.kdims[0] * ky;

				for (long kx = 0; kx < (long)config.kdims[0]; kx++) {

					long ox = (long)ix - kx;
					long offset_x = config.N_out_elements_k_only * ox + config.NC * kx;
					long index = c 	+ offset_x + offset_y + offset_z;

					if ((0 <= ox) && ((int)config.odims[0] > ox))
						result = cuCaddf(result, src[index]);
				}
			}
		}

		dst[i] = result;
	}
}

__global__ static void kern_im2col_valid_no_dil_str_transp(struct im2col_descriptor_uint32 config, cuFloatComplex* dst, const cuFloatComplex* src)
{
	int start = threadIdx.x + blockDim.x * blockIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (long i = start; i < config.N_in_elements; i += stride) {

		uint32_t i_cur = i;
		uint32_t i_new;

#ifdef USE_FXDIV
		i_new = cuda_fxdiv_quotient_uint32_t(i_cur, config.div_NC);
#else
		i_new = i_cur / config.NC;
#endif
		uint32_t c = i_cur - i_new * config.NC;

		i_cur = i_new;

#ifdef USE_FXDIV
		i_new = cuda_fxdiv_quotient_uint32_t(i_cur, config.div_idims[0]);
#else
		i_new = i_cur / config.idims[0];
#endif
		uint32_t ix = i_cur - i_new * config.idims[0];

		i_cur = i_new;

#ifdef USE_FXDIV
		i_new = cuda_fxdiv_quotient_uint32_t(i_cur, config.div_idims[1]);
#else
		i_new = i_cur / config.idims[1];
#endif

		uint32_t iy = i_cur - i_new * config.idims[1];
		uint32_t iz = i_new;

		cuFloatComplex result = dst[i];

		for (int kz = 0; kz < (int)config.kdims[2]; kz++) {

			int oz = (int)iz - kz;

			if ((0 > oz) || ((int)config.odims[2] <= oz))
				continue;

			long offset_z = config.N_out_elements_k_only * config.odims[0] * config.odims[1] * oz + config.NC * config.kdims[0] * config.kdims[1] * kz;

			for(int ky = 0; ky < (int)config.kdims[1]; ky++) {

				int oy = (int)iy - ky;

				if ((0 > oy) || ((int)config.odims[1] <= oy))
					continue;

				long offset_y = config.N_out_elements_k_only * config.odims[0] * oy + config.NC * config.kdims[0] * ky;

				for (int kx = 0; kx < (int)config.kdims[0]; kx++) {

					int ox = (int)ix - kx;

					long offset_x = config.N_out_elements_k_only * ox + config.NC * kx;

					long index = c 	+ offset_x + offset_y + offset_z;

					if ((0 <= ox) && ((int)config.odims[0] > ox))
						result = cuCaddf(result, src[index]);
				}
			}
		}

		dst[i] = result;
	}
}

#ifdef NON_DETERMINISTIC

__global__ static void kern_im2col_valid_transp(struct im2col_descriptor_uint32 config, cuFloatComplex* dst, const cuFloatComplex* src)
{
	int start = threadIdx.x + blockDim.x * blockIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (long i = start; i < config.N_out_elements; i += stride) {

		uint32_t i0 = i;
		uint32_t i_new = i;
		long in_index = 0;

		if (1 < config.NC) {

#ifdef USE_FXDIV
			i_new = cuda_fxdiv_quotient_uint32_t(i0, config.div_NC);
#else
			i_new = i0 / config.NC;
#endif
			in_index = (i0 - config.NC * i_new) * config.istrs_NC;
			i0 = i_new;
		}

		for (int j = 0; j < config.N_conv_dims; j++) {

#ifdef USE_FXDIV
			i_new = cuda_fxdiv_quotient_uint32_t(i0, config.div_kdims[j]);
#else
			i_new = i0 / config.kdims[j];
#endif
			in_index += config.istrs_kdims[j] * (i0 - config.kdims[j] * i_new);
			i0 = i_new;
		}

		for (int j = 0; j < config.N_conv_dims - 1; j++) {

#ifdef USE_FXDIV
			i_new = cuda_fxdiv_quotient_uint32_t(i0, config.div_odims[j]);
#else
			i_new = i0 / config.odims[j];
#endif
			in_index += config.istrs_odims[j] * (i0 - config.odims[j] * i_new);
			i0 = i_new;
		}

		in_index += i0 * config.istrs_odims[config.N_conv_dims - 1];

		atomicAdd(&(dst[in_index].x), src[i].x);
		atomicAdd(&(dst[in_index].y), src[i].y);
	}
}

__global__ static void kern_im2col_valid_transp(struct im2col_descriptor_uint64 config, cuFloatComplex* dst, const cuFloatComplex* src)
{
	int start = threadIdx.x + blockDim.x * blockIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (long i = start; i < config.N_out_elements; i += stride) {

		uint64_t i0 = i;
		uint64_t i_new = i;
		long in_index = 0;

		if (1 < config.NC) {

#ifdef USE_FXDIV
			i_new = cuda_fxdiv_quotient_uint64_t(i0, config.div_NC);
#else
			i_new = i0 / config.NC;
#endif
			in_index = (i0 - config.NC * i_new) * config.istrs_NC;
			i0 = i_new;
		}

		for (int j = 0; j < config.N_conv_dims; j++) {

#ifdef USE_FXDIV
			i_new = cuda_fxdiv_quotient_uint64_t(i0, config.div_kdims[j]);
#else
			i_new = i0 / config.kdims[j];
#endif
			in_index += config.istrs_kdims[j] * (i0 - config.kdims[j] * i_new);
			i0 = i_new;
		}

		for (int j = 0; j < config.N_conv_dims - 1; j++) {

#ifdef USE_FXDIV
			i_new = cuda_fxdiv_quotient_uint64_t(i0, config.div_odims[j]);
#else
			i_new = i0 / config.odims[j];
#endif
			in_index += config.istrs_odims[j] * (i0 - config.odims[j] * i_new);
			i0 = i_new;
		}

		in_index += i0 * config.istrs_odims[config.N_conv_dims - 1];

		atomicAdd(&(dst[in_index].x), src[i].x);
		atomicAdd(&(dst[in_index].y), src[i].y);
	}
}
#endif


/* *
 * Transposed/adjoint of cuda im2col

#include <cuda.h>
#include <cuComplex.h>

#include "misc/misc.h"

#include "num/gpu_conv.h"
#include "num/multind.h"

