| 40 | } |
| 41 | |
| 42 | void Compile(XlaOpKernelContext* ctx) override { |
| 43 | const TensorShape logits_shape = ctx->InputShape(0); |
| 44 | OP_REQUIRES(ctx, TensorShapeUtils::IsVectorOrHigher(logits_shape), |
| 45 | errors::InvalidArgument("logits must have >= 1 dimension, got ", |
| 46 | logits_shape.DebugString())); |
| 47 | |
| 48 | // Major dimensions are batch dimensions, minor dimension is the class |
| 49 | // dimension. |
| 50 | std::vector<int64> batch_dims(logits_shape.dims() - 1); |
| 51 | std::iota(batch_dims.begin(), batch_dims.end(), 0); |
| 52 | const int kClassDim = logits_shape.dims() - 1; |
| 53 | |
| 54 | const DataType type = input_type(0); |
| 55 | const xla::PrimitiveType xla_type = ctx->input_xla_type(0); |
| 56 | auto logits = ctx->Input(0); |
| 57 | xla::XlaOp softmax; |
| 58 | if (is_on_gpu_) { |
| 59 | softmax = xla::Softmax(logits, kClassDim, log_); |
| 60 | } else { |
| 61 | xla::XlaBuilder* const b = ctx->builder(); |
| 62 | const xla::XlaComputation& max_func = *ctx->GetOrCreateMax(type); |
| 63 | |
| 64 | // Find the max in each batch, resulting in a tensor of shape [batch] |
| 65 | auto logits_max = |
| 66 | xla::Reduce(logits, xla::MinValue(b, xla_type), max_func, {kClassDim}); |
| 67 | // Subtract the max in batch b from every element in batch b. Broadcasts |
| 68 | // along the batch dimension. |
| 69 | auto shifted_logits = xla::Sub(logits, logits_max, batch_dims); |
| 70 | auto exp_shifted = xla::Exp(shifted_logits); |
| 71 | const DataType accumulation_type = XlaHelpers::SumAccumulationType(type); |
| 72 | xla::PrimitiveType xla_accumulation_type; |
| 73 | OP_REQUIRES_OK(ctx, DataTypeToPrimitiveType(accumulation_type, |
| 74 | &xla_accumulation_type)); |
| 75 | auto converted = |
| 76 | xla::ConvertElementType(exp_shifted, xla_accumulation_type); |
| 77 | auto reduce = |
| 78 | xla::Reduce(converted, xla::Zero(b, xla_accumulation_type), |
| 79 | *ctx->GetOrCreateAdd(accumulation_type), {kClassDim}); |
| 80 | auto sum = XlaHelpers::ConvertElementType(reduce, type); |
| 81 | softmax = |
| 82 | log_ |
| 83 | // softmax = shifted_logits - log(sum(exp(shifted_logits))) |
| 84 | ? xla::Sub(shifted_logits, xla::Log(sum), batch_dims) |
| 85 | // softmax = exp(shifted_logits) / sum(exp(shifted_logits)) |
| 86 | : xla::Div(exp_shifted, sum, batch_dims); |
| 87 | } |
| 88 | ctx->SetOutput(0, softmax); |
| 89 | } |
| 90 | |
| 91 | private: |
| 92 | bool log_; |
nothing calls this directly
no test coverage detected