 	
 	_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,
+	BF_CHECK( set_fft_load_callback(in->dtype, _nbit, _handle, _do_fftshift ^ _do_fftshift_before,
 	                                callback_data, &_using_load_callback) );
 	
 	if( tmp_storage_size ) {

	
	// Set callback data needed for applying fftshift
	h_callback_data->inverse = _real_out || (!_real_in && inverse);
	h_callback_data->do_fftshift = _do_fftshift;
	h_callback_data->ndim = _axes.size();
	for( int d=0; d<h_callback_data->ndim; ++d ) {
		h_callback_data->shape[d]    = in->shape[_axes[d]];

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