[xla:gpu] NFC: Rename runtime3 to runtime

PiperOrigin-RevId: 606063012

third_party/xla/xla/service/gpu/BUILD

@@ -353,28 +353,28 @@ cc_library(
         "//xla/service/gpu/fusions:tiling_util",
         "//xla/service/gpu/kernels:custom_kernel",
         "//xla/service/gpu/kernels:topk_custom_kernel",
-        "//xla/service/gpu/runtime3:command_buffer_cmd",
-        "//xla/service/gpu/runtime3:command_buffer_cmd_emitter",
-        "//xla/service/gpu/runtime3:command_buffer_thunk",
-        "//xla/service/gpu/runtime3:conditional_thunk",
-        "//xla/service/gpu/runtime3:convolution_thunk",
-        "//xla/service/gpu/runtime3:copy_thunk",
-        "//xla/service/gpu/runtime3:custom_call_thunk",
-        "//xla/service/gpu/runtime3:fft_thunk",
-        "//xla/service/gpu/runtime3:fused_mha_thunk",
-        "//xla/service/gpu/runtime3:gemm_thunk",
-        "//xla/service/gpu/runtime3:infeed_thunk",
-        "//xla/service/gpu/runtime3:kernel_thunk",
-        "//xla/service/gpu/runtime3:nccl_all_gather_thunk",
-        "//xla/service/gpu/runtime3:nccl_all_reduce_thunk",
-        "//xla/service/gpu/runtime3:nccl_all_to_all_thunk",
-        "//xla/service/gpu/runtime3:norm_thunk",
-        "//xla/service/gpu/runtime3:outfeed_thunk",
-        "//xla/service/gpu/runtime3:replica_id_thunk",
-        "//xla/service/gpu/runtime3:send_recv_thunk",
-        "//xla/service/gpu/runtime3:sequential_thunk",
-        "//xla/service/gpu/runtime3:wait_for_streams_thunk",
-        "//xla/service/gpu/runtime3:while_thunk",
+        "//xla/service/gpu/runtime:command_buffer_cmd",
+        "//xla/service/gpu/runtime:command_buffer_cmd_emitter",
+        "//xla/service/gpu/runtime:command_buffer_thunk",
+        "//xla/service/gpu/runtime:conditional_thunk",
+        "//xla/service/gpu/runtime:convolution_thunk",
+        "//xla/service/gpu/runtime:copy_thunk",
+        "//xla/service/gpu/runtime:custom_call_thunk",
+        "//xla/service/gpu/runtime:fft_thunk",
+        "//xla/service/gpu/runtime:fused_mha_thunk",
+        "//xla/service/gpu/runtime:gemm_thunk",
+        "//xla/service/gpu/runtime:infeed_thunk",
+        "//xla/service/gpu/runtime:kernel_thunk",
+        "//xla/service/gpu/runtime:nccl_all_gather_thunk",
+        "//xla/service/gpu/runtime:nccl_all_reduce_thunk",
+        "//xla/service/gpu/runtime:nccl_all_to_all_thunk",
+        "//xla/service/gpu/runtime:norm_thunk",
+        "//xla/service/gpu/runtime:outfeed_thunk",
+        "//xla/service/gpu/runtime:replica_id_thunk",
+        "//xla/service/gpu/runtime:send_recv_thunk",
+        "//xla/service/gpu/runtime:sequential_thunk",
+        "//xla/service/gpu/runtime:wait_for_streams_thunk",
+        "//xla/service/gpu/runtime:while_thunk",
         "//xla/service/llvm_ir:buffer_assignment_util",
         "//xla/service/llvm_ir:dynamic_update_slice_util",
         "//xla/service/llvm_ir:fused_ir_emitter",
@@ -424,10 +424,10 @@ cc_library(
         "@local_tsl//tsl/protobuf:dnn_proto_cc",
     ] + if_gpu_is_configured([
         ":ir_emitter_triton",
-        "//xla/service/gpu/runtime3:cholesky_thunk",
-        "//xla/service/gpu/runtime3:cub_sort_thunk",
-        "//xla/service/gpu/runtime3:gpublas_lt_matmul_thunk",
-        "//xla/service/gpu/runtime3:triangular_solve_thunk",
+        "//xla/service/gpu/runtime:cholesky_thunk",
+        "//xla/service/gpu/runtime:cub_sort_thunk",
+        "//xla/service/gpu/runtime:gpublas_lt_matmul_thunk",
+        "//xla/service/gpu/runtime:triangular_solve_thunk",
     ]) + if_rocm_is_configured([
         "@local_config_rocm//rocm:rocm_headers",
     ]),
@@ -1243,7 +1243,7 @@ cc_library(
         "//xla/service:stream_pool",
         "//xla/service:xla_debug_info_manager",
         "//xla/service/gpu:backend_configs_cc",
-        "//xla/service/gpu/runtime3:annotation",
+        "//xla/service/gpu/runtime:annotation",
         "//xla/stream_executor",
         "//xla/stream_executor:device_description",
         "//xla/stream_executor:device_memory",
@@ -2283,7 +2283,7 @@ cc_library(
         "//xla:xla_data_proto_cc",
         "//xla/hlo/ir:hlo",
         "//xla/service:hlo_pass",
-        "//xla/service/gpu/runtime3:cub_sort_thunk",
+        "//xla/service/gpu/runtime:cub_sort_thunk",
         "@com_google_absl//absl/container:flat_hash_set",
         "@com_google_absl//absl/strings",
         "@com_google_absl//absl/types:span",
@@ -3183,9 +3183,9 @@ cc_library(
         "//xla/service:hlo_ordering",
         "//xla/service:hlo_proto_cc",
         "//xla/service:logical_buffer",
-        "//xla/service/gpu/runtime3:conditional_thunk",
-        "//xla/service/gpu/runtime3:sequential_thunk",
-        "//xla/service/gpu/runtime3:while_thunk",
+        "//xla/service/gpu/runtime:conditional_thunk",
+        "//xla/service/gpu/runtime:sequential_thunk",
+        "//xla/service/gpu/runtime:while_thunk",
         "//xla/service/llvm_ir:llvm_util",
         "//xla/stream_executor",
         "//xla/stream_executor:device_description",

third_party/xla/xla/service/gpu/compile_module_to_llvm_ir.cc

@@ -68,9 +68,9 @@ limitations under the License.
 #include "xla/service/gpu/ir_emitter_context.h"
 #include "xla/service/gpu/ir_emitter_unnested.h"
 #include "xla/service/gpu/metrics.h"
-#include "xla/service/gpu/runtime3/conditional_thunk.h"
-#include "xla/service/gpu/runtime3/sequential_thunk.h"
-#include "xla/service/gpu/runtime3/while_thunk.h"
+#include "xla/service/gpu/runtime/conditional_thunk.h"
+#include "xla/service/gpu/runtime/sequential_thunk.h"
+#include "xla/service/gpu/runtime/while_thunk.h"
 #include "xla/service/gpu/thunk.h"
 #include "xla/service/hlo_dataflow_analysis.h"
 #include "xla/service/hlo_ordering.h"

third_party/xla/xla/service/gpu/fusions/BUILD

@@ -45,7 +45,7 @@ cc_library(
         "//xla/service:buffer_assignment",
         "//xla/service/gpu:ir_emitter_context",
         "//xla/service/gpu:thunk",
-        "//xla/service/gpu/runtime3:copy_thunk",
+        "//xla/service/gpu/runtime:copy_thunk",
         "@llvm-project//mlir:IR",
     ],
 )
@@ -81,9 +81,9 @@ cc_library(
         "//xla/service/gpu:thunk",
         "//xla/service/gpu/kernels:custom_kernel",
         "//xla/service/gpu/kernels:custom_kernel_fusion",
-        "//xla/service/gpu/runtime3:custom_call_thunk",
-        "//xla/service/gpu/runtime3:gemm_thunk",
-        "//xla/service/gpu/runtime3:kernel_thunk",
+        "//xla/service/gpu/runtime:custom_call_thunk",
+        "//xla/service/gpu/runtime:gemm_thunk",
+        "//xla/service/gpu/runtime:kernel_thunk",
         "@com_google_absl//absl/log",
         "@com_google_absl//absl/status",
         "@com_google_absl//absl/strings",
@@ -143,7 +143,7 @@ cc_library(
         "//xla/service/gpu:thunk",
         "//xla/service/gpu/model:indexing_analysis",
         "//xla/service/gpu/model:indexing_map",
-        "//xla/service/gpu/runtime3:kernel_thunk",
+        "//xla/service/gpu/runtime:kernel_thunk",
         "//xla/service/llvm_ir:ir_array",
         "//xla/service/llvm_ir:llvm_util",
         "//xla/stream_executor:device_description",
@@ -330,7 +330,7 @@ cc_library(
         "//xla/service/gpu:launch_dimensions",
         "//xla/service/gpu:matmul_utils",
         "//xla/service/gpu:triton_fusion_analysis",
-        "//xla/service/gpu/runtime3:kernel_thunk",
+        "//xla/service/gpu/runtime:kernel_thunk",
         "//xla/service/llvm_ir:llvm_util",
         "@com_google_absl//absl/log",
         "@com_google_absl//absl/log:check",
@@ -373,7 +373,7 @@ cc_library(
         "//xla/service/gpu:gpu_executable",
         "//xla/service/gpu:ir_emitter_context",
         "//xla/service/gpu:thunk",
-        "//xla/service/gpu/runtime3:memset_thunk",
+        "//xla/service/gpu/runtime:memset_thunk",
         "@com_google_absl//absl/algorithm:container",
         "@com_google_absl//absl/types:span",
         "@llvm-project//mlir:IR",
@@ -411,7 +411,7 @@ cc_library(
         "//xla/service/gpu:reduction_utils",
         "//xla/service/gpu:target_util",
         "//xla/service/gpu:thunk",
-        "//xla/service/gpu/runtime3:kernel_thunk",
+        "//xla/service/gpu/runtime:kernel_thunk",
         "//xla/service/llvm_ir:fused_ir_emitter",
         "//xla/service/llvm_ir:ir_array",
         "//xla/service/llvm_ir:kernel_support_library",

third_party/xla/xla/service/gpu/fusions/copy.cc

@@ -20,7 +20,7 @@ limitations under the License.
 #include "xla/mlir_hlo/lhlo/IR/lhlo_ops.h"
 #include "xla/service/gpu/fusions/fusion_emitter.h"
 #include "xla/service/gpu/ir_emitter_context.h"
-#include "xla/service/gpu/runtime3/copy_thunk.h"
+#include "xla/service/gpu/runtime/copy_thunk.h"
 #include "xla/service/gpu/thunk.h"
 #include "xla/statusor.h"
 

third_party/xla/xla/service/gpu/fusions/custom.cc

@@ -54,9 +54,9 @@ limitations under the License.
 #include "xla/service/gpu/kernels/custom_kernel.h"
 #include "xla/service/gpu/kernels/custom_kernel_fusion.h"
 #include "xla/service/gpu/matmul_utils.h"
-#include "xla/service/gpu/runtime3/custom_call_thunk.h"
-#include "xla/service/gpu/runtime3/gemm_thunk.h"
-#include "xla/service/gpu/runtime3/kernel_thunk.h"
+#include "xla/service/gpu/runtime/custom_call_thunk.h"
+#include "xla/service/gpu/runtime/gemm_thunk.h"
+#include "xla/service/gpu/runtime/kernel_thunk.h"
 #include "xla/service/gpu/thunk.h"
 #include "xla/shape.h"
 #include "xla/shape_util.h"

third_party/xla/xla/service/gpu/fusions/fusion_emitter.cc

@@ -50,7 +50,7 @@ limitations under the License.
 #include "xla/service/gpu/kernel_reuse_cache.h"
 #include "xla/service/gpu/launch_dimensions.h"
 #include "xla/service/gpu/model/indexing_map.h"
-#include "xla/service/gpu/runtime3/kernel_thunk.h"
+#include "xla/service/gpu/runtime/kernel_thunk.h"
 #include "xla/service/gpu/target_util.h"
 #include "xla/service/llvm_ir/ir_array.h"
 #include "xla/service/llvm_ir/llvm_util.h"

third_party/xla/xla/service/gpu/fusions/reduction.cc

@@ -71,7 +71,7 @@ limitations under the License.
 #include "xla/service/gpu/launch_dimensions.h"
 #include "xla/service/gpu/parallel_loop_emitter.h"
 #include "xla/service/gpu/reduction_utils.h"
-#include "xla/service/gpu/runtime3/kernel_thunk.h"
+#include "xla/service/gpu/runtime/kernel_thunk.h"
 #include "xla/service/gpu/target_util.h"
 #include "xla/service/gpu/thunk.h"
 #include "xla/service/llvm_ir/fused_ir_emitter.h"

third_party/xla/xla/service/gpu/fusions/thunk_util.cc

@@ -30,7 +30,7 @@ limitations under the License.
 #include "xla/literal.h"
 #include "xla/service/buffer_assignment.h"
 #include "xla/service/gpu/ir_emitter_context.h"
-#include "xla/service/gpu/runtime3/memset_thunk.h"
+#include "xla/service/gpu/runtime/memset_thunk.h"
 #include "xla/service/gpu/thunk.h"
 #include "xla/shape.h"
 #include "xla/shape_util.h"

third_party/xla/xla/service/gpu/fusions/triton.cc

@@ -36,7 +36,7 @@ limitations under the License.
 #include "xla/service/gpu/kernel_reuse_cache.h"
 #include "xla/service/gpu/launch_dimensions.h"
 #include "xla/service/gpu/matmul_utils.h"
-#include "xla/service/gpu/runtime3/kernel_thunk.h"
+#include "xla/service/gpu/runtime/kernel_thunk.h"
 #include "xla/service/gpu/triton_fusion_analysis.h"
 #include "xla/service/llvm_ir/llvm_util.h"
 #include "xla/statusor.h"

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

@@ -46,7 +46,7 @@ limitations under the License.
 #include "xla/service/gpu/gpu_executable_run_options.h"
 #include "xla/service/gpu/nccl_clique.h"
 #include "xla/service/gpu/nccl_clique_key.h"
-#include "xla/service/gpu/runtime3/annotation.h"
+#include "xla/service/gpu/runtime/annotation.h"
 #include "xla/service/gpu/stream_executor_util.h"
 #include "xla/service/gpu/thunk.h"
 #include "xla/service/hlo_parser.h"

third_party/xla/xla/service/gpu/gpu_executable.h

@@ -38,7 +38,7 @@ limitations under the License.
 #include "xla/service/executable.h"
 #include "xla/service/gpu/buffer_allocations.h"
 #include "xla/service/gpu/ir_emission_utils.h"
-#include "xla/service/gpu/runtime3/annotation.h"
+#include "xla/service/gpu/runtime/annotation.h"
 #include "xla/service/gpu/thunk.h"
 #include "xla/service/hlo_execution_profile.h"
 #include "xla/service/hlo_module_config.h"

third_party/xla/xla/service/gpu/gpu_sort_rewriter.cc

@@ -32,7 +32,7 @@ limitations under the License.
 #include "xla/hlo/ir/hlo_instructions.h"
 #include "xla/hlo/ir/hlo_module.h"
 #include "xla/service/gpu/cublas_cudnn.h"
-#include "xla/service/gpu/runtime3/cub_sort_thunk.h"
+#include "xla/service/gpu/runtime/cub_sort_thunk.h"
 #include "xla/shape.h"
 #include "xla/shape_util.h"
 #include "xla/statusor.h"

third_party/xla/xla/service/gpu/ir_emitter_unnested.cc

@@ -121,28 +121,28 @@ limitations under the License.
 #include "xla/service/gpu/nccl_recv_thunk.h"
 #include "xla/service/gpu/nccl_send_thunk.h"
 #include "xla/service/gpu/parallel_loop_emitter.h"
-#include "xla/service/gpu/runtime3/command_buffer_cmd.h"
-#include "xla/service/gpu/runtime3/command_buffer_cmd_emitter.h"
-#include "xla/service/gpu/runtime3/command_buffer_thunk.h"
-#include "xla/service/gpu/runtime3/conditional_thunk.h"
-#include "xla/service/gpu/runtime3/convolution_thunk.h"
-#include "xla/service/gpu/runtime3/copy_thunk.h"
-#include "xla/service/gpu/runtime3/custom_call_thunk.h"
-#include "xla/service/gpu/runtime3/fft_thunk.h"
-#include "xla/service/gpu/runtime3/fused_mha_thunk.h"
-#include "xla/service/gpu/runtime3/gemm_thunk.h"
-#include "xla/service/gpu/runtime3/infeed_thunk.h"
-#include "xla/service/gpu/runtime3/kernel_thunk.h"
-#include "xla/service/gpu/runtime3/nccl_all_gather_thunk.h"
-#include "xla/service/gpu/runtime3/nccl_all_reduce_thunk.h"
-#include "xla/service/gpu/runtime3/nccl_all_to_all_thunk.h"
-#include "xla/service/gpu/runtime3/norm_thunk.h"
-#include "xla/service/gpu/runtime3/outfeed_thunk.h"
-#include "xla/service/gpu/runtime3/replica_id_thunk.h"
-#include "xla/service/gpu/runtime3/send_recv_thunk.h"
-#include "xla/service/gpu/runtime3/sequential_thunk.h"
-#include "xla/service/gpu/runtime3/wait_for_streams_thunk.h"
-#include "xla/service/gpu/runtime3/while_thunk.h"
+#include "xla/service/gpu/runtime/command_buffer_cmd.h"
+#include "xla/service/gpu/runtime/command_buffer_cmd_emitter.h"
+#include "xla/service/gpu/runtime/command_buffer_thunk.h"
+#include "xla/service/gpu/runtime/conditional_thunk.h"
+#include "xla/service/gpu/runtime/convolution_thunk.h"
+#include "xla/service/gpu/runtime/copy_thunk.h"
+#include "xla/service/gpu/runtime/custom_call_thunk.h"
+#include "xla/service/gpu/runtime/fft_thunk.h"
+#include "xla/service/gpu/runtime/fused_mha_thunk.h"
+#include "xla/service/gpu/runtime/gemm_thunk.h"
+#include "xla/service/gpu/runtime/infeed_thunk.h"
+#include "xla/service/gpu/runtime/kernel_thunk.h"
+#include "xla/service/gpu/runtime/nccl_all_gather_thunk.h"
+#include "xla/service/gpu/runtime/nccl_all_reduce_thunk.h"
+#include "xla/service/gpu/runtime/nccl_all_to_all_thunk.h"
+#include "xla/service/gpu/runtime/norm_thunk.h"
+#include "xla/service/gpu/runtime/outfeed_thunk.h"
+#include "xla/service/gpu/runtime/replica_id_thunk.h"
+#include "xla/service/gpu/runtime/send_recv_thunk.h"
+#include "xla/service/gpu/runtime/sequential_thunk.h"
+#include "xla/service/gpu/runtime/wait_for_streams_thunk.h"
+#include "xla/service/gpu/runtime/while_thunk.h"
 #include "xla/service/gpu/thunk.h"
 #include "xla/service/llvm_ir/buffer_assignment_util.h"
 #include "xla/service/llvm_ir/fused_ir_emitter.h"
@@ -172,14 +172,14 @@ limitations under the License.
 #include "tsl/protobuf/dnn.pb.h"
 
 #if GOOGLE_CUDA || TF_HIPBLASLT
-#include "xla/service/gpu/runtime3/gpublas_lt_matmul_thunk.h"
+#include "xla/service/gpu/runtime/gpublas_lt_matmul_thunk.h"
 #endif  // GOOGLE_CUDA || TF_HIPBLASLT
 
 #if GOOGLE_CUDA || TENSORFLOW_USE_ROCM
 #include "xla/service/gpu/ir_emitter_triton.h"
-#include "xla/service/gpu/runtime3/cholesky_thunk.h"
-#include "xla/service/gpu/runtime3/cub_sort_thunk.h"
-#include "xla/service/gpu/runtime3/triangular_solve_thunk.h"
+#include "xla/service/gpu/runtime/cholesky_thunk.h"
+#include "xla/service/gpu/runtime/cub_sort_thunk.h"
+#include "xla/service/gpu/runtime/triangular_solve_thunk.h"
 #endif  // GOOGLE_CUDA || TENSORFLOW_USE_ROCM
 
 namespace xla {

third_party/xla/xla/service/gpu/ir_emitter_unnested.h

@@ -40,7 +40,7 @@ limitations under the License.
 #include "xla/service/gpu/hlo_fusion_analysis.h"
 #include "xla/service/gpu/ir_emitter.h"
 #include "xla/service/gpu/nccl_collective_thunk.h"
-#include "xla/service/gpu/runtime3/send_recv_thunk.h"
+#include "xla/service/gpu/runtime/send_recv_thunk.h"
 #include "xla/service/gpu/thunk.h"
 #include "xla/service/llvm_ir/ir_array.h"
 #include "xla/service/llvm_ir/llvm_util.h"

third_party/xla/xla/service/gpu/runtime3/annotation.cc → third_party/xla/xla/service/gpu/runtime/annotation.cc

@@ -13,7 +13,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#include "xla/service/gpu/runtime3/annotation.h"
+#include "xla/service/gpu/runtime/annotation.h"
 
 #include <cstddef>
 #include <cstring>

third_party/xla/xla/service/gpu/runtime3/annotation.h → third_party/xla/xla/service/gpu/runtime/annotation.h

@@ -13,8 +13,8 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#ifndef XLA_SERVICE_GPU_RUNTIME3_ANNOTATION_H_
-#define XLA_SERVICE_GPU_RUNTIME3_ANNOTATION_H_
+#ifndef XLA_SERVICE_GPU_RUNTIME_ANNOTATION_H_
+#define XLA_SERVICE_GPU_RUNTIME_ANNOTATION_H_
 
 #include <string>
 #include <string_view>
@@ -90,4 +90,4 @@ const ModuleAnnotations* GetCurrentModuleAnnotations();
 
 }  // namespace xla::gpu
 
-#endif  // XLA_SERVICE_GPU_RUNTIME3_ANNOTATION_H_
+#endif  // XLA_SERVICE_GPU_RUNTIME_ANNOTATION_H_

third_party/xla/xla/service/gpu/runtime3/cholesky_thunk.cc → third_party/xla/xla/service/gpu/runtime/cholesky_thunk.cc

@@ -13,7 +13,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#include "xla/service/gpu/runtime3/cholesky_thunk.h"
+#include "xla/service/gpu/runtime/cholesky_thunk.h"
 
 #include <complex>
 #include <memory>

third_party/xla/xla/service/gpu/runtime3/cholesky_thunk.h → third_party/xla/xla/service/gpu/runtime/cholesky_thunk.h

@@ -13,8 +13,8 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#ifndef XLA_SERVICE_GPU_RUNTIME3_CHOLESKY_THUNK_H_
-#define XLA_SERVICE_GPU_RUNTIME3_CHOLESKY_THUNK_H_
+#ifndef XLA_SERVICE_GPU_RUNTIME_CHOLESKY_THUNK_H_
+#define XLA_SERVICE_GPU_RUNTIME_CHOLESKY_THUNK_H_
 
 #include <optional>
 
@@ -80,4 +80,4 @@ absl::Status RunCholesky(const se::GpuAsmOpts& asm_opts, PrimitiveType type,
 }  // namespace gpu
 }  // namespace xla
 
-#endif  // XLA_SERVICE_GPU_RUNTIME3_CHOLESKY_THUNK_H_
+#endif  // XLA_SERVICE_GPU_RUNTIME_CHOLESKY_THUNK_H_

third_party/xla/xla/service/gpu/runtime3/command_buffer_allocations.cc → third_party/xla/xla/service/gpu/runtime/command_buffer_allocations.cc

@@ -13,7 +13,7 @@ See the License for the specific language governing permissions and
 limitations under the License.
 ==============================================================================*/
 
-#include "xla/service/gpu/runtime3/command_buffer_allocations.h"
+#include "xla/service/gpu/runtime/command_buffer_allocations.h"
 
 #include <utility>