Merge pull request #43633 from ROCmSoftwarePlatform/google_upstream_r21_rocm_updates_200923

[ROCm] Porting changes to support newer ROCm versions to r2.1

tensorflow/compiler/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc

@@ -542,10 +542,9 @@ static std::vector<string> GetROCDLPaths(int amdgpu_version,
                                          const string& rocdl_dir_path) {
   // AMDGPU version-neutral bitcodes.
   static std::vector<string>* rocdl_filenames = new std::vector<string>(
-      {"hc.amdgcn.bc", "opencl.amdgcn.bc", "ocml.amdgcn.bc", "ockl.amdgcn.bc",
+      {"hc.bc", "opencl.bc", "ocml.bc", "ockl.bc", "oclc_finite_only_off.bc",
-       "oclc_finite_only_off.amdgcn.bc", "oclc_daz_opt_off.amdgcn.bc",
+       "oclc_daz_opt_off.bc", "oclc_correctly_rounded_sqrt_on.bc",
-       "oclc_correctly_rounded_sqrt_on.amdgcn.bc",
+       "oclc_unsafe_math_off.bc"});
-       "oclc_unsafe_math_off.amdgcn.bc"});
 
   // Construct full path to ROCDL bitcode libraries.
   std::vector<string> result;
@@ -556,7 +555,7 @@ static std::vector<string> GetROCDLPaths(int amdgpu_version,
   // Add AMDGPU version-specific bitcodes.
   result.push_back(tensorflow::io::JoinPath(
       rocdl_dir_path,
-      absl::StrCat("oclc_isa_version_", amdgpu_version, ".amdgcn.bc")));
+      absl::StrCat("oclc_isa_version_", amdgpu_version, ".bc")));
   return result;
 }
 

tensorflow/core/kernels/reduction_gpu_kernels.cu.h

