| 131 | } |
| 132 | |
| 133 | void Compile(XlaOpKernelContext* ctx) override { |
| 134 | xla::XlaBuilder* b = ctx->builder(); |
| 135 | TensorShape elem_shape = ctx->InputShape(1); |
| 136 | |
| 137 | XlaResource* resource; |
| 138 | OP_REQUIRES_OK(ctx, ctx->GetResourceInput(0, &resource)); |
| 139 | |
| 140 | // Initializes the Stack, if the element shape was not already known. |
| 141 | OP_REQUIRES_OK(ctx, MaybeInitializeStack(b, resource, dtype_, elem_shape)); |
| 142 | |
| 143 | xla::XlaOp ta = xla::GetTupleElement(resource->value(), 0); |
| 144 | xla::XlaOp index = xla::GetTupleElement(resource->value(), 1); |
| 145 | xla::XlaOp value = ctx->Input(1); |
| 146 | |
| 147 | // start_indices of the DynamicUpdateSlice are [index, 0, 0, ..., 0]. |
| 148 | std::vector<xla::XlaOp> start_indices(elem_shape.dims() + 1, |
| 149 | xla::ConstantR0<int32>(b, 0)); |
| 150 | start_indices[0] = index; |
| 151 | |
| 152 | TensorShape slice_shape = elem_shape; |
| 153 | slice_shape.InsertDim(0, 1LL); |
| 154 | auto update = xla::Reshape(value, slice_shape.dim_sizes()); |
| 155 | |
| 156 | // TODO(phawkins): We don't check the index is in bounds --- there is no |
| 157 | // error mechanism in XLA. |
| 158 | OP_REQUIRES_OK(ctx, |
| 159 | resource->SetValue(xla::Tuple( |
| 160 | b, {xla::DynamicUpdateSlice(ta, update, start_indices), |
| 161 | xla::Add(index, xla::ConstantR0<int32>(b, 1))}))); |
| 162 | |
| 163 | ctx->SetOutput(0, value); |
| 164 | } |
| 165 | |
| 166 | private: |
| 167 | DataType dtype_; |
nothing calls this directly
no test coverage detected