Build computation the backprop into input 'data'. Where input: grad_output is of dimension [batch, dim_0, ...dim_n, channel] ratio is of dimension [batch, dim_0, ...dim_n, 2] gather_indices is of dimension [batch, dim_0, ...dim_n, 3] data_shape is of dimension [batch, x(width), y(height), channel] Output: scatter-add to each 2x2 grad_data neighbor: grad_data[fx, fy, chan] += output_grad * dx * dy
| 246 | // where (dx, dy) is (1 - ratio). If (dx, dy) is out of bound, then the their |
| 247 | // contribution is 0 to 'grad_data'. |
| 248 | XlaOp CalculateGradData(XlaOpKernelContext* ctx, XlaOp grad_output, XlaOp ratio, |
| 249 | XlaOp gather_indices, XlaOp warp, |
| 250 | xla::PrimitiveType warp_type, TensorShape warp_shape, |
| 251 | int64 last_warp_dim, int64 data_channels, |
| 252 | xla::Shape data_shape) { |
| 253 | // Weights tensor has dimension [batch, dim_0, ... dim_n, 4]. |
| 254 | auto weights = BilinearWeights(ctx, ratio, warp_shape, warp_type); |
| 255 | |
| 256 | auto warp_dims = warp_shape.dim_sizes(); |
| 257 | std::vector<int64> warp_dims_without_last_dims(warp_dims.begin(), |
| 258 | warp_dims.end() - 1); |
| 259 | |
| 260 | std::vector<int64> reshaped_weights_dims = warp_dims_without_last_dims; |
| 261 | // Reshape the last dimension of size 4 to two dimensions [2, 2]. |
| 262 | reshaped_weights_dims.push_back(2); |
| 263 | reshaped_weights_dims.push_back(2); |
| 264 | std::vector<int64> reshape_dims(warp_shape.dims()); |
| 265 | std::iota(reshape_dims.begin(), reshape_dims.end(), 0); |
| 266 | // The dimension is [batch, dim_0,..., dim_n, 2, 2]. |
| 267 | auto reshaped_weights = xla::Reshape(weights, /*dimensions=*/reshape_dims, |
| 268 | /*new_sizes=*/reshaped_weights_dims); |
| 269 | |
| 270 | std::vector<int64> weights_with_channels_dims = reshaped_weights_dims; |
| 271 | weights_with_channels_dims.push_back(data_channels); |
| 272 | std::vector<int64> reshaped_weights_indices(reshaped_weights_dims.size()); |
| 273 | std::iota(reshaped_weights_indices.begin(), reshaped_weights_indices.end(), |
| 274 | 0); |
| 275 | |
| 276 | // Set out of bound weights to 0. |
| 277 | // The dimension of the reshaped_weight: [batch, dim_0, ...dim_n, 2, 2]. |
| 278 | std::vector<int64> reshaped_result_dims(warp_dims.begin(), |
| 279 | warp_dims.end() - 1); |
| 280 | reshaped_result_dims.push_back(2); |
| 281 | reshaped_result_dims.push_back(2); |
| 282 | std::vector<int64> broadcasted_dims(warp_dims.size() - 1); |
| 283 | std::iota(broadcasted_dims.begin(), broadcasted_dims.end(), 0); |
| 284 | reshaped_weights = BoundSamples(ctx, warp, warp_type, warp_shape, |
| 285 | reshaped_result_dims, broadcasted_dims, |
| 286 | last_warp_dim, data_shape, reshaped_weights); |
| 287 | |
| 288 | // The dimension is [batch, dim_0, ..., dim_n, 2, 2, data_channel]. |
| 289 | auto broadcast_reshaped_weights = xla::BroadcastInDim( |
| 290 | reshaped_weights, weights_with_channels_dims, reshaped_weights_indices); |
| 291 | |
| 292 | std::vector<int64> grad_output_indices(warp_dims_without_last_dims.size()); |
| 293 | std::iota(grad_output_indices.begin(), grad_output_indices.end(), 0); |
| 294 | grad_output_indices.push_back(weights_with_channels_dims.size() - 1); |
| 295 | XlaOp broadcast_grad_output = xla::BroadcastInDim( |
| 296 | grad_output, weights_with_channels_dims, grad_output_indices); |
| 297 | |
| 298 | auto grad_output_multiply_weights = |
| 299 | broadcast_grad_output * broadcast_reshaped_weights; |
| 300 | |
| 301 | auto grad_data = xla::ConstantLiteral( |
| 302 | ctx->builder(), xla::Literal::CreateFromShape(data_shape)); |
| 303 | |
| 304 | // Pad grad data then slice it back. |
| 305 | // |
no test coverage detected