Skip to content

Commit 5cf7d2d

Browse files
atalmanfacebook-github-bot
atalman
authored and
facebook-github-bot
committed
Revert "Enable ROCm RNN-T Loss (#2485)" (#3586)
Summary: This reverts commit c593961. Unblock 2.1.0 rc Pull Request resolved: #3586 Reviewed By: osalpekar Differential Revision: D48842032 Pulled By: atalman fbshipit-source-id: bbdf9e45c9aa5fde00f315a2ff491ed050bc1707
1 parent bf77b2a commit 5cf7d2d

File tree

15 files changed

+5
-157
lines changed

15 files changed

+5
-157
lines changed

‎.gitmodules‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +0,0 @@
1-
[submodule "third_party/hipify_torch"]
2-
path = third_party/hipify_torch
3-
url = https://github.com/ROCmSoftwarePlatform/hipify_torch

‎CMakeLists.txt‎

Lines changed: 0 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -76,11 +76,6 @@ if(USE_ROCM)
7676
if(NOT PYTORCH_FOUND_HIP)
7777
set(USE_ROCM OFF)
7878
endif()
79-
80-
if(CMAKE_VERSION VERSION_LESS 3.21.0)
81-
message("Need at least CMake 3.21.0 to compile ROCm support.")
82-
set(USE_ROCM OFF)
83-
endif()
8479
endif()
8580

8681
if(USE_CUDA)
@@ -95,11 +90,6 @@ if(USE_CUDA)
9590
)
9691
endif()
9792

98-
if(USE_ROCM)
99-
enable_language(HIP)
100-
endif()
101-
102-
find_package(Torch REQUIRED)
10393
include(cmake/TorchAudioHelper.cmake)
10494

10595
# https://github.com/pytorch/pytorch/issues/54174

‎third_party/hipify_torch‎

Lines changed: 0 additions & 1 deletion
This file was deleted.

‎torchaudio/csrc/CMakeLists.txt‎

