Flexible implementation of the discrete Fourier transform. All transform types (FFT, IFFT, RFFT, and IRFFT) are supported, as well as the arbitrary rank and length of each dimension of the transform, and arbitrary layouts of the input and output literals. The input literal in operand 0 provides input data, which must be complex64 for FFT, IFFT, IRFFT transforms and float for RFFT. The transform i
| 1360 | // O(N0*N1*...*Nn*(log(N0)+log(N1)+...+log(Nn)) in the best case. |
| 1361 | // |
| 1362 | Status HloEvaluator::HandleFft(HloInstruction* fft) { |
| 1363 | const FftType fft_type = fft->fft_type(); |
| 1364 | std::vector<int64> fft_lengths = fft->fft_length(); |
| 1365 | const int64 fft_rank = fft_lengths.size(); |
| 1366 | const Literal& input_literal = GetEvaluatedLiteralFor(fft->operand(0)); |
| 1367 | const Shape& input_shape = input_literal.shape(); |
| 1368 | const Shape& output_shape = fft->shape(); |
| 1369 | Literal output_literal = Literal::CreateFromShape(output_shape); |
| 1370 | |
| 1371 | // Make fft_lengths[0] the minor-most dimension. |
| 1372 | absl::c_reverse(fft_lengths); |
| 1373 | |
| 1374 | TF_RETURN_IF_ERROR(CheckParameters(input_shape, output_shape, fft_rank, |
| 1375 | fft_type, fft_lengths)); |
| 1376 | |
| 1377 | const auto fft_strides = ComputeStrides(fft_lengths); |
| 1378 | |
| 1379 | // Working set size. |
| 1380 | const int64 fft_size = fft_strides[fft_rank]; |
| 1381 | |
| 1382 | if (fft_size > 0) { |
| 1383 | // Linearized working data set. |
| 1384 | std::vector<complex128> data(fft_size); |
| 1385 | |
| 1386 | // Temporary buffer allocated once and used in 1D sweeps. For dimension |
| 1387 | // length values that are powers of 2, the buffer should be twice as large. |
| 1388 | int64 buffer_size = 0; |
| 1389 | for (auto len : fft_lengths) { |
| 1390 | int64 size = IsPowerOfTwo(static_cast<uint64>(len)) ? len * 2 : len; |
| 1391 | buffer_size = std::max(buffer_size, size); |
| 1392 | } |
| 1393 | std::vector<complex128> buffer(buffer_size); |
| 1394 | |
| 1395 | // Sizes of each axis of input and output literals. |
| 1396 | const auto input_lengths = GetDimensionLengths(input_literal); |
| 1397 | const auto output_lengths = GetDimensionLengths(output_literal); |
| 1398 | |
| 1399 | // Strides for generating linearized indices into multidimensional arrays. |
| 1400 | const auto input_strides = ComputeStrides(input_lengths, input_literal); |
| 1401 | const auto output_strides = ComputeStrides(output_lengths, output_literal); |
| 1402 | |
| 1403 | // Visit all elements in the dimensions with ranks above the FFT rank. For |
| 1404 | // each such element invoke the transform. Use separate indices for the |
| 1405 | // input and the output to allow different layouts. |
| 1406 | auto base_case = [&](int64 axis, int64 output_index, int64 input_index, |
| 1407 | bool within_src_bounds) { |
| 1408 | if (axis == fft_rank - 1) { |
| 1409 | // Base case: copy the data from the input literal, apply the |
| 1410 | // transform, and copy the result to the output literal. |
| 1411 | CHECK(within_src_bounds); |
| 1412 | bool input_is_zero = |
| 1413 | CopyDataFromInput(input_literal, input_index, fft_rank, fft_type, |
| 1414 | fft_size, fft_lengths, fft_strides, input_lengths, |
| 1415 | input_strides, absl::MakeSpan(data)); |
| 1416 | if (!input_is_zero) { |
| 1417 | // Make 1D sweeps along each transform axis. |
| 1418 | Sweep(fft_rank, fft_type, fft_lengths, fft_strides, |
| 1419 | absl::MakeSpan(data), absl::MakeSpan(buffer)); |
nothing calls this directly
no test coverage detected