@@ -388,7 +388,7 @@ __global__ __launch_bounds__(1024) void ColumnReduceKernel(
     //  -         =
     //            =
     const int numRowsThisBlock =
-        min(blockDim.y, num_rows - blockIdx.y * blockDim.y);
+        min(int(blockDim.y), num_rows - blockIdx.y * blockDim.y);
 
     for (int row = 1; row < numRowsThisBlock; ++row) {
       value_type t = partial_sums[threadIdx.x * (TF_RED_WARPSIZE + 1) + row];

tensorflow/core/platform/default/rocm_rocdl_path.cc

@@ -37,7 +37,7 @@ string RocmRoot() {
 
 string RocdlRoot() {
 #if TENSORFLOW_COMPILER_IS_HIP_CLANG
-  return tensorflow::io::JoinPath(tensorflow::RocmRoot(), "lib");
+  return tensorflow::io::JoinPath(tensorflow::RocmRoot(), "amdgcn/bitcode");
 #else
   return tensorflow::io::JoinPath(tensorflow::RocmRoot(), "hcc/lib");
 #endif

tensorflow/core/platform/rocm_rocdl_path_test.cc

@@ -27,7 +27,7 @@ TEST(RocmRocdlPathTest, ROCDLPath) {
   VLOG(2) << "ROCm-Deivce-Libs root = " << RocdlRoot();
   std::vector<string> rocdl_files;
   TF_EXPECT_OK(Env::Default()->GetMatchingPaths(
-      io::JoinPath(RocdlRoot(), "*.amdgcn.bc"), &rocdl_files));
+      io::JoinPath(RocdlRoot(), "*.bc"), &rocdl_files));
   EXPECT_LT(0, rocdl_files.size());
 }
 #endif

tensorflow/tools/ci_build/Dockerfile.rocm

@@ -3,10 +3,10 @@
 FROM ubuntu:bionic
 MAINTAINER Jeff Poznanovic <jeffrey.poznanovic@amd.com>
 
-ARG ROCM_DEB_REPO=http://repo.radeon.com/rocm/apt/3.7/
+ARG ROCM_DEB_REPO=http://repo.radeon.com/rocm/apt/3.9/
 ARG ROCM_BUILD_NAME=xenial
 ARG ROCM_BUILD_NUM=main
-ARG ROCM_PATH=/opt/rocm-3.7.0
+ARG ROCM_PATH=/opt/rocm-3.9.0
 
 ENV DEBIAN_FRONTEND noninteractive
 ENV TF_NEED_ROCM 1

tensorflow/tools/ci_build/linux/rocm/run_cc_core.sh

@@ -28,7 +28,7 @@ echo "Bazel will use ${N_BUILD_JOBS} concurrent build job(s) and ${N_TEST_JOBS}
 echo ""
 
 # First positional argument (if any) specifies the ROCM_INSTALL_DIR
-ROCM_INSTALL_DIR=/opt/rocm-3.7.0
+ROCM_INSTALL_DIR=/opt/rocm-3.9.0
 if [[ -n $1 ]]; then
     ROCM_INSTALL_DIR=$1
 fi

tensorflow/tools/ci_build/linux/rocm/run_csb_tests.sh

@@ -28,7 +28,7 @@ echo "Bazel will use ${N_BUILD_JOBS} concurrent build job(s) and ${N_TEST_JOBS}
 echo ""
 
 # First positional argument (if any) specifies the ROCM_INSTALL_DIR
-ROCM_INSTALL_DIR=/opt/rocm-3.7.0
+ROCM_INSTALL_DIR=/opt/rocm-3.9.0
 if [[ -n $1 ]]; then
     ROCM_INSTALL_DIR=$1
 fi

tensorflow/tools/ci_build/linux/rocm/run_py3_core.sh

@@ -28,7 +28,7 @@ echo "Bazel will use ${N_BUILD_JOBS} concurrent build job(s) and ${N_TEST_JOBS}
 echo ""
 
 # First positional argument (if any) specifies the ROCM_INSTALL_DIR
-ROCM_INSTALL_DIR=/opt/rocm-3.7.0
+ROCM_INSTALL_DIR=/opt/rocm-3.9.0
 if [[ -n $1 ]]; then
     ROCM_INSTALL_DIR=$1
 fi

tensorflow/tools/ci_build/xla/linux/rocm/run_py3.sh

@@ -30,7 +30,7 @@ export PYTHON_BIN_PATH=`which python3`
 export CC_OPT_FLAGS='-mavx'
 
 export TF_NEED_ROCM=1
-export ROCM_PATH=/opt/rocm-3.3.0
+export ROCM_PATH=/opt/rocm-3.9.0
 export TF_GPU_COUNT=${N_GPUS}
 
 yes "" | $PYTHON_BIN_PATH configure.py

third_party/gpus/crosstool/clang/bin/crosstool_wrapper_driver_rocm.tpl

@@ -186,6 +186,7 @@ def InvokeHipcc(argv, log=False):
   # of link time. This allows the default host compiler (gcc) be used as the
   # linker for TensorFlow on ROCm platform.
   hipccopts += ' -fno-gpu-rdc '
+  hipccopts += ' -fcuda-flush-denormals-to-zero '
   hipccopts += undefines
   hipccopts += defines
   hipccopts += std_options

third_party/gpus/rocm_configure.bzl

@@ -204,6 +204,7 @@ def _rocm_include_path(repository_ctx, rocm_config):
     inc_dirs.append(rocm_config.rocm_toolkit_path + "/llvm/lib/clang/9.0.0/include")
     inc_dirs.append(rocm_config.rocm_toolkit_path + "/llvm/lib/clang/10.0.0/include")
     inc_dirs.append(rocm_config.rocm_toolkit_path + "/llvm/lib/clang/11.0.0/include")
+    inc_dirs.append(rocm_config.rocm_toolkit_path + "/llvm/lib/clang/12.0.0/include")
 
     # Add rocrand and hiprand headers
     inc_dirs.append(rocm_config.rocm_toolkit_path + "/rocrand/include")