From d0115737b71831d2e694d84c3b56a3fb0c8fd890 Mon Sep 17 00:00:00 2001 From: Jason Furmanek Date: Wed, 6 Dec 2023 18:29:12 +0000 Subject: [PATCH] Initial commit to resolve merge conflicts --- third_party/xla/xla/service/gpu/BUILD | 4 -- .../xla/service/gpu/buffer_comparator_test.cc | 6 --- .../xla/stream_executor/device_description.h | 42 ------------------- .../xla/stream_executor/rocm/hip_blas_lt.cc | 19 --------- .../xla/stream_executor/rocm/hip_blas_lt.h | 5 --- third_party/xla/xla/tests/BUILD | 4 -- 6 files changed, 80 deletions(-) diff --git a/third_party/xla/xla/service/gpu/BUILD b/third_party/xla/xla/service/gpu/BUILD index cb5e399b6e5937..a010fba0a0c1cc 100644 --- a/third_party/xla/xla/service/gpu/BUILD +++ b/third_party/xla/xla/service/gpu/BUILD @@ -5035,14 +5035,10 @@ xla_cc_test( xla_cc_test( name = "determinism_test", srcs = ["determinism_test.cc"], -<<<<<<< HEAD - tags = tf_cuda_tests_tags() + ["no_rocm"], -======= local_defines = if_cuda_is_configured(["GOOGLE_CUDA=1"]) + if_rocm_is_configured([ "TENSORFLOW_USE_ROCM=1", ]), tags = tf_gpu_tests_tags(), ->>>>>>> db579439eef970657f5ddbf05dc9b798cb748c51 deps = [ ":autotuner_util", "//xla:literal", diff --git a/third_party/xla/xla/service/gpu/buffer_comparator_test.cc b/third_party/xla/xla/service/gpu/buffer_comparator_test.cc index fbd2163c587d0f..77d1b5d783ba44 100644 --- a/third_party/xla/xla/service/gpu/buffer_comparator_test.cc +++ b/third_party/xla/xla/service/gpu/buffer_comparator_test.cc @@ -39,11 +39,6 @@ namespace { class BufferComparatorTest : public testing::Test { protected: BufferComparatorTest() -<<<<<<< HEAD - : platform_( - se::MultiPlatformManager::PlatformWithName(PLATFORM).ValueOrDie()), - stream_exec_(platform_->ExecutorForDevice(0).value()) {} -======= #if GOOGLE_CUDA : platform_(se::MultiPlatformManager::PlatformWithName("CUDA").value()), #elif TENSORFLOW_USE_ROCM @@ -51,7 +46,6 @@ class BufferComparatorTest : public testing::Test { #endif stream_exec_(platform_->ExecutorForDevice(0).value()) { } ->>>>>>> db579439eef970657f5ddbf05dc9b798cb748c51 // Take floats only for convenience. Still uses ElementType internally. template diff --git a/third_party/xla/xla/stream_executor/device_description.h b/third_party/xla/xla/stream_executor/device_description.h index 2cba18f555360b..85cac5a9ee075d 100644 --- a/third_party/xla/xla/stream_executor/device_description.h +++ b/third_party/xla/xla/stream_executor/device_description.h @@ -159,25 +159,15 @@ class RocmComputeCapability { return absl::StrJoin(kSupportedGfxVersions, ", "); } -<<<<<<< HEAD - bool has_nhwc_layout_support() const { -======= bool gfx9_mi100_or_later() const { ->>>>>>> db579439eef970657f5ddbf05dc9b798cb748c51 static constexpr absl::string_view kList[] = {"gfx908", "gfx90a", "gfx940", "gfx941", "gfx942"}; return absl::c_count(kList, gfx_version()) != 0; } -<<<<<<< HEAD - bool has_bf16_dtype_support() const { - static constexpr absl::string_view kList[] = {"gfx908", "gfx90a", "gfx940", - "gfx941", "gfx942"}; -======= bool gfx9_mi200_or_later() const { static constexpr absl::string_view kList[] = {"gfx90a", "gfx940", "gfx941", "gfx942"}; ->>>>>>> db579439eef970657f5ddbf05dc9b798cb748c51 return absl::c_count(kList, gfx_version()) != 0; } @@ -190,25 +180,6 @@ class RocmComputeCapability { bool has_bf16_dtype_support() const { return gfx9_mi100_or_later(); } bool has_fast_fp16_support() const { -<<<<<<< HEAD - static constexpr absl::string_view kList[] = {"gfx906", "gfx908", "gfx90a", - "gfx940", "gfx941", "gfx942", - "gfx1030", "gfx1100"}; - return absl::c_count(kList, gfx_version()) != 0; - } - - bool has_mfma_instr_support() const { - static constexpr absl::string_view kList[] = {"gfx908", "gfx90a", "gfx940", - "gfx941", "gfx942"}; - return absl::c_count(kList, gfx_version()) != 0; - } - - bool has_fp16_atomics_support() const { - // TODO(rocm): Check. This should be the same as has_fast_fp16_support(). - static constexpr absl::string_view kList[] = {"gfx90a", "gfx940", "gfx941", - "gfx942"}; - return absl::c_count(kList, gfx_version()) != 0; -======= return gfx9_mi100_or_later() || navi21() || navi31(); } @@ -217,7 +188,6 @@ class RocmComputeCapability { bool has_fp16_atomics_support() const { // TODO(rocm): Check. This should be the same as has_fast_fp16_support(). return gfx9_mi200_or_later(); ->>>>>>> db579439eef970657f5ddbf05dc9b798cb748c51 } bool fence_before_barrier() const { @@ -240,17 +210,6 @@ class RocmComputeCapability { std::string gcn_arch_name_ = "gfx000"; // default to invalid arch. static constexpr absl::string_view kSupportedGfxVersions[]{ -<<<<<<< HEAD - "gfx900", // MI25 - "gfx906", // MI50 / MI60 - "gfx908", // MI100 - "gfx90a", // MI200 - "gfx940", // MI300 - "gfx941", // MI300 - "gfx942", // MI300 - "gfx1030", // Navi21 - "gfx1100" // Navi31 -======= "gfx900", // MI25 "gfx906", // MI50 / MI60 "gfx908", // MI100 @@ -258,7 +217,6 @@ class RocmComputeCapability { "gfx940", "gfx941", "gfx942", "gfx1030", // Navi21 "gfx1100" // Navi31 ->>>>>>> db579439eef970657f5ddbf05dc9b798cb748c51 }; }; diff --git a/third_party/xla/xla/stream_executor/rocm/hip_blas_lt.cc b/third_party/xla/xla/stream_executor/rocm/hip_blas_lt.cc index 925c4082caa572..262a3a5c3122f9 100644 --- a/third_party/xla/xla/stream_executor/rocm/hip_blas_lt.cc +++ b/third_party/xla/xla/stream_executor/rocm/hip_blas_lt.cc @@ -478,24 +478,6 @@ tsl::Status BlasLt::MatmulPlan::ExecuteOnStream( } // Other data types: -<<<<<<< HEAD - TYPED_MATMUL(float, HIP_R_16BF, HIP_R_16BF, HIP_R_16BF, - HIP_R_16BF) - TYPED_MATMUL(float, HIP_R_16F, HIP_R_16F, HIP_R_16F, - HIP_R_16F) - TYPED_MATMUL(float, HIP_R_16BF, HIP_R_16BF, HIP_R_32F, - HIP_R_32F) - TYPED_MATMUL(float, HIP_R_16F, HIP_R_16F, HIP_R_32F, - HIP_R_32F) - TYPED_MATMUL(float, HIP_R_32F, HIP_R_32F, HIP_R_32F, - HIP_R_32F) - TYPED_MATMUL(double, HIP_R_64F, HIP_R_64F, HIP_R_64F, - HIP_R_64F) - TYPED_MATMUL(complex64, HIP_C_32F, HIP_C_32F, HIP_C_32F, - HIP_C_32F) - TYPED_MATMUL(complex128, HIP_C_64F, HIP_C_64F, HIP_C_64F, - HIP_C_64F) -======= TYPED_MATMUL(float, HIP_R_16BF, HIP_R_16BF, HIP_R_16BF, HIP_R_16BF) TYPED_MATMUL(float, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIP_R_16F) TYPED_MATMUL(float, HIP_R_16BF, HIP_R_16BF, HIP_R_32F, HIP_R_32F) @@ -504,7 +486,6 @@ tsl::Status BlasLt::MatmulPlan::ExecuteOnStream( TYPED_MATMUL(double, HIP_R_64F, HIP_R_64F, HIP_R_64F, HIP_R_64F) TYPED_MATMUL(complex64, HIP_C_32F, HIP_C_32F, HIP_C_32F, HIP_C_32F) TYPED_MATMUL(complex128, HIP_C_64F, HIP_C_64F, HIP_C_64F, HIP_C_64F) ->>>>>>> db579439eef970657f5ddbf05dc9b798cb748c51 #undef TYPED_MATMUL diff --git a/third_party/xla/xla/stream_executor/rocm/hip_blas_lt.h b/third_party/xla/xla/stream_executor/rocm/hip_blas_lt.h index 6f9530720f279f..0e253c7d8062e2 100644 --- a/third_party/xla/xla/stream_executor/rocm/hip_blas_lt.h +++ b/third_party/xla/xla/stream_executor/rocm/hip_blas_lt.h @@ -72,12 +72,7 @@ class BlasLt : public gpu::BlasLt { hipblasLtMatmulDesc_t get() const { return handle_.get(); } private: -<<<<<<< HEAD - MatmulDesc(hipblasLtMatmulDesc_t handle, - hipblasComputeType_t compute_type, -======= MatmulDesc(hipblasLtMatmulDesc_t handle, hipblasComputeType_t compute_type, ->>>>>>> db579439eef970657f5ddbf05dc9b798cb748c51 hipDataType datatype) : handle_(handle, wrap::hipblasLtMatmulDescDestroy), compute_type_(compute_type), diff --git a/third_party/xla/xla/tests/BUILD b/third_party/xla/xla/tests/BUILD index 51aa59abd76466..6ee3b50a90fdea 100644 --- a/third_party/xla/xla/tests/BUILD +++ b/third_party/xla/xla/tests/BUILD @@ -777,13 +777,9 @@ xla_test( xla_test( name = "array_elementwise_ops_test", srcs = ["array_elementwise_ops_test.cc"], -<<<<<<< HEAD - tags = ["no_rocm"], -======= local_defines = if_cuda_is_configured(["GOOGLE_CUDA=1"]) + if_rocm_is_configured([ "TENSORFLOW_USE_ROCM=1", ]), ->>>>>>> db579439eef970657f5ddbf05dc9b798cb748c51 shard_count = 25, deps = [ ":client_library_test_base",