 		          BF_STATUS_INSUFFICIENT_STORAGE);
 	}
 	BF_CHECK_CUFFT( cufftSetWorkArea(_handle, tmp_storage) );
-	// TODO: Try to reduce repetition in here
+	
+	CallbackData h_callback_data;
+	// WAR for CUFFT insisting that pointer be aligned to sizeof(cufftComplex)
+	int alignment = (_nbit == 32 ?
+	                 sizeof(cufftComplex) :
+	                 sizeof(cufftDoubleComplex));
+	h_callback_data.ptr_offset = (uintptr_t)idata % sizeof(cufftComplex);
+	*(char**)&idata -= h_callback_data.ptr_offset;
+	CallbackData* d_callback_data = thrust::raw_pointer_cast(&_dv_callback_data[0]);
+	cudaMemcpyAsync(d_callback_data, &h_callback_data, sizeof(CallbackData),
+	                cudaMemcpyHostToDevice, g_cuda_stream);
+	
+	BF_ASSERT((uintptr_t)idata % alignment == 0, BF_STATUS_UNSUPPORTED_STRIDE);
+	BF_ASSERT((uintptr_t)odata % alignment == 0, BF_STATUS_UNSUPPORTED_STRIDE);
+	
 	if( !_real_in && !_real_out ) {
 		int direction = inverse ? CUFFT_INVERSE : CUFFT_FORWARD;
 		if( _nbit == 32 ) {
-			BF_ASSERT((uintptr_t)idata % sizeof(cufftComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
-			BF_ASSERT((uintptr_t)odata % sizeof(cufftComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
 			BF_CHECK_CUFFT( cufftExecC2C(_handle, (cufftComplex*)idata, (cufftComplex*)odata, direction) );
 		} else if( _nbit == 64 ) {
-			BF_ASSERT((uintptr_t)idata % sizeof(cufftDoubleComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
-			BF_ASSERT((uintptr_t)odata % sizeof(cufftDoubleComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
 			BF_CHECK_CUFFT( cufftExecZ2Z(_handle, (cufftDoubleComplex*)idata, (cufftDoubleComplex*)odata, direction) );
 		} else {
 			BF_FAIL("Supported data types", BF_STATUS_UNSUPPORTED_DTYPE);
 		}
 	} else if( _real_in && !_real_out ) {
 		if( _nbit == 32 ) {
-			BF_ASSERT((uintptr_t)idata % sizeof(cufftComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
-			BF_ASSERT((uintptr_t)odata % sizeof(cufftComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
 			BF_CHECK_CUFFT( cufftExecR2C(_handle, (cufftReal*)idata, (cufftComplex*)odata) );
 		} else if( _nbit == 64 ) {
-			BF_ASSERT((uintptr_t)idata % sizeof(cufftDoubleComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
-			BF_ASSERT((uintptr_t)odata % sizeof(cufftDoubleComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
 			BF_CHECK_CUFFT( cufftExecD2Z(_handle, (cufftDoubleReal*)idata, (cufftDoubleComplex*)odata) );
 		} else {
 			BF_FAIL("Supported data types", BF_STATUS_UNSUPPORTED_DTYPE);
 		}
 	} else if( !_real_in && _real_out ) {
 		if( _nbit == 32 ) {
-			BF_ASSERT((uintptr_t)idata % sizeof(cufftComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
-			BF_ASSERT((uintptr_t)odata % sizeof(cufftComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
 			BF_CHECK_CUFFT( cufftExecC2R(_handle, (cufftComplex*)idata, (cufftReal*)odata) );
 		} else if( _nbit == 64 ) {
-			BF_ASSERT((uintptr_t)idata % sizeof(cufftDoubleComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
-			BF_ASSERT((uintptr_t)odata % sizeof(cufftDoubleComplex) == 0, BF_STATUS_UNSUPPORTED_STRIDE);
 			BF_CHECK_CUFFT( cufftExecZ2D(_handle, (cufftDoubleComplex*)idata, (cufftDoubleReal*)odata) );
 		} else {
 			BF_FAIL("Supported data types", BF_STATUS_UNSUPPORTED_DTYPE);

		          bifrost_status(cufft_ret)); \
	} while(0)

__device__
cufftComplex callback_load_ci4(void*  dataIn,
                               size_t offset,
                               void*  callerInfo,
                               void*  sharedPointer) {
	int8_t packed = ((int8_t*)dataIn)[offset];
	int8_t real = packed & 0xF0;
	int8_t imag = packed << 4;

		               BF_STATUS_DEVICE_ERROR );
		BF_CHECK_CUFFT( cufftXtSetCallback(_handle,
		                                   (void**)&callback_load_r_hptr,
		                                   CUFFT_CB_LD_REAL,
		                                   (void**)&callback_data) );
		break;
	}
	case BF_DTYPE_I16: {
