| 26 | explicit CrossOp(OpKernelConstruction* context) : XlaOpKernel(context) {} |
| 27 | |
| 28 | void Compile(XlaOpKernelContext* ctx) override { |
| 29 | TensorShape in0_shape = ctx->InputShape(0); |
| 30 | TensorShape in1_shape = ctx->InputShape(1); |
| 31 | OP_REQUIRES(ctx, in0_shape == in1_shape, |
| 32 | errors::InvalidArgument("Both inputs must be of same shape: ", |
| 33 | in0_shape.DebugString(), " vs. ", |
| 34 | in1_shape.DebugString())); |
| 35 | OP_REQUIRES(ctx, in0_shape.dims() >= 1, |
| 36 | errors::InvalidArgument("Input must be at least 1D", |
| 37 | in0_shape.DebugString())); |
| 38 | |
| 39 | auto inner_dim = in0_shape.dim_size(in0_shape.dims() - 1); |
| 40 | OP_REQUIRES(ctx, inner_dim == 3, |
| 41 | errors::FailedPrecondition( |
| 42 | "Cross-products are only defined for 3-element vectors.")); |
| 43 | |
| 44 | // in0 is a [...,X,Y,Z,3] |
| 45 | // in1 is the same shape as in0 |
| 46 | // So slice 0 is: in0[...,:,:,:,0:1] |
| 47 | // So slice 1 is: in0[...,:,:,:,1:2] |
| 48 | // So slice 2 is: in0[...,:,:,:,2:3] |
| 49 | |
| 50 | std::vector<int64> starts(in0_shape.dims(), 0); |
| 51 | std::vector<int64> limits; |
| 52 | for (auto dim_size : in0_shape.dim_sizes()) { |
| 53 | limits.push_back(dim_size); |
| 54 | } |
| 55 | std::vector<int64> strides(in0_shape.dims(), 1); |
| 56 | |
| 57 | xla::XlaBuilder* b = ctx->builder(); |
| 58 | auto in0 = ctx->Input(0); |
| 59 | auto in1 = ctx->Input(1); |
| 60 | starts.back() = 0; |
| 61 | limits.back() = 1; |
| 62 | auto u1 = xla::Slice(in0, starts, limits, strides); |
| 63 | auto v1 = xla::Slice(in1, starts, limits, strides); |
| 64 | starts.back() = 1; |
| 65 | limits.back() = 2; |
| 66 | auto u2 = xla::Slice(in0, starts, limits, strides); |
| 67 | auto v2 = xla::Slice(in1, starts, limits, strides); |
| 68 | starts.back() = 2; |
| 69 | limits.back() = 3; |
| 70 | auto u3 = xla::Slice(in0, starts, limits, strides); |
| 71 | auto v3 = xla::Slice(in1, starts, limits, strides); |
| 72 | |
| 73 | auto s1 = xla::Sub(xla::Mul(u2, v3), xla::Mul(u3, v2)); |
| 74 | auto s2 = xla::Sub(xla::Mul(u3, v1), xla::Mul(u1, v3)); |
| 75 | auto s3 = xla::Sub(xla::Mul(u1, v2), xla::Mul(u2, v1)); |
| 76 | auto output = xla::ConcatInDim(b, {s1, s2, s3}, in0_shape.dims() - 1); |
| 77 | |
| 78 | ctx->SetOutput(0, output); |
| 79 | } |
| 80 | |
| 81 | private: |
| 82 | TF_DISALLOW_COPY_AND_ASSIGN(CrossOp); |
nothing calls this directly
no test coverage detected