How to write custom XLA op for tensorflow
@pranavladkat,
Have you got the chance to have a look at the below official document link. Please try and let me know if it serves the purpose. 
https://www.tensorflow.org/xla/custom_call
Thank you! Hi, I have followed it, however this guide lacks several steps such as compilation and fails to provide an end to end working example. I tried following how tensorflow implements its ops by looking at the code, but I'm not able to get past op registration error which I pasted above. It'll be great help if that can be clarified.  Hi, I'm simply looking for better documentation or an end-to-end example which demonstrate writing custom XLA op for TF. Any pointers in the right direction are appreciated! It appears that we require native tf op to be registered along with XLA op to address this issue. This is not mentioned in the documentation.
<details><summary>Click to expand!</summary> 
 
 ### Issue Type

Documentation Feature Request

### Have you reproduced the bug with TF nightly?

Yes

### Source

binary

### Tensorflow Version

2.12.0

### Custom Code

Yes

### OS Platform and Distribution

Linux

### Mobile device

_No response_

### Python version

3.10

### Bazel version

_No response_

### GCC/Compiler version

_No response_

### CUDA/cuDNN version

_No response_

### GPU model and memory

_No response_

### Current Behaviour?

Hi,

I have custom tensorflow ops implemented in C++ [using this guide](https://www.tensorflow.org/guide/create_op) which are invoked during training. They work fine as long as XLA is not used, however I'm unable to get them working (and even a dummy custom xla op). 

The test op I have written is mentioned below and compiled by following the guide linked above. The compilation succeeds but I get following runtime error when the module with custom op is imported:

```
Non-OK-status: lookup_status status: NOT_FOUND: Op type not registered 'XlaTestOp' in binary running on <hostname>. Make sure the Op and Kernel are registered in the binary running in this process. 
```

I'm not able to find any more info how to resolve this issue. 



### Standalone code to reproduce the issue

```shell
C++ XLA op implementation:



#include "tensorflow/compiler/tf2xla/shape_util.h"
#include "tensorflow/compiler/tf2xla/xla_compiler.h"
#include "tensorflow/compiler/tf2xla/xla_op_kernel.h"
#include "tensorflow/compiler/tf2xla/xla_op_registry.h"
#include "tensorflow/compiler/xla/client/xla_builder.h"
#include "tensorflow/compiler/xla/service/custom_call_target_registry.h"
#include "tensorflow/compiler/xla/service/hlo.pb.h"
#include "tensorflow/compiler/xla/shape.h"
#include "tensorflow/compiler/xla/xla_data.pb.h"
#include "tensorflow/core/framework/op_kernel.h"
#include "tensorflow/core/lib/hash/hash.h"
#include "tensorflow/core/platform/human_readable_json.h"

void test_op(cudaStream_t stream, void** buffers, const char* opaque,
                     size_t opaque_len) {
  std::cout << "executing test_op()\n";
}



class XLATestOp : public XlaOpKernel {
public:
  explicit XLATestOp(OpKernelConstruction* ctx) : XlaOpKernel(ctx) { }

  void Compile(XlaOpKernelContext* ctx) override {

    ::xla::XlaBuilder* const builder = ctx->builder();
    ::xla::XlaOp input = ctx->Input(0);

    // use output aliasing to reuse input buffer for output
    std::vector<std::pair<::xla::ShapeIndex,
        std::pair<int64, ::xla::ShapeIndex>>> output_operand_aliasing = {
        {::xla::ShapeIndex{}, {0, ::xla::ShapeIndex{}}}
    };

    ::xla::Shape output_shape = builder->GetShape(input).value();

    ::xla::XlaOp output_op = builder->ReportErrorOrReturn(
        ::xla::CustomCall(
            builder,
            "test_op",
            {input},
            output_shape,
            /*opaque=*/ "",
            /*has_side_effect=*/false,
            output_operand_aliasing,
            /*literal=*/nullptr
        )
    );

    ctx->SetOutput(0, output_op);
  }
};


REGISTER_XLA_OP(Name("XlaTestOp"), XLATestOp);
XLA_REGISTER_CUSTOM_CALL_TARGET(test_op, "CUDA");
```
```


### Relevant log output

_No response_</details>