 	                &d_out,
 	                &ostride,
 	                &obatchstride};
-	cudaLaunchKernel((void*)fdmt_init_kernel<InType,OutType>,
-	                 grid, block,
-	                 &args[0], 0, stream);
+	BF_CHECK_CUDA_EXCEPTION(
+		cudaLaunchKernel((void*)fdmt_init_kernel<InType,OutType>,
+		                 grid, block,
+		                 &args[0], 0, stream),
+		BF_STATUS_INTERNAL_ERROR);
 }
 
 template<typename DType>

		in  =  &in_flattened;
		BF_ASSERT(in_flattened.ndim == out_flattened.ndim,
		          BF_STATUS_INTERNAL_ERROR);
		BF_ASSERT(in_flattened.ndim == 3, BF_STATUS_UNSUPPORTED_SHAPE);
		BF_ASSERT_EXCEPTION(in_flattened.shape[0] == out_flattened.shape[0],
		                    BF_STATUS_INVALID_SHAPE);

				ostride_cur += reverse_time ? +1 : -1;
				obatchstride_cur = obatchstride;
			}
			launch_fdmt_exec_kernel(ntime, nrow, nbatch, (step==nstep-1), reverse_time,
			                        _d_step_delays  + step*_plan_stride,
			                        _d_step_srcrows + step*_plan_stride,
			                        d_ibuf, _buffer_stride, _batch_stride,
			                        d_obuf, ostride_cur, obatchstride_cur,
			                        _stream);
			std::swap(d_ibuf, d_obuf);
		}
		BF_CHECK_CUDA_EXCEPTION(cudaGetLastError(), BF_STATUS_INTERNAL_ERROR);
	}
	void set_stream(cudaStream_t stream) {
		_stream = stream;
