[xla:ffi] NFC: Remove function name from XLA_FFI binding

FFI module itself defines under what name FFI binding should be exported to the Xla runtime and ffi binding name was not really used anywhere.

PiperOrigin-RevId: 501955340

Author: ezhulenev

xla/runtime/ffi/ffi_api.h

@@ -146,12 +146,11 @@ class Ffi {
  public:
   virtual ~Ffi() = default;
 
-  virtual std::string_view name() const = 0;
   virtual XLA_FFI_Error* operator()(const XLA_FFI_Api* api,
                                     XLA_FFI_ExecutionContext* ctx, void** args,
                                     void** attrs, void** rets) const = 0;
 
-  static FfiBinding<> Bind(std::string name);
+  static FfiBinding<> Binding();
 
   template <typename T>
   static bool Isa(const XLA_FFI_Api* api, XLA_FFI_TypeId type_id);
@@ -510,32 +509,29 @@ class FfiBinding {
 
   template <typename Fn>
   std::unique_ptr<FfiHandler<Fn, Ts...>> To(Fn fn) {
-    return std::unique_ptr<FfiHandler<Fn, Ts...>>(new FfiHandler<Fn, Ts...>(
-        std::forward<Fn>(fn), std::move(name_), std::move(attrs_)));
+    return std::unique_ptr<FfiHandler<Fn, Ts...>>(
+        new FfiHandler<Fn, Ts...>(std::forward<Fn>(fn), std::move(attrs_)));
   }
 
  private:
   template <typename...>
   friend class FfiBinding;
   friend class Ffi;
 
-  explicit FfiBinding(std::string name) : name_(std::move(name)) {
+  explicit FfiBinding() {
     static_assert(sizeof...(Ts) == 0, "ffi arguments must be empty");
   }
 
   template <typename... TTs>
   FfiBinding(FfiBinding<TTs...>&& other)  // NOLINT
-      : name_(std::move(other.name_)), attrs_(std::move(other.attrs_)) {}
+      : attrs_(std::move(other.attrs_)) {}
 
   FfiBinding(FfiBinding&) = delete;
 
-  std::string name_;                // ffi name
   std::vector<std::string> attrs_;  // names of bound attributes
 };
 
-inline FfiBinding<> Ffi::Bind(std::string name) {
-  return FfiBinding<>(std::move(name));
-}
+inline FfiBinding<> Ffi::Binding() { return FfiBinding<>(); }
 
 //===----------------------------------------------------------------------===//
 // Helpers for decoding opaque arguments and attributes' memory.
@@ -798,8 +794,6 @@ class FfiHandler : public Ffi {
   }
 
  public:
-  std::string_view name() const final { return name_; }
-
   XLA_FFI_Error* operator()(const XLA_FFI_Api* api,
                             XLA_FFI_ExecutionContext* ctx, void** args,
                             void** attrs, void** rets) const final {
@@ -867,9 +861,8 @@ class FfiHandler : public Ffi {
     return ToError(api, FfiStatus::Ok());
   }
 
-  FfiHandler(Fn fn, std::string name, std::vector<std::string> attrs)
+  FfiHandler(Fn fn, std::vector<std::string> attrs)
       : fn_(std::move(fn)),
-        name_(std::move(name)),
         attrs_(std::move(attrs)),
         attrs_idx_(attrs_.size()) {
     // Sort attributes names.
@@ -886,7 +879,6 @@ class FfiHandler : public Ffi {
 
   Fn fn_;
 
-  std::string name_;
   std::vector<std::string> attrs_;
 
   // A mapping from the attribute index to its index in the lexicographically

xla/runtime/ffi_test.cc

@@ -156,7 +156,7 @@ struct TestModule : public ffi::StatefulModule<TestModuleState> {
   // Function that tests that we can successfully decode various kinds of
   // attributes attached to custom calls.
   XLA_FFI_DEFINE_FUNCTION(FFI_AttrsDecoding, AttrsDecoding,
-                          ffi::Ffi::Bind("ffi.attrs")
+                          ffi::Ffi::Binding()
                               .State<TestModuleState>()  // state
                               .Attr<std::string_view>("str")
                               .Attr<float>("f32")
@@ -173,7 +173,7 @@ struct TestModule : public ffi::StatefulModule<TestModuleState> {
   // Function that tests that we can successfully decode various kinds of
   // arguments passed to custom calls.
   XLA_FFI_DEFINE_FUNCTION(FFI_Fill, Fill,
-                          ffi::Ffi::Bind("ffi.fill")
+                          ffi::Ffi::Binding()
                               .State<TestModuleState>()  // state
                               .Arg<int32_t>()            // arg0
                               .Arg<ffi::BufferArg>()     // arg1

xla/service/gpu/custom_call_test.cc

@@ -417,12 +417,11 @@ struct TestFfiModule : ffi::StatelessModule {
             {{"ffi.always_fail", FFI_AlwaysFail}, {"ffi.memcpy", FFI_Memcpy}}) {
   }
 
-  XLA_FFI_DEFINE_FUNCTION(
-      FFI_AlwaysFail, AlwaysFail,
-      ffi::Ffi::Bind("ffi.always_fail").Arg<ffi::StridedBufferArg>());
+  XLA_FFI_DEFINE_FUNCTION(FFI_AlwaysFail, AlwaysFail,
+                          ffi::Ffi::Binding().Arg<ffi::StridedBufferArg>());
 
   XLA_FFI_DEFINE_FUNCTION(FFI_Memcpy, Memcpy,
-                          ffi::Ffi::Bind("ffi.memcpy")
+                          ffi::Ffi::Binding()
                               .Stream<se::gpu::GpuStreamHandle>()
                               .Arg<ffi::StridedBufferArg>()
                               .Arg<ffi::StridedBufferArg>());

xla/tests/custom_call_test.cc

@@ -371,7 +371,7 @@ struct TestFfiModule : ffi::StatelessModule {
                         {{"ffi.add_const", FFI_AddConst}}) {}
 
   XLA_FFI_DEFINE_FUNCTION(FFI_AddConst, AddConst,
-                          ffi::Ffi::Bind("ffi.add_const")
+                          ffi::Ffi::Binding()
                               .Arg<ffi::StridedBufferArg>()
                               .Arg<ffi::StridedBufferArg>()
                               .Attr<float>("cst"));

xla/tests/xla_ffi_test.cc

@@ -57,7 +57,7 @@ struct TestModule : public runtime::ffi::StatefulModule<TestModuleState> {
 
   // XLA runtime binding for the C++ function.
   XLA_FFI_DEFINE_FUNCTION(FFI_Impl, Impl,
-                          Ffi::Bind("test.ffi")
+                          Ffi::Binding()
                               .State<TestModuleState>()
                               .Arg<StridedBufferArg>()
                               .Arg<StridedBufferArg>()