diff --git a/tensorflow/core/distributed_runtime/master_test.cc b/tensorflow/core/distributed_runtime/master_test.cc index 5c2f17e31f819d..e8b00ae9dca443 100644 --- a/tensorflow/core/distributed_runtime/master_test.cc +++ b/tensorflow/core/distributed_runtime/master_test.cc @@ -389,8 +389,8 @@ TEST_F(MasterTest, EigenProblem) { TF_CHECK_OK(CreateSession(def, &handle, &initial_version)); // Temps supporting the computation of the convergence condition. - const Eigen::array sum_along_dim(0); - const Eigen::array matrix_transpose({1, 0}); + const Eigen::array sum_along_dim{0}; + const Eigen::array matrix_transpose{1, 0}; Tensor x(DT_FLOAT, TensorShape({2, 1})); Tensor y(DT_FLOAT, TensorShape({2, 1})); Eigen::Tensor y_square_sum; diff --git a/tensorflow/core/kernels/gather_nd_op_gpu.cu.cc b/tensorflow/core/kernels/gather_nd_op_gpu.cu.cc index 227cd311c244ee..c26f5bdf492e39 100644 --- a/tensorflow/core/kernels/gather_nd_op_gpu.cu.cc +++ b/tensorflow/core/kernels/gather_nd_op_gpu.cu.cc @@ -39,11 +39,14 @@ __global__ void GatherSliceOpKernel( const auto indices_i = indices + IXDIM * loc; bool out_of_bounds = false; Index offset = 0; + // Avoid empty std::array access, which fails to compile on GPU. + if constexpr (IXDIM > 0) { #pragma unroll - for (int j = 0; j < IXDIM; ++j) { - const Index index_j = ldg(indices_i + j); - out_of_bounds |= !FastBoundsCheck(index_j, batch_indices[j]); - offset += batch_strides[j] * index_j; + for (int j = 0; j < IXDIM; ++j) { + const Index index_j = ldg(indices_i + j); + out_of_bounds |= !FastBoundsCheck(index_j, batch_indices[j]); + offset += batch_strides[j] * index_j; + } } // TODO(ebrevdo): // This is the only part that depends on the offset. The part diff --git a/tensorflow/core/kernels/image/adjust_contrast_op.cc b/tensorflow/core/kernels/image/adjust_contrast_op.cc index 7cef95b9479022..df8650ebfed515 100644 --- a/tensorflow/core/kernels/image/adjust_contrast_op.cc +++ b/tensorflow/core/kernels/image/adjust_contrast_op.cc @@ -248,6 +248,7 @@ class AdjustContrastOpv2 : public AdjustContrastOpV2Base { TTypes::Tensor mean_flat(&mean(0, 0), mean.size()); TTypes::Tensor summation_scratch(&scratch(0, 0, 0), scratch.size()); + using Eigen::DenseIndex; typedef Eigen::array Index; const int64_t plane_size = image_size * channels; // Since the number of channels in the early layers is often small, a @@ -255,10 +256,10 @@ class AdjustContrastOpv2 : public AdjustContrastOpV2Base { // This algorithm repeatedly folds each image plane by half, until // only one set of channels remains. for (int64_t i = 0; i < batch; i++) { - auto input_plane = - input_flat.slice(Index(i * plane_size), Index(plane_size)); - auto summation_plane = - summation_scratch.slice(Index(i * plane_size), Index(plane_size)); + auto input_plane = input_flat.slice(Index{DenseIndex(i * plane_size)}, + Index{DenseIndex(plane_size)}); + auto summation_plane = summation_scratch.slice( + Index{DenseIndex(i * plane_size)}, Index{DenseIndex(plane_size)}); int64_t remaining_size = image_size; int round = 0; // Sum the input(i, :, k) into mean(i, k). Repeatedly splits the input @@ -289,26 +290,29 @@ class AdjustContrastOpv2 : public AdjustContrastOpV2Base { if (round == 0) { // In the first round, sum the left side and right side of the input // array into the summation area. - summation_plane.slice(Index(0), Index(right_size * channels)) = - input_plane.slice(Index(left_size * channels), - Index(right_size * channels)) + - input_plane.slice(Index(0), Index(right_size * channels)); + summation_plane.slice(Index{0}, + Index{DenseIndex(right_size * channels)}) = + input_plane.slice(Index{DenseIndex(left_size * channels)}, + Index{DenseIndex(right_size * channels)}) + + input_plane.slice(Index{0}, + Index{DenseIndex(right_size * channels)}); if (left_size > right_size) { DCHECK_EQ(left_size - right_size, 1); // Copy over the remaining column if the remaining_size is odd. // This also handles the case where image_size == 1. - summation_plane.slice(Index(right_size * channels), - Index(channels)) = - input_plane.slice(Index(right_size * channels), - Index(channels)); + summation_plane.slice(Index{DenseIndex(right_size * channels)}, + Index{DenseIndex(channels)}) = + input_plane.slice(Index{DenseIndex(right_size * channels)}, + Index{DenseIndex(channels)}); } } else { // For all the remaining rounds, add the second half of the inputs // into the first half of the inputs. With the flat structure and // large size, this utilizes vectorization between components. - summation_plane.slice(Index(0), Index(right_size * channels)) += - summation_plane.slice(Index(left_size * channels), - Index(right_size * channels)); + summation_plane.slice(Index{0}, + Index{DenseIndex(right_size * channels)}) += + summation_plane.slice(Index{DenseIndex(left_size * channels)}, + Index{DenseIndex(right_size * channels)}); } remaining_size = left_size; round++; @@ -316,9 +320,11 @@ class AdjustContrastOpv2 : public AdjustContrastOpV2Base { const float mean_scaling = 1.0f / image_size; // The first channels elements in summation_plane now holds the summation. // Scale it with image_size and copy over to the means. - auto mean_plane = mean_flat.slice(Index(i * channels), Index(channels)); + auto mean_plane = mean_flat.slice(Index{DenseIndex(i * channels)}, + Index{DenseIndex(channels)}); mean_plane = - summation_plane.slice(Index(0), Index(channels)) * mean_scaling; + summation_plane.slice(Index{0}, Index{DenseIndex(channels)}) * + mean_scaling; } } diff --git a/tensorflow/core/kernels/parameterized_truncated_normal_op_gpu.cu.cc b/tensorflow/core/kernels/parameterized_truncated_normal_op_gpu.cu.cc index b826564437c0a1..e7b76653dc329e 100644 --- a/tensorflow/core/kernels/parameterized_truncated_normal_op_gpu.cu.cc +++ b/tensorflow/core/kernels/parameterized_truncated_normal_op_gpu.cu.cc @@ -132,8 +132,6 @@ __global__ void __launch_bounds__(1024) (normMax >= T(0.))) || ((normMax > kStdDevsInsideBoundsToUseRandnSampler) && (normMin <= T(0.)))) { - Eigen::array n; - int numIterations = 0; while (numIterations < kMaxIterations) { const auto randn = normal_dist(&gen); diff --git a/tensorflow/core/kernels/random_binomial_op.cc b/tensorflow/core/kernels/random_binomial_op.cc index 8fceaf70c0dbbb..98118b78eb5b58 100644 --- a/tensorflow/core/kernels/random_binomial_op.cc +++ b/tensorflow/core/kernels/random_binomial_op.cc @@ -187,8 +187,6 @@ struct RandomBinomialFunctor { &gen, &output](int64_t start_output, int64_t limit_output) { // Vectorized intermediate calculations for uniform rejection sampling. // We always generate at most 4 samples. - Eigen::array z; - Eigen::array g; const bool should_bcast = bcast.IsBroadcastingRequired(); const auto& counts_batch_indices = bcast.x_batch_indices(); const auto& probs_batch_indices = bcast.y_batch_indices(); diff --git a/tensorflow/core/kernels/sparse_tensor_dense_matmul_op.cc b/tensorflow/core/kernels/sparse_tensor_dense_matmul_op.cc index 04aff711362552..d8b203dec3f40d 100644 --- a/tensorflow/core/kernels/sparse_tensor_dense_matmul_op.cc +++ b/tensorflow/core/kernels/sparse_tensor_dense_matmul_op.cc @@ -310,7 +310,7 @@ Status SparseTensorDenseMatMulImpl( if (ADJ_B) { // Perform transpose and conjugation on B once, since we chip out B's // columns in the nnz loop. - Eigen::array shuffle(1, 0); // preserve dimension order + Eigen::array shuffle{1, 0}; // preserve dimension order Eigen::Tensor col_major_conj_b = b.swap_layout().shuffle(shuffle).conjugate(); LOOP_NNZ(col_major_conj_b); diff --git a/third_party/xla/third_party/tsl/tsl/framework/convolution/eigen_spatial_convolutions_test.cc b/third_party/xla/third_party/tsl/tsl/framework/convolution/eigen_spatial_convolutions_test.cc index 48e0379bce1466..9f589c549901e9 100644 --- a/third_party/xla/third_party/tsl/tsl/framework/convolution/eigen_spatial_convolutions_test.cc +++ b/third_party/xla/third_party/tsl/tsl/framework/convolution/eigen_spatial_convolutions_test.cc @@ -1036,10 +1036,6 @@ static void PackLhsHelper(::testing::benchmark::State& state, reshape_dims[0] = filter_count; reshape_dims[1] = input_depth * filter_rows * filter_cols; - // We are going to contract along the 'in_depth * filter_rows * filter_cols`. - nocontract_t nocontract_dim = {0}; - contract_t contract_dim = {1}; - // These values computed using the algorithm in TensorContraction.h, with // 'nocontract_dim' and 'contract_dim' values specified above. nocontract_t nocontract_strides = {1};