 	_axes.assign(axes, axes+rank);
 	_do_fftshift = do_fftshift;
 	_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,
 	                                callback_data, &_using_load_callback) );

	void* idata = in->data;
	void* odata = out->data;
	
	CallbackData h_callback_data;
	// WAR for CUFFT insisting that pointer be aligned to sizeof(cufftComplex)
	int alignment = (_nbit == 32 ?
	                 sizeof(cufftComplex) :

	//         so we really need to set the callback here instead. Not sure
	//         how expensive it is to set the callback.
	if( _using_load_callback ) {
		h_callback_data->ptr_offset = (uintptr_t)idata % sizeof(cufftComplex);
		*(char**)&idata -= h_callback_data->ptr_offset;
	}
	
	// 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]];
		int itype_nbyte = BF_DTYPE_NBYTE(in->dtype);
		h_callback_data->istrides[d] = in->strides[_axes[d]] / itype_nbyte;
		h_callback_data->inembed[d]  =
			(_axes[d] > 0 ?
			 in->strides[_axes[d]-1] / in->strides[_axes[d]] :
			 in->shape[_axes[d]]);
	}
	
	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);