Lines changed: 1 addition & 54 deletions
Original file line numberDiff line numberDiff line change
@@ -1,29 +1,6 @@
11
################################################################################
22
# libtorchaudio
33
################################################################################
4-
5-
if(USE_ROCM)
6-
list (APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
7-
FIND_PACKAGE(HIP REQUIRED)
8-
MESSAGE(STATUS "hip found ${ROCM_FOUND}")
9-
10-
list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/third_party/hipify_torch/cmake")
11-
include(Hipify)
12-
13-
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
14-
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
15-
set(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE)
16-
list( APPEND CMAKE_INSTALL_RPATH "/opt/rocm/llvm/lib" )
17-
set(OPENMP_LIBRARIES "/opt/rocm/llvm/lib/")
18-
set(OpenMP_CXX "${CMAKE_CXX_COMPILER}")
19-
set(OpenMP_CXX_FLAGS "-fopenmp=libomp")
20-
#set(OpenMP_CXX_LIB_NAMES "omp")
21-
set(OpenMP_omp_LIBRARY omp)
22-
find_package(OpenMP REQUIRED)
23-
24-
endif()
25-
26-
274
set(
285
sources
296
lfilter.cpp
@@ -62,37 +39,6 @@ if(BUILD_RNNT)
6239
rnnt/gpu/compute.cu
6340
)
6441
endif()
65-
66-
if (USE_ROCM)
67-
hipify(CUDA_SOURCE_DIR ${PROJECT_SOURCE_DIR}/torchaudio/csrc/rnnt/gpu HIP_SOURCE_DIR ${PROJECT_SOURCE_DIR}/torchaudio/csrc/rnnt/hip)
68-
if ( NOT HIP_ADD_LIBRARY_FOUND )
69-
list(APPEND CMAKE_MODULE_PATH /opt/rocm/hip/cmake)
70-
find_package(HIP REQUIRED)
71-
endif()
72-
73-
list(
74-
APPEND
75-
sources
76-
rnnt/hip/compute_alphas.hip
77-
rnnt/hip/compute_betas.hip
78-
rnnt/hip/compute.hip
79-
)
80-
endif()
81-
endif()
82-
83-
if(USE_ROCM)
84-
list(
85-
APPEND
86-
additional_libs
87-
hip::host
88-
hip::device
89-
/opt/rocm/llvm/lib/libomp.so
90-
)
91-
list(
92-
APPEND
93-
compile_definitions
94-
USE_ROCM
95-
)
9642
endif()
9743

9844
if(BUILD_RIR)
@@ -141,6 +87,7 @@ endif()
14187
#------------------------------------------------------------------------------#
14288
# END OF CUSTOMIZATION LOGICS
14389
#------------------------------------------------------------------------------#
90+
14491
torchaudio_library(
14592
libtorchaudio
14693
"${sources}"

‎torchaudio/csrc/rnnt/gpu/compute.cu‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,6 @@
11
#include <c10/cuda/CUDAStream.h>
22
#include <torch/types.h>
3-
#ifdef __HIP_PLATFORM_AMD__
4-
#include <torchaudio/csrc/rnnt/hip/gpu_transducer_hip.h>
5-
#else
63
#include <torchaudio/csrc/rnnt/gpu/gpu_transducer.h>
7-
#endif
84

95
namespace torchaudio {
106
namespace rnnt {

‎torchaudio/csrc/rnnt/gpu/compute_alphas.cu‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,6 @@
11
#include <c10/cuda/CUDAStream.h>
22
#include <torch/types.h>
3-
#ifdef __HIP_PLATFORM_AMD__
4-
#include <torchaudio/csrc/rnnt/hip/gpu_transducer_hip.h>
5-
#else
63
#include <torchaudio/csrc/rnnt/gpu/gpu_transducer.h>
7-
#endif
84

95
namespace torchaudio {
106
namespace rnnt {

‎torchaudio/csrc/rnnt/gpu/compute_betas.cu‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,6 @@
11
#include <c10/cuda/CUDAStream.h>
22
#include <torch/types.h>
3-
#ifdef __HIP_PLATFORM_AMD__
4-
#include <torchaudio/csrc/rnnt/hip/gpu_transducer_hip.h>
5-
#else
63
#include <torchaudio/csrc/rnnt/gpu/gpu_transducer.h>
7-
#endif
84

95
namespace torchaudio {
106
namespace rnnt {

‎torchaudio/csrc/rnnt/gpu/gpu_kernel_utils.cuh‎

Lines changed: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -2,11 +2,7 @@
22

33
#ifdef USE_CUDA
44

5-
#ifdef __HIP_PLATFORM_AMD__
6-
#include <torchaudio/csrc/rnnt/hip/math_hip.cuh>
7-
#else
85
#include <torchaudio/csrc/rnnt/gpu/math.cuh>
9-
#endif
106

117
namespace torchaudio {
128
namespace rnnt {
@@ -43,11 +39,7 @@ __global__ void ReduceMax2D(
4339

4440
CAST_DTYPE shf;
4541
for (int stride = (WARP_SIZE >> 1); stride > 0; stride >>= 1) {
46-
#ifdef __HIP_PLATFORM_AMD__
47-
shf = __shfl_down(val, stride);
48-
#else
4942
shf = __shfl_down_sync(0xFFFFFFFF, val, stride);
50-
#endif
5143
if (threadIdx.x < stride && threadIdx.x + stride < dim) {
5244
if (shf > val) {
5345
val = shf;
@@ -89,11 +81,7 @@ __global__ void ReduceLogSumExpGivenMax2D(
8981

9082
CAST_DTYPE shf;
9183
for (int stride = (WARP_SIZE >> 1); stride > 0; stride >>= 1) {
92-
#ifdef __HIP_PLATFORM_AMD__
93-
shf = __shfl_down(val, stride);
94-
#else
9584
shf = __shfl_down_sync(0xFFFFFFFF, val, stride);
96-
#endif
9785
if (threadIdx.x < stride && threadIdx.x + stride < dim) {
9886
val = val + shf;
9987
}

‎torchaudio/csrc/rnnt/gpu/gpu_kernels.cuh‎

Lines changed: 0 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -4,15 +4,9 @@
44

55
#include <cassert>
66

7-
#ifdef __HIP_PLATFORM_AMD__
8-
#include <torchaudio/csrc/rnnt/hip/kernel_utils.h>
9-
#include <torchaudio/csrc/rnnt/hip/kernels.h>
10-
#include <torchaudio/csrc/rnnt/hip/math_hip.cuh>
11-
#else
127
#include <torchaudio/csrc/rnnt/gpu/kernel_utils.h>
138
#include <torchaudio/csrc/rnnt/gpu/kernels.h>
149
#include <torchaudio/csrc/rnnt/gpu/math.cuh>
15-
#endif
1610

1711
namespace torchaudio {
1812
namespace rnnt {
@@ -132,11 +126,7 @@ __device__ void ComputeAlphas(
132126

133127
#pragma unroll
134128
for (int i = 1; i < warpSize; i <<= 1) {
135-
#ifdef __HIP_PLATFORM_AMD__
136-
val = __shfl_up(skip_prob, i);
137-
#else
138129
val = __shfl_up_sync(0xffffffff, skip_prob, i);
139-
#endif
140130
if (i <= threadIdx.x) {
141131
skip_prob = skip_prob + val;
142132
}
@@ -160,11 +150,7 @@ __device__ void ComputeAlphas(
160150
CAST_DTYPE out = val;
161151

162152
for (int i = 1; i < warpSize; ++i) {
163-
#ifdef __HIP_PLATFORM_AMD__
164-
val = __shfl_up(val, 1);
165-
#else
166153
val = __shfl_up_sync(0xffffffff, val, 1);
167-
#endif
168154
if (i == threadIdx.x) {
169155
val = math::lse(val + skip_prob, emit);
170156
out = val;
@@ -239,11 +225,7 @@ __device__ void ComputeBetasCosts(
239225

240226
#pragma unroll
241227
for (int i = 1; i < warpSize; i <<= 1) {
242-
#ifdef __HIP_PLATFORM_AMD__
243-
val = __shfl_up(skip_prob, i);
244-
#else
245228
val = __shfl_up_sync(0xffffffff, skip_prob, i);
246-
#endif
247229
if (i <= threadIdx.x) {
248230
skip_prob = skip_prob + val;
249231
}
@@ -266,11 +248,7 @@ __device__ void ComputeBetasCosts(
266248
CAST_DTYPE out = val;
267249

268250
for (int i = 1; i < warpSize; ++i) {
269-
#ifdef __HIP_PLATFORM_AMD__
270-
val = __shfl_up(val, 1);
271-
#else
272251
val = __shfl_up_sync(0xffffffff, val, 1);
273-
#endif
274252
if (i == threadIdx.x) {
275253
val = math::lse(val + skip_prob, emit);
276254
out = val;

‎torchaudio/csrc/rnnt/gpu/gpu_transducer.h‎

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -3,13 +3,8 @@
33
#ifdef USE_CUDA
44

55
#include <torchaudio/csrc/rnnt/workspace.h>
6-
#ifdef __HIP_PLATFORM_AMD__
7-
#include <torchaudio/csrc/rnnt/hip/gpu_kernel_utils_hip.cuh>
8-
#include <torchaudio/csrc/rnnt/hip/gpu_kernels_hip.cuh>
9-
#else
106
#include <torchaudio/csrc/rnnt/gpu/gpu_kernel_utils.cuh>
117
#include <torchaudio/csrc/rnnt/gpu/gpu_kernels.cuh>
12-
#endif
138

149
namespace torchaudio {
1510
namespace rnnt {

0 commit comments

Comments
 (0)