[XLA:GPU] Change SDC names to more descriptive ones

And dump the log proto into file called buffer_debug_log rather than sdc_log

Changes to implementation details:
- Renames:
  - SdcLogProto -> BufferDebugLogProto
  - SdcLog -> BufferDebugLog
  - SdcBufferId -> ThunkBufferId
  - SdcThunk -> BuffersChecksumThunk
  - SdcXorChecksumKernel -> BufferDebugXorChecksumKernel
- move BufferDebugLog to stream_executor/gpu from stream_executor/cuda as it's not CUDA-specific

PiperOrigin-RevId: 820186034

Author: Marcin Radomski

xla/backends/gpu/runtime/BUILD

@@ -2895,11 +2895,11 @@ cc_library(
     srcs = ["thunk_checksum_tracing_pass.cc"],
     hdrs = ["thunk_checksum_tracing_pass.h"],
     deps = [
+        ":buffers_checksum_thunk",
         ":custom_call_thunk",
-        ":sdc_buffer_id",
-        ":sdc_thunk",
         ":sequential_thunk",
         ":thunk",
+        ":thunk_buffer_id",
         ":thunk_pass_pipeline",
         "//xla:shape_util",
         "//xla:xla_data_proto_cc",
@@ -2911,7 +2911,7 @@ cc_library(
         "//xla/service:dump",
         "//xla/stream_executor:device_description",
         "//xla/stream_executor:stream",
-        "//xla/stream_executor/cuda:sdc_log",
+        "//xla/stream_executor/gpu:buffer_debug_log",
         "//xla/tsl/platform:statusor",
         "@com_google_absl//absl/base:nullability",
         "@com_google_absl//absl/container:flat_hash_map",
@@ -2927,11 +2927,11 @@ xla_cc_test(
     name = "thunk_checksum_tracing_pass_test",
     srcs = ["thunk_checksum_tracing_pass_test.cc"],
     deps = [
+        ":buffers_checksum_thunk",
         ":custom_call_thunk",
-        ":sdc_buffer_id",
-        ":sdc_thunk",
         ":sequential_thunk",
         ":thunk",
+        ":thunk_buffer_id",
         ":thunk_checksum_tracing_pass",
         ":thunk_id",
         ":thunk_pass_pipeline",
@@ -2994,21 +2994,21 @@ xla_test(
 )
 
 cc_library(
-    name = "sdc_thunk",
-    srcs = ["sdc_thunk.cc"],
-    hdrs = ["sdc_thunk.h"],
+    name = "buffers_checksum_thunk",
+    srcs = ["buffers_checksum_thunk.cc"],
+    hdrs = ["buffers_checksum_thunk.h"],
     deps = [
-        ":sdc_buffer_id",
         ":thunk",
+        ":thunk_buffer_id",
         "//xla/service:buffer_assignment",
         "//xla/stream_executor:device_memory",
         "//xla/stream_executor:launch_dim",
         "//xla/stream_executor:stream_executor_h",
         "//xla/stream_executor/cuda:cuda_compute_capability",
         "//xla/stream_executor/cuda:cuda_platform_id",
-        "//xla/stream_executor/cuda:sdc_log",
+        "//xla/stream_executor/gpu:buffer_debug_log",
+        "//xla/stream_executor/gpu:buffer_debug_xor_checksum_kernel",
         "//xla/stream_executor/gpu:gpu_kernel_registry",
-        "//xla/stream_executor/gpu:sdc_xor_checksum_kernel",
         "//xla/tsl/platform:errors",
         "//xla/tsl/platform:statusor",
         "@com_google_absl//absl/container:flat_hash_map",
@@ -3020,18 +3020,18 @@ cc_library(
 )
 
 xla_test(
-    name = "sdc_thunk_test",
-    srcs = ["sdc_thunk_test.cc"],
+    name = "buffers_checksum_thunk_test",
+    srcs = ["buffers_checksum_thunk_test.cc"],
     backends = ["gpu"],
     tags = [
         "cuda-only",
         "gpu",
     ],
     deps = [
-        ":sdc_buffer_id",
-        ":sdc_log_structs",
-        ":sdc_thunk",
+        ":buffer_debug_log_structs",
+        ":buffers_checksum_thunk",
         ":thunk",
+        ":thunk_buffer_id",
         ":thunk_id",
         "//xla/service:buffer_assignment",
         "//xla/service:executable",
@@ -3042,21 +3042,21 @@ xla_test(
         "//xla/stream_executor:platform_manager",
         "//xla/stream_executor:stream",
         "//xla/stream_executor:stream_executor_memory_allocator",
-        "//xla/stream_executor/cuda:sdc_log",
+        "//xla/stream_executor/gpu:buffer_debug_log",
         "//xla/tsl/lib/core:status_test_util",
         "//xla/tsl/platform:statusor",
         "@com_google_googletest//:gtest_main",
     ],
 )
 
 tf_proto_library(
-    name = "sdc_proto",
-    srcs = ["sdc.proto"],
+    name = "buffer_debug_log_proto",
+    srcs = ["buffer_debug_log.proto"],
 )
 
 cc_library(
-    name = "sdc_buffer_id",
-    hdrs = ["sdc_buffer_id.h"],
+    name = "thunk_buffer_id",
+    hdrs = ["thunk_buffer_id.h"],
     compatible_with = get_compatible_with_portable(),
     deps = [
         ":thunk_id",
@@ -3067,10 +3067,10 @@ cc_library(
 )
 
 xla_cc_test(
-    name = "sdc_buffer_id_test",
-    srcs = ["sdc_buffer_id_test.cc"],
+    name = "thunk_buffer_id_test",
+    srcs = ["thunk_buffer_id_test.cc"],
     deps = [
-        ":sdc_buffer_id",
+        ":thunk_buffer_id",
         ":thunk_id",
         "//xla/tsl/platform:statusor",
         "@com_google_absl//absl/status",
@@ -3080,10 +3080,10 @@ xla_cc_test(
 )
 
 cc_library(
-    name = "sdc_log_structs",
-    hdrs = ["sdc_log_structs.h"],
+    name = "buffer_debug_log_structs",
+    hdrs = ["buffer_debug_log_structs.h"],
     compatible_with = get_compatible_with_portable(),
     deps = [
-        ":sdc_buffer_id",
+        ":thunk_buffer_id",
     ],
 )

xla/backends/gpu/runtime/sdc.proto …kends/gpu/runtime/buffer_debug_log.protoxla/backends/gpu/runtime/sdc.proto renamed to xla/backends/gpu/runtime/buffer_debug_log.proto xla/backends/gpu/runtime/sdc.proto renamed to xla/backends/gpu/runtime/buffer_debug_log.proto

@@ -17,7 +17,7 @@ syntax = "proto3";
 
 package xla.gpu;
 
-message SdcLogEntryProto {
+message BufferDebugLogEntryProto {
   // The ID of the thunk that produced this entry, as returned by
   // ThunkInfo::thunk_id().
   uint64 thunk_id = 1;
@@ -30,7 +30,8 @@ message SdcLogEntryProto {
   uint32 checksum = 3;
 }
 
-message SdcLogProto {
-  // The list of entries in the SDC log.
-  repeated SdcLogEntryProto entries = 1;
+// A dump of a `BufferDebugLog` contents.
+message BufferDebugLogProto {
+  // The list of entries in the log.
+  repeated BufferDebugLogEntryProto entries = 1;
 }

…a/backends/gpu/runtime/sdc_log_structs.h …s/gpu/runtime/buffer_debug_log_structs.hxla/backends/gpu/runtime/sdc_log_structs.h renamed to xla/backends/gpu/runtime/buffer_debug_log_structs.h xla/backends/gpu/runtime/sdc_log_structs.h renamed to xla/backends/gpu/runtime/buffer_debug_log_structs.h

@@ -13,57 +13,59 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#ifndef XLA_BACKENDS_GPU_RUNTIME_SDC_LOG_STRUCTS_H_
-#define XLA_BACKENDS_GPU_RUNTIME_SDC_LOG_STRUCTS_H_
+#ifndef XLA_BACKENDS_GPU_RUNTIME_BUFFER_DEBUG_LOG_STRUCTS_H_
+#define XLA_BACKENDS_GPU_RUNTIME_BUFFER_DEBUG_LOG_STRUCTS_H_
 
 #include <cstddef>
 #include <cstdint>
 #include <tuple>
 
-#include "xla/backends/gpu/runtime/sdc_buffer_id.h"
+#include "xla/backends/gpu/runtime/thunk_buffer_id.h"
 
 namespace xla::gpu {
 
-struct SdcLogEntry {
+struct BufferDebugLogEntry {
   // An ID that uniquely identifies a thunk and its specific input or output
   // buffer.
-  SdcBufferId entry_id;
+  ThunkBufferId entry_id;
   uint32_t checksum;
 
   template <typename Sink>
-  friend void AbslStringify(Sink& sink, const SdcLogEntry& entry) {
+  friend void AbslStringify(Sink& sink, const BufferDebugLogEntry& entry) {
     absl::Format(&sink, "{entry_id: %v, checksum: %u}", entry.entry_id,
                  entry.checksum);
   }
 
-  bool operator==(const SdcLogEntry& other) const {
+  bool operator==(const BufferDebugLogEntry& other) const {
     return std::tie(entry_id, checksum) ==
            std::tie(other.entry_id, other.checksum);
   }
 
-  bool operator!=(const SdcLogEntry& other) const { return !(*this == other); }
+  bool operator!=(const BufferDebugLogEntry& other) const {
+    return !(*this == other);
+  }
 };
 
 // The struct layout must match on both host and device.
-static_assert(_Alignof(SdcLogEntry) == _Alignof(uint32_t));
-static_assert(sizeof(SdcLogEntry) == sizeof(uint32_t) * 2);
-static_assert(offsetof(SdcLogEntry, entry_id) == 0);
-static_assert(offsetof(SdcLogEntry, checksum) == sizeof(uint32_t));
+static_assert(_Alignof(BufferDebugLogEntry) == _Alignof(uint32_t));
+static_assert(sizeof(BufferDebugLogEntry) == sizeof(uint32_t) * 2);
+static_assert(offsetof(BufferDebugLogEntry, entry_id) == 0);
+static_assert(offsetof(BufferDebugLogEntry, checksum) == sizeof(uint32_t));
 
-struct SdcLogHeader {
-  // The first entry in `SdcLogEntry` following the header that has not
+struct BufferDebugLogHeader {
+  // The first entry in `BufferDebugLogEntry` following the header that has not
   // been written to. May be bigger than `capacity` if the log was truncated.
   uint32_t write_idx;
-  // The number of `SdcLogEntry` structs the log can hold.
+  // The number of `BufferDebugLogEntry` structs the log can hold.
   uint32_t capacity;
 };
 
 // The struct layout must match on both host and device.
-static_assert(_Alignof(SdcLogHeader) == _Alignof(uint32_t));
-static_assert(sizeof(SdcLogHeader) == sizeof(uint32_t) * 2);
-static_assert(offsetof(SdcLogHeader, write_idx) == 0);
-static_assert(offsetof(SdcLogHeader, capacity) == sizeof(uint32_t));
+static_assert(_Alignof(BufferDebugLogHeader) == _Alignof(uint32_t));
+static_assert(sizeof(BufferDebugLogHeader) == sizeof(uint32_t) * 2);
+static_assert(offsetof(BufferDebugLogHeader, write_idx) == 0);
+static_assert(offsetof(BufferDebugLogHeader, capacity) == sizeof(uint32_t));
 
 }  // namespace xla::gpu
 
-#endif  // XLA_BACKENDS_GPU_RUNTIME_SDC_LOG_STRUCTS_H_
+#endif  // XLA_BACKENDS_GPU_RUNTIME_BUFFER_DEBUG_LOG_STRUCTS_H_

xla/backends/gpu/runtime/sdc_thunk.cc …ds/gpu/runtime/buffers_checksum_thunk.ccxla/backends/gpu/runtime/sdc_thunk.cc renamed to xla/backends/gpu/runtime/buffers_checksum_thunk.cc xla/backends/gpu/runtime/sdc_thunk.cc renamed to xla/backends/gpu/runtime/buffers_checksum_thunk.cc

@@ -13,7 +13,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#include "xla/backends/gpu/runtime/sdc_thunk.h"
+#include "xla/backends/gpu/runtime/buffers_checksum_thunk.h"
 
 #include <cstdint>
 #include <string>
@@ -24,10 +24,10 @@ limitations under the License.
 #include "xla/backends/gpu/runtime/thunk.h"
 #include "xla/stream_executor/cuda/cuda_compute_capability.h"
 #include "xla/stream_executor/cuda/cuda_platform_id.h"
-#include "xla/stream_executor/cuda/sdc_log.h"
 #include "xla/stream_executor/device_memory.h"
+#include "xla/stream_executor/gpu/buffer_debug_log.h"
+#include "xla/stream_executor/gpu/buffer_debug_xor_checksum_kernel.h"
 #include "xla/stream_executor/gpu/gpu_kernel_registry.h"
-#include "xla/stream_executor/gpu/sdc_xor_checksum_kernel.h"
 #include "xla/stream_executor/launch_dim.h"
 #include "xla/stream_executor/stream_executor.h"
 #include "xla/tsl/platform/errors.h"
@@ -37,63 +37,66 @@ namespace xla::gpu {
 
 namespace se = stream_executor;
 
-absl::Status SdcThunk::Initialize(const InitializeParams& params) {
+absl::Status BuffersDebugChecksumThunk::Initialize(
+    const InitializeParams& params) {
   if (params.executor->GetPlatform()->id() != se::cuda::kCudaPlatformId) {
-    VLOG(1) << "[SDC LOG] Not supported on non-CUDA platforms, skipping";
+    VLOG(1)
+        << "Buffer checksumming not supported on non-CUDA platforms, skipping";
     return absl::OkStatus();
   }
   if (!params.executor->GetDeviceDescription()
            .cuda_compute_capability()
            .IsAtLeastPascal()) {
-    VLOG(1) << "[SDC LOG] Not supported on CUDA architectures older than "
-               "Pascal due to missing atomic fetch_add with system scope, "
-               "skipping";
+    VLOG(1)
+        << "Buffer checksumming not supported on CUDA architectures older than "
+           "Pascal due to missing atomic fetch_add with system scope, skipping";
     return absl::OkStatus();
   }
 
   se::gpu::GpuKernelRegistry registry =
       se::gpu::GpuKernelRegistry::GetGlobalRegistry();
   TF_ASSIGN_OR_RETURN(
-      kernel_,
-      registry.LoadKernel<se::gpu::SdcXorChecksumKernel>(params.executor));
+      kernel_, registry.LoadKernel<se::gpu::BufferDebugXorChecksumKernel>(
+                   params.executor));
 
-  VLOG(1) << "[SDC LOG] SDC kernel loaded";
+  VLOG(1) << "Checksum kernel loaded";
   return absl::OkStatus();
 }
 
-absl::Status SdcThunk::ExecuteOnStream(const ExecuteParams& params) {
+absl::Status BuffersDebugChecksumThunk::ExecuteOnStream(
+    const ExecuteParams& params) {
   se::StreamExecutor* executor = params.stream->parent();
   if (!kernel_.has_value()) {
     // Initialize didn't load the kernel. This can happen when we're running on
     // an unsupported platform.
-    VLOG(1) << "[SDC LOG] SDC kernel not loaded, skipping";
+    VLOG(1) << "Checksum kernel not loaded, skipping";
     return absl::OkStatus();
   }
 
-  VLOG(1) << "[SDC LOG] SdcThunk::ExecuteOnStream";
+  VLOG(1) << "BuffersDebugChecksumThunk::ExecuteOnStream";
 
   const se::ThreadDim thread_dim(
       executor->GetDeviceDescription().threads_per_block_limit(), 1, 1);
 
   se::DeviceMemory<uint8_t> log_ptr(
       params.buffer_allocations->GetDeviceAddress(log_slice_));
-  se::cuda::SdcLog sdc_log =
-      se::cuda::SdcLog::FromDeviceMemoryUnchecked(log_ptr);
+  se::cuda::BufferDebugLog buffer_debug_log =
+      se::cuda::BufferDebugLog::FromDeviceMemoryUnchecked(log_ptr);
 
   for (const auto& [entry_id, buffer] : buffers_) {
     se::DeviceMemory<uint8_t> device_buffer(
         params.buffer_allocations->GetDeviceAddress(buffer));
 
-    TF_RETURN_IF_ERROR(
-        kernel_->Launch(thread_dim, se::BlockDim(1, 1, 1), params.stream,
-                        entry_id, device_buffer, device_buffer.size(),
-                        sdc_log.GetDeviceHeader(), sdc_log.GetDeviceEntries()));
+    TF_RETURN_IF_ERROR(kernel_->Launch(
+        thread_dim, se::BlockDim(1, 1, 1), params.stream, entry_id,
+        device_buffer, device_buffer.size(), buffer_debug_log.GetDeviceHeader(),
+        buffer_debug_log.GetDeviceEntries()));
   }
 
   return absl::OkStatus();
 }
 
-std::string SdcThunk::ToString(int indent) const {
+std::string BuffersDebugChecksumThunk::ToString(int indent) const {
   std::string result;
   absl::StrAppend(&result, ", buffers = ", buffers_.size());
   for (const auto& [buffer_id, buffer] : buffers_) {