A helper function containing the common part of several kernels below. Precondition: 'algorithm' and 'shape' are compile-time constants.
| 169 | // A helper function containing the common part of several kernels below. |
| 170 | // Precondition: 'algorithm' and 'shape' are compile-time constants. |
| 171 | Status CompileImpl( |
| 172 | XlaOpKernelContext* ctx, int state_input_idx, int alg_input_idx, |
| 173 | int shape_input_idx, |
| 174 | std::function<SamplerReturnType(Algorithm, xla::XlaOp, xla::XlaOp, |
| 175 | TensorShape)> const& sampler) { |
| 176 | auto alg_shape = ctx->InputShape(alg_input_idx); |
| 177 | if (alg_shape.dims() != 0) { |
| 178 | return errors::InvalidArgument("algorithm must be of shape [], not ", |
| 179 | alg_shape.DebugString()); |
| 180 | } |
| 181 | xla::Literal alg_literal; |
| 182 | TF_RETURN_IF_ERROR(ctx->ConstantInput(alg_input_idx, &alg_literal)); |
| 183 | auto alg = alg_literal.Get<Algorithm>({}); |
| 184 | if (!(alg == RNG_ALG_THREEFRY || alg == RNG_ALG_PHILOX)) { |
| 185 | return errors::InvalidArgument("Unsupported algorithm id: ", alg); |
| 186 | } |
| 187 | |
| 188 | xla::XlaOp var; |
| 189 | TensorShape var_shape; |
| 190 | TF_RETURN_IF_ERROR(ctx->ReadVariableInput( |
| 191 | state_input_idx, STATE_ELEMENT_DTYPE, &var_shape, &var)); |
| 192 | TF_RETURN_IF_ERROR(CheckStateShape(alg, var_shape)); |
| 193 | TensorShape shape; |
| 194 | TF_RETURN_IF_ERROR(ctx->ConstantInputAsShape(shape_input_idx, &shape)); |
| 195 | xla::XlaOp state; |
| 196 | xla::XlaOp key; |
| 197 | std::tie(state, key) = StateAndKeyFromVariable(alg, var); |
| 198 | auto status_or_value = sampler(alg, state, key, shape); |
| 199 | if (!status_or_value.ok()) { |
| 200 | return status_or_value.status(); |
| 201 | } |
| 202 | xla::RngOutput value_state = status_or_value.ConsumeValueOrDie(); |
| 203 | state = value_state.state; |
| 204 | ctx->SetOutput(0, value_state.value); |
| 205 | var = StateAndKeyToVariable(alg, state, key); |
| 206 | xla::PrimitiveType state_element_type; |
| 207 | TF_RETURN_IF_ERROR( |
| 208 | DataTypeToPrimitiveType(STATE_ELEMENT_DTYPE, &state_element_type)); |
| 209 | var = BitcastConvertType(var, state_element_type); |
| 210 | TF_RETURN_IF_ERROR( |
| 211 | ctx->AssignVariable(state_input_idx, STATE_ELEMENT_DTYPE, var)); |
| 212 | return Status::OK(); |
| 213 | } |
| 214 | |
| 215 | class StatefulUniformOp : public XlaOpKernel { |
| 216 | public: |
no test coverage detected