 	BF_CHECK_CUFFT( cufftSetWorkArea(_handle, tmp_storage) );
 	void* idata = in->data;
 	void* odata = out->data;
+	
+	// Apply a pre-call fftshift to deal with c2r problems in CUDA 10.1+
+#if CUDA_VERSION >= 10100
+	if( _do_fftshift_before ) {
+		/* Pre-call fftshift here */
+	}
+#endif
 	
 	// TODO: This sync is needed to ensure that the previous h2d copy of
 	//         h_callback_data has finished before we overwrite it.

	size_t           _workspace_size;
	std::vector<int> _axes;
	bool             _do_fftshift;
	bool             _using_load_callback;
	thrust::device_vector<char> _dv_tmp_storage;
	thrust::device_vector<CallbackData> _dv_callback_data;

	
	_axes.assign(axes, axes+rank);
	_do_fftshift = do_fftshift;
	_do_fftshift_before = false;
#if CUDA_VERSION >= 10100
	// Figure out what to do with c2r with fftshift on CUDA 10.1+
	if( type == CUFFT_C2R || type == CUFFT_Z2D ) {
		_do_fftshift_before = true;
	}
#endif
	_dv_callback_data.resize(1);
	_hv_callback_data.resize(1);
	CallbackData* callback_data = thrust::raw_pointer_cast(&_dv_callback_data[0]);
	BF_CHECK( set_fft_load_callback(in->dtype, _nbit, _handle, _do_fftshift ^ _do_fftshift_before,
	                                callback_data, &_using_load_callback) );
	
	if( tmp_storage_size ) {
