[XLA:GPU] rename thunk_checksum_tracing_pass to thunk_buffer_debug_pass

It is pure mechanical move cl.

The goal is to use the pass for all the buffer debug checks. We have checksum and nan_counter kernels at the moment.

PiperOrigin-RevId: 825602375

third_party/xla/xla/backends/gpu/runtime/BUILD

@@ -2878,9 +2878,9 @@ xla_test(
 )
 
 cc_library(
-    name = "thunk_checksum_tracing_pass",
-    srcs = ["thunk_checksum_tracing_pass.cc"],
-    hdrs = ["thunk_checksum_tracing_pass.h"],
+    name = "thunk_buffer_debug_pass",
+    srcs = ["thunk_buffer_debug_pass.cc"],
+    hdrs = ["thunk_buffer_debug_pass.h"],
     deps = [
         ":buffers_checksum_thunk",
         ":custom_call_thunk",
@@ -2911,15 +2911,15 @@ cc_library(
 )
 
 xla_cc_test(
-    name = "thunk_checksum_tracing_pass_test",
-    srcs = ["thunk_checksum_tracing_pass_test.cc"],
+    name = "thunk_buffer_debug_pass_test",
+    srcs = ["thunk_buffer_debug_pass_test.cc"],
     deps = [
         ":buffers_checksum_thunk",
         ":custom_call_thunk",
         ":sequential_thunk",
         ":thunk",
+        ":thunk_buffer_debug_pass",
         ":thunk_buffer_id",
-        ":thunk_checksum_tracing_pass",
         ":thunk_id",
         ":thunk_pass_pipeline",
         "//xla:literal_util",

third_party/xla/xla/backends/gpu/runtime/thunk_buffer_debug_pass.cc

@@ -13,7 +13,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#include "xla/backends/gpu/runtime/thunk_checksum_tracing_pass.h"
+#include "xla/backends/gpu/runtime/thunk_buffer_debug_pass.h"
 
 #include <cstddef>
 #include <cstring>
@@ -165,12 +165,12 @@ XLA_FFI_DEFINE_HANDLER_SYMBOL(
 
 }  // namespace
 
-absl::StatusOr<bool> ThunkChecksumTracingPass::Run(
+absl::StatusOr<bool> ThunkBufferDebugPass::Run(
     SequentialThunk* root_thunk, const DebugOptions& debug_options,
     const HloModule* absl_nullable hlo_module,
     const se::DeviceDescription& device_info,
     ThunkPassBufferAllocator& allocator) {
-  VLOG(1) << "ThunkChecksumTracingPass running";
+  VLOG(1) << "ThunkBufferDebugPass running";
   if (hlo_module == nullptr) {
     // We need the HLO module to dump the buffer debug log proto to a file. If
     // it's not available, there's no point in doing extra work.

third_party/xla/xla/backends/gpu/runtime/thunk_buffer_debug_pass.h

@@ -13,8 +13,8 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#ifndef XLA_BACKENDS_GPU_RUNTIME_THUNK_CHECKSUM_TRACING_PASS_H_
-#define XLA_BACKENDS_GPU_RUNTIME_THUNK_CHECKSUM_TRACING_PASS_H_
+#ifndef XLA_BACKENDS_GPU_RUNTIME_THUNK_BUFFER_DEBUG_PASS_H_
+#define XLA_BACKENDS_GPU_RUNTIME_THUNK_BUFFER_DEBUG_PASS_H_
 
 #include "absl/base/nullability.h"
 #include "absl/status/statusor.h"
@@ -27,12 +27,12 @@ limitations under the License.
 namespace xla {
 namespace gpu {
 
-// Adds checksum tracing to thunks.
-class ThunkChecksumTracingPass : public ThunkPassInterface {
+// Adds buffer debug tracing to thunks.
+class ThunkBufferDebugPass : public ThunkPassInterface {
  public:
-  ThunkChecksumTracingPass() = default;
+  ThunkBufferDebugPass() = default;
 
-  absl::string_view name() const override { return "thunk-checksum-tracing"; }
+  absl::string_view name() const override { return "thunk-buffer-debug"; }
 
   absl::StatusOr<bool> Run(SequentialThunk* root_thunk,
                            const DebugOptions& debug_options,
@@ -44,4 +44,4 @@ class ThunkChecksumTracingPass : public ThunkPassInterface {
 }  // namespace gpu
 }  // namespace xla
 
-#endif  // XLA_BACKENDS_GPU_RUNTIME_THUNK_CHECKSUM_TRACING_PASS_H_
+#endif  // XLA_BACKENDS_GPU_RUNTIME_THUNK_BUFFER_DEBUG_PASS_H_

third_party/xla/xla/backends/gpu/runtime/thunk_buffer_debug_pass_test.cc

@@ -13,7 +13,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#include "xla/backends/gpu/runtime/thunk_checksum_tracing_pass.h"
+#include "xla/backends/gpu/runtime/thunk_buffer_debug_pass.h"
 
 #include <cstdint>
 #include <memory>
@@ -85,7 +85,7 @@ class FakeThunk : public Thunk {
   BufferUses buffer_uses_;
 };
 
-TEST(ThunkChecksumTracingPassTest, IsNoOpWhenHloModuleIsNull) {
+TEST(ThunkBufferDebugPassTest, IsNoOpWhenHloModuleIsNull) {
   DebugOptions debug_options;
   debug_options.set_xla_gpu_experimental_enable_checksum_tracing_on_thunks(
       true);
@@ -101,7 +101,7 @@ TEST(ThunkChecksumTracingPassTest, IsNoOpWhenHloModuleIsNull) {
   auto root_thunk =
       std::make_unique<SequentialThunk>(Thunk::ThunkInfo(), std::move(thunks));
 
-  ThunkChecksumTracingPass pass;
+  ThunkBufferDebugPass pass;
   TF_ASSERT_OK_AND_ASSIGN(
       bool changed, pass.Run(root_thunk.get(), debug_options,
                              /*hlo_module=*/nullptr, device_info, allocator));
@@ -109,14 +109,14 @@ TEST(ThunkChecksumTracingPassTest, IsNoOpWhenHloModuleIsNull) {
   EXPECT_THAT(root_thunk->thunks(), ElementsAre(Pointer(fake_thunk_ptr)));
 }
 
-TEST(ThunkChecksumTracingPassTest, InsertsBuffersDebugChecksumThunks) {
+TEST(ThunkBufferDebugPassTest, InsertsBuffersDebugChecksumThunks) {
   static constexpr ThunkId kTestThunkId = ThunkId(123);
   DebugOptions debug_options;
   debug_options.set_xla_gpu_experimental_enable_checksum_tracing_on_thunks(
       true);
   se::DeviceDescription device_info;
   FakeThunkPassBufferAllocator allocator;
-  // The callbacks created by ThunkChecksumTracingPass require a HloModule with
+  // The callbacks created by ThunkBufferDebugPass require a HloModule with
   // a non-null entry computation.
   auto builder = HloComputation::Builder("entry");
   HloInstruction* root = builder.AddInstruction(
@@ -152,7 +152,7 @@ TEST(ThunkChecksumTracingPassTest, InsertsBuffersDebugChecksumThunks) {
   auto root_thunk =
       std::make_unique<SequentialThunk>(Thunk::ThunkInfo(), std::move(thunks));
 
-  ThunkChecksumTracingPass pass;
+  ThunkBufferDebugPass pass;
   TF_ASSERT_OK_AND_ASSIGN(bool changed,
                           pass.Run(root_thunk.get(), debug_options, &hlo_module,
                                    device_info, allocator));

third_party/xla/xla/service/gpu/BUILD

@@ -738,7 +738,7 @@ cc_library(
         "//xla/backends/gpu/runtime:nvshmem_collective_thunk",
         "//xla/backends/gpu/runtime:sequential_thunk",
         "//xla/backends/gpu/runtime:thunk",
-        "//xla/backends/gpu/runtime:thunk_checksum_tracing_pass",
+        "//xla/backends/gpu/runtime:thunk_buffer_debug_pass",
         "//xla/backends/gpu/runtime:thunk_pass_pipeline",
         "//xla/core/collectives:clique_key",
         "//xla/hlo/ir:hlo",

third_party/xla/xla/service/gpu/gpu_executable.cc

@@ -45,7 +45,7 @@ limitations under the License.
 #include "xla/backends/gpu/runtime/nvshmem_collective_thunk.h"
 #include "xla/backends/gpu/runtime/sequential_thunk.h"
 #include "xla/backends/gpu/runtime/thunk.h"
-#include "xla/backends/gpu/runtime/thunk_checksum_tracing_pass.h"
+#include "xla/backends/gpu/runtime/thunk_buffer_debug_pass.h"
 #include "xla/backends/gpu/runtime/thunk_pass_pipeline.h"
 #include "xla/core/collectives/clique_key.h"
 #include "xla/executable_run_options.h"
@@ -177,7 +177,7 @@ static absl::Status RunThunkPasses(const DebugOptions& debug_options,
                                    ThunkPassBufferAllocator& allocator) {
   ThunkPassPipeline pipeline("thunk-passes");
   if (debug_options.xla_gpu_experimental_enable_checksum_tracing_on_thunks()) {
-    pipeline.AddPass(std::make_unique<ThunkChecksumTracingPass>());
+    pipeline.AddPass(std::make_unique<ThunkBufferDebugPass>());
   }
   if (debug_options.xla_gpu_experimental_enable_command_buffer_on_thunks()) {
     pipeline.AddPass(std::make_unique<CommandBufferConversionPass>(