merge amd-staging into amd-mainline

for ROCm HIP 6.1.1

Change-Id: Ib2273b022b1ebb3b3ee1e347cb00780b6c2055a1

Author: emankov

CHANGELOG.md

@@ -3,6 +3,20 @@
 Documentation for HIPIFY is available at
 [https://rocmdocs.amd.com/projects/HIPIFY/en/latest/](https://rocmdocs.amd.com/projects/HIPIFY/en/latest/).
 
+## HIPIFY for ROCm 6.1.1
+
+### Additions
+
+* LLVM 18.1.2 support
+* cuDNN 9.0.0 support
+* New options:
+  * `--clang-resource-directory` to specify the clang resource path - the path to the parent folder for the `include` folder that
+    contains `__clang_cuda_runtime_wrapper.h` and other header files used during the hipification process
+
+### Fixes
+
+* Clang resource files used during hipification are now searchable and also can be specified by the `--clang-resource-directory` option
+
 ## HIPIFY for ROCm 6.1.0
 
 ### Additions
@@ -14,13 +28,10 @@ Documentation for HIPIFY is available at
 * Full `rocSPARSE` support
 * New options:
   * `--amap` to hipify as much as possible, ignoring `--default-preprocessor` behavior
-  * `--clang-resource-directory` to specify the clang resource path - the path to the parent folder for the `include` folder that
-    contains `__clang_cuda_runtime_wrapper.h` and other header files used during the hipification process
 
 ### Fixes
 
 * Code blocks skipped by the Preprocessor are not hipified anymore under the `--default-preprocessor` option
-* Clang resource files used during hipification are now searchable and also can be specified by the `--clang-resource-directory` option
 
 ## HIPIFY for ROCm 6.0.2
 

README.md

@@ -3,28 +3,6 @@
 HIPIFY is a set of tools that you can use to automatically translate CUDA source code into portable
 [HIP](https://github.com/ROCm/HIP) C++.
 
-## Table of contents
-
-<!-- toc -->
-
-* [hipify-clang](./docs/hipify-clang.md#hipify-clang)
-  * [Dependencies](./docs/hipify-clang.md#hipify-clang-dependencies)
-  * [Usage](./docs/hipify-clang.md#hipify-clang-usage)
-    * [JSON Compilation Database](./docs/hipify-clang.md#hipify-clang-using-json-compilation-database)
-    * [Hipification Statistics](./docs/hipify-clang.md#hipify-clang-hipification-statistics)
-  * [Building](./docs/hipify-clang.md#hipify-clang-building)
-  * [Testing](./docs/hipify-clang.md#hipify-clang-testing)
-    * [Linux](./docs/hipify-clang.md#hipify-clang-linux-testing)
-    * [Windows](./docs/hipify-clang.md#hipify-clang-windows-testing)
-* [hipify-perl](./docs/hipify-perl.md#hipify-perl)
-  * [Usage](./docs/hipify-perl.md#hipify-perl-usage)
-  * [Building](./docs/hipify-perl.md#hipify-perl-building)
-* Related: [hipify_torch](https://github.com/ROCmSoftwarePlatform/hipify_torch)
-* [Supported CUDA APIs](./docs/supported_apis.md#supported-cuda-apis)
-* [Documentation](#documentation)
-
-<!-- tocstop -->
-
 ## Documentation
 
 Documentation for HIPIFY is available at

bin/hipify-perl

@@ -2,7 +2,7 @@
 
 hipify-clang and hipify-perl are tools to translate CUDA source code into portable HIP C++ automatically.
 
-* [hipify-clang](hipify-clang.md)
+* [hipify-clang](hipify-clang.rst)
 * [hipify-perl](hipify-perl.md)
 * [Supported CUDA APIs](supported_apis.md)
 

docs/pics/csv_statistics.png renamed to docs/data/csv_statistics.png

@@ -1 +1 @@
-rocm-docs-core==0.35.0
+rocm-docs-core==0.38.0

patches/patch_for_clang_10.0.0_bug_47332.zip renamed to docs/data/patches/patch_for_clang_10.0.0_bug_47332.zip

@@ -84,9 +84,7 @@ pygments==2.15.0
     #   pydata-sphinx-theme
     #   sphinx
 pyjwt[crypto]==2.6.0
-    # via
-    #   pygithub
-    #   pyjwt
+    # via pygithub
 pynacl==1.5.0
     # via pygithub
 pytz==2023.3.post1
@@ -100,7 +98,7 @@ requests==2.31.0
     # via
     #   pygithub
     #   sphinx
-rocm-docs-core==0.35.0
+rocm-docs-core==0.38.0
     # via -r requirements.in
 smmap==5.0.0
     # via gitdb

patches/patch_for_clang_10.0.1_bug_47332.zip renamed to docs/data/patches/patch_for_clang_10.0.1_bug_47332.zip

@@ -2276,7 +2276,7 @@ clang::SourceRange getWriteRange(clang::SourceManager &SM, const clang::SourceRa
   // If the range is contained within a macro, update the macro definition.
   // Otherwise, use the file location and hope for the best.
   if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) {
-    return {SM.getExpansionLoc(begin), SM.getExpansionLoc(end)};
+    return {SM.getFileLoc(begin), SM.getFileLoc(end)};
   }
   return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)};
 }

patches/patch_for_clang_11.0.0_bug_47332.zip renamed to docs/data/patches/patch_for_clang_11.0.0_bug_47332.zip

@@ -536,6 +536,7 @@ std::string Statistics::getCudaVersion(const cudaVersions& ver) {
     case CUDNN_895: return "8.9.5";
     case CUDNN_896: return "8.9.6";
     case CUDNN_897: return "8.9.7";
+    case CUDNN_900: return "9.0.0";
   }
   return "";
 }

patches/patch_for_clang_14.0.0_bug_54609.zip renamed to docs/data/patches/patch_for_clang_14.0.0_bug_54609.zip

@@ -290,7 +290,8 @@ enum cudaVersions {
   CUDNN_895 = 895,
   CUDNN_896 = 896,
   CUDNN_897 = 897,
-  CUDNN_LATEST = CUDNN_897,
+  CUDNN_900 = 900,
+  CUDNN_LATEST = CUDNN_900,
 };
 
 enum hipVersions {

patches/patch_for_clang_14.0.1_bug_54609.zip renamed to docs/data/patches/patch_for_clang_14.0.1_bug_54609.zip

@@ -137,6 +137,9 @@ if config.cuda_version_major >= 12:
     config.excludes.append('cusparse2rocsparse_9020_12000.cu')
     config.excludes.append('cusparse2rocsparse_before_12000.cu')
 
+if config.cudnn_version_major >= 9:
+    config.excludes.append('cudnn2miopen_before_9000.cu')
+
 if config.llvm_version_major < 8:
     config.excludes.append('cd_intro.cu')
 

patches/patch_for_clang_14.0.2_bug_54609.zip renamed to docs/data/patches/patch_for_clang_14.0.2_bug_54609.zip

@@ -24,6 +24,14 @@ if sys.platform in ['win32']:
 else:
     config.cuda_sdk_root = config.cuda_root
 
+config.cudnn_version_major = int("0")
+if config.cuda_dnn_root and config.cuda_dnn_root != "OFF":
+    cudnn_ver_file = config.cuda_dnn_root + '/include/cudnn_version.h';
+    if os.path.isfile(cudnn_ver_file):
+        with open(cudnn_ver_file) as f:
+            if 'CUDNN_MAJOR 9' in f.read():
+                config.cudnn_version_major = int("9")
+
 # Support substitution of the tools and libs dirs with user parameters. This is
 # used when we can't determine the tool dir at configuration time.
 try:

patches/patch_for_clang_14.0.3_bug_54609.zip renamed to docs/data/patches/patch_for_clang_14.0.3_bug_54609.zip

@@ -0,0 +1,28 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --hip-kernel-execution-syntax %clang_args
+// CHECK: #include <hip/hip_runtime.h>
+#include <stdio.h>
+
+#define a0     -3.0124472f
+#define a1      1.7383092f
+#define a2     -0.2796695f
+#define a3      0.0547837f
+#define a4     -0.0073118f
+
+#define CHECK(call)                                                            \
+{                                                                              \
+    const cudaError_t error = call;                                            \
+    if (error != cudaSuccess)                                                  \
+    {                                                                          \
+        fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__);                 \
+        fprintf(stderr, "code: %d, reason: %s\n", error,                       \
+                cudaGetErrorString(error));                                    \
+    }                                                                          \
+}
+
+__device__ __constant__ float coef[5];
+
+int main() {
+    const float h_coef[] = {a0, a1, a2, a3, a4};
+    // CHECK: CHECK( hipMemcpyToSymbol( HIP_SYMBOL(coef), h_coef, 5 * sizeof(float) ));
+    CHECK( cudaMemcpyToSymbol( coef, h_coef, 5 * sizeof(float) ));
+}

patches/patch_for_clang_14.0.4_bug_54609.zip renamed to docs/data/patches/patch_for_clang_14.0.4_bug_54609.zip

@@ -1,4 +1,4 @@
-// RUN: %run_test hipify "%s" "%t" %hipify_args 4 --skip-excluded-preprocessor-conditional-blocks --experimental --roc --miopen %clang_args -D__CUDA_API_VERSION_INTERNAL
+// RUN: %run_test hipify "%s" "%t" %hipify_args 5 --skip-excluded-preprocessor-conditional-blocks --experimental --roc --miopen --amap %clang_args -D__CUDA_API_VERSION_INTERNAL
 
 // CHECK: #include <hip/hip_runtime.h>
 #include <cuda_runtime.h>
@@ -647,24 +647,13 @@ int main() {
   int hiddenSize = 0;
   int layer = 0;
 
-  // TODO [#837]: Insert miopenRNNBiasMode_t* biasMode in the hipified miopenGetRNNDescriptor_V2 after miopenRNNMode_t* rnnMode: will need variable declaration
-  // CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnGetRNNDescriptor_v6(cudnnHandle_t handle, cudnnRNNDescriptor_t rnnDesc, int* hiddenSize, int* numLayers, cudnnDropoutDescriptor_t* dropoutDesc, cudnnRNNInputMode_t* inputMode, cudnnDirectionMode_t* direction, cudnnRNNMode_t* cellMode, cudnnRNNAlgo_t* algo, cudnnDataType_t* mathPrec);
-  // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenGetRNNDescriptor_V2(miopenRNNDescriptor_t rnnDesc, int* hiddenSize, int* layer, miopenDropoutDescriptor_t* dropoutDesc, miopenRNNInputMode_t* inputMode, miopenRNNDirectionMode_t* dirMode, miopenRNNMode_t* rnnMode, miopenRNNBiasMode_t* biasMode, miopenRNNAlgo_t* algoMode, miopenDataType_t* dataType);
-  // CHECK: status = miopenGetRNNDescriptor_V2(RNNDescriptor, &hiddenSize, &layer, &DropoutDescriptor, &RNNInputMode, &DirectionMode, &RNNMode, &RNNAlgo, &dataType);
-  status = cudnnGetRNNDescriptor_v6(handle, RNNDescriptor, &hiddenSize, &layer, &DropoutDescriptor, &RNNInputMode, &DirectionMode, &RNNMode, &RNNAlgo, &dataType);
-
   // NOTE: cudnnSetRNNDescriptor - removed after cuDNN 7.6.5
   // NOTE: cudnnSetRNNDescriptor_v5 - removed after cuDNN 7.6.5
   // TODO: add cudnnSetRNNDescriptor -> miopenSetRNNDescriptor_V2 mapping after implementing cuDNN versioning in tests
 
-  // TODO [#837]: Insert miopenRNNBiasMode_t biasMode in the hipified miopenSetRNNDescriptor_V2 after miopenRNNMode_t rnnMode: will need variable declaration
-  // CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnSetRNNDescriptor_v6(cudnnHandle_t handle, cudnnRNNDescriptor_t rnnDesc, const int hiddenSize, const int numLayers, cudnnDropoutDescriptor_t dropoutDesc, cudnnRNNInputMode_t inputMode, cudnnDirectionMode_t direction, cudnnRNNMode_t cellMode, cudnnRNNAlgo_t algo, cudnnDataType_t mathPrec);
-  // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenSetRNNDescriptor_V2(miopenRNNDescriptor_t rnnDesc, const int hsize, const int nlayers, miopenDropoutDescriptor_t dropoutDesc, miopenRNNInputMode_t inMode, miopenRNNDirectionMode_t direction, miopenRNNMode_t rnnMode, miopenRNNBiasMode_t biasMode, miopenRNNAlgo_t algo, miopenDataType_t dataType);
-  // CHECK: status = miopenSetRNNDescriptor_V2(RNNDescriptor, hiddenSize, layer, DropoutDescriptor, RNNInputMode, DirectionMode, RNNMode, RNNAlgo, dataType);
-  status = cudnnSetRNNDescriptor_v6(handle, RNNDescriptor, hiddenSize, layer, DropoutDescriptor, RNNInputMode, DirectionMode, RNNMode, RNNAlgo, dataType);
-
   int seqLength = 0;
 
+#if CUDNN_MAJOR < 9
   // CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnGetRNNWorkspaceSize(cudnnHandle_t handle, const cudnnRNNDescriptor_t rnnDesc, const int seqLength, const cudnnTensorDescriptor_t* xDesc, size_t* sizeInBytes);
   // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenGetRNNWorkspaceSize(miopenHandle_t handle, const miopenRNNDescriptor_t rnnDesc, const int sequenceLen, const miopenTensorDescriptor_t* xDesc, size_t* numBytes);
   // CHECK: status = miopenGetRNNWorkspaceSize(handle, RNNDescriptor, seqLength, &xD, &workSpaceSizeInBytes);
@@ -680,6 +669,11 @@ int main() {
   // CHECK: status = miopenGetRNNParamsSize(handle, RNNDescriptor, xD, &workSpaceSizeInBytes, dataType);
   status = cudnnGetRNNParamsSize(handle, RNNDescriptor, xD, &workSpaceSizeInBytes, dataType);
 
+  // CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnRNNForwardInference(cudnnHandle_t handle, const cudnnRNNDescriptor_t rnnDesc, const int seqLength, const cudnnTensorDescriptor_t* xDesc, const void* x, const cudnnTensorDescriptor_t hxDesc, const void* hx, const cudnnTensorDescriptor_t cxDesc, const void* cx, const cudnnFilterDescriptor_t wDesc, const void* w, const cudnnTensorDescriptor_t* yDesc, void* y, const cudnnTensorDescriptor_t hyDesc, void* hy, const cudnnTensorDescriptor_t cyDesc, void* cy, void* workSpace, size_t workSpaceSizeInBytes);
+  // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenRNNForwardInference(miopenHandle_t handle, miopenRNNDescriptor_t rnnDesc, const int sequenceLen, const miopenTensorDescriptor_t* xDesc, const void* x, const miopenTensorDescriptor_t hxDesc, const void* hx, const miopenTensorDescriptor_t cxDesc, const void* cx, const miopenTensorDescriptor_t wDesc, const void* w, const miopenTensorDescriptor_t* yDesc, void* y, const miopenTensorDescriptor_t hyDesc, void* hy, const miopenTensorDescriptor_t cyDesc, void* cy, void* workSpace, size_t workSpaceNumBytes);
+  // CHECK: status = miopenRNNForwardInference(handle, RNNDescriptor, seqLength, &xD, x, hxD, hx, cxD, cx, filterDescriptor, W, &yD, y, hyD, hy, cyD, cy, workSpace, workSpaceSizeInBytes);
+  status = cudnnRNNForwardInference(handle, RNNDescriptor, seqLength, &xD, x, hxD, hx, cxD, cx, filterDescriptor, W, &yD, y, hyD, hy, cyD, cy, workSpace, workSpaceSizeInBytes);
+
   // CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnRNNForwardTraining(cudnnHandle_t handle, const cudnnRNNDescriptor_t rnnDesc, const int seqLength, const cudnnTensorDescriptor_t* xDesc, const void* x, const cudnnTensorDescriptor_t hxDesc, const void* hx, const cudnnTensorDescriptor_t cxDesc, const void* cx, const cudnnFilterDescriptor_t wDesc, const void* w, const cudnnTensorDescriptor_t* yDesc, void* y, const cudnnTensorDescriptor_t hyDesc, void* hy, const cudnnTensorDescriptor_t cyDesc, void* cy, void* workSpace, size_t workSpaceSizeInBytes, void* reserveSpace, size_t reserveSpaceSizeInBytes);
   // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenRNNForwardTraining(miopenHandle_t handle, const miopenRNNDescriptor_t rnnDesc, const int sequenceLen, const miopenTensorDescriptor_t* xDesc, const void* x, const miopenTensorDescriptor_t hxDesc, const void* hx, const miopenTensorDescriptor_t cxDesc, const void* cx, const miopenTensorDescriptor_t wDesc, const void* w, const miopenTensorDescriptor_t* yDesc, void* y, const miopenTensorDescriptor_t hyDesc, void* hy, const miopenTensorDescriptor_t cyDesc, void* cy, void* workSpace, size_t workSpaceNumBytes, void* reserveSpace, size_t reserveSpaceNumBytes);
   // CHECK: status = miopenRNNForwardTraining(handle, RNNDescriptor, seqLength, &xD, x, hxD, hx, cxD, cx, filterDescriptor, W, &yD, y, hyD, hy, cyD, cy, workSpace, workSpaceSizeInBytes, reserveSpace, reserveSpaceNumBytes);
@@ -690,15 +684,18 @@ int main() {
   // CHECK: status = miopenRNNBackwardData(handle, RNNDescriptor, seqLength, &yD, y, &dyD, dy, dhyD, dhy, dcyD, dcy, filterDescriptor, W, hxD, hx, cxD, cx, &dxD, dx, dhxD, dhx, dcxD, dcx, workSpace, workSpaceSizeInBytes, &reserveSpace, reserveSpaceNumBytes);
   status = cudnnRNNBackwardData(handle, RNNDescriptor, seqLength, &yD, y, &dyD, dy, dhyD, dhy, dcyD, dcy, filterDescriptor, W, hxD, hx, cxD, cx, &dxD, dx, dhxD, dhx, dcxD, dcx, workSpace, workSpaceSizeInBytes, &reserveSpace, reserveSpaceNumBytes);
 
-  // CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnRNNBackwardWeights(cudnnHandle_t handle, const cudnnRNNDescriptor_t rnnDesc, const int seqLength, const cudnnTensorDescriptor_t* xDesc, const void* x, const cudnnTensorDescriptor_t hxDesc, const void* hx, const cudnnTensorDescriptor_t* yDesc, const void* y, const void* workSpace, size_t workSpaceSizeInBytes, const cudnnFilterDescriptor_t dwDesc, void* dw, const void* reserveSpace, size_t reserveSpaceSizeInBytes);
-  // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenRNNBackwardWeights(miopenHandle_t handle, const miopenRNNDescriptor_t rnnDesc, const int sequenceLen, const miopenTensorDescriptor_t* xDesc, const void* x, const miopenTensorDescriptor_t hxDesc, const void* hx, const miopenTensorDescriptor_t* yDesc, const void* y, const miopenTensorDescriptor_t dwDesc, void* dw, void* workSpace, size_t workSpaceNumBytes, const void* reserveSpace, size_t reserveSpaceNumBytes);
-  // CHECK: status = miopenRNNBackwardWeights(handle, RNNDescriptor, seqLength, &xD, x, hxD, hx, &yD, y, filterDescriptor, dw, workSpace, workSpaceSizeInBytes, &reserveSpace, reserveSpaceNumBytes);
-  status = cudnnRNNBackwardWeights(handle, RNNDescriptor, seqLength, &xD, x, hxD, hx, &yD, y, workSpace, workSpaceSizeInBytes, filterDescriptor, dw, &reserveSpace, reserveSpaceNumBytes);
+  // TODO [#837]: Insert int* blank_label_id, bool* apply_softmax_layer in the hipified miopenGetCTCLossDescriptor: will need variable declaration
+  // CUDA: cudnnStatus_t CUDNNWINAPI cudnnGetCTCLossDescriptor(cudnnCTCLossDescriptor_t ctcLossDesc, cudnnDataType_t* compType);
+  // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenGetCTCLossDescriptor(miopenCTCLossDescriptor_t ctcLossDesc, miopenDataType_t* dataType, int* blank_label_id, bool* apply_softmax_layer);
+  // CHECK: status = miopenGetCTCLossDescriptor(CTCLossDescriptor, &dataType);
+  status = cudnnGetCTCLossDescriptor(CTCLossDescriptor, &dataType);
 
-  // CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnRNNForwardInference(cudnnHandle_t handle, const cudnnRNNDescriptor_t rnnDesc, const int seqLength, const cudnnTensorDescriptor_t* xDesc, const void* x, const cudnnTensorDescriptor_t hxDesc, const void* hx, const cudnnTensorDescriptor_t cxDesc, const void* cx, const cudnnFilterDescriptor_t wDesc, const void* w, const cudnnTensorDescriptor_t* yDesc, void* y, const cudnnTensorDescriptor_t hyDesc, void* hy, const cudnnTensorDescriptor_t cyDesc, void* cy, void* workSpace, size_t workSpaceSizeInBytes);
-  // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenRNNForwardInference(miopenHandle_t handle, miopenRNNDescriptor_t rnnDesc, const int sequenceLen, const miopenTensorDescriptor_t* xDesc, const void* x, const miopenTensorDescriptor_t hxDesc, const void* hx, const miopenTensorDescriptor_t cxDesc, const void* cx, const miopenTensorDescriptor_t wDesc, const void* w, const miopenTensorDescriptor_t* yDesc, void* y, const miopenTensorDescriptor_t hyDesc, void* hy, const miopenTensorDescriptor_t cyDesc, void* cy, void* workSpace, size_t workSpaceNumBytes);
-  // CHECK: status = miopenRNNForwardInference(handle, RNNDescriptor, seqLength, &xD, x, hxD, hx, cxD, cx, filterDescriptor, W, &yD, y, hyD, hy, cyD, cy, workSpace, workSpaceSizeInBytes);
-  status = cudnnRNNForwardInference(handle, RNNDescriptor, seqLength, &xD, x, hxD, hx, cxD, cx, filterDescriptor, W, &yD, y, hyD, hy, cyD, cy, workSpace, workSpaceSizeInBytes);
+  // TODO [#837]: Insert int blank_label_id, bool apply_softmax_layer in the hipified miopenSetCTCLossDescriptor: will need variable declaration
+  // CUDA: cudnnStatus_t CUDNNWINAPI cudnnSetCTCLossDescriptor(cudnnCTCLossDescriptor_t ctcLossDesc, cudnnDataType_t compType);
+  // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenSetCTCLossDescriptor(miopenCTCLossDescriptor_t ctcLossDesc, miopenDataType_t dataType, const int blank_label_id, bool apply_softmax_layer);
+  // CHECK: status = miopenSetCTCLossDescriptor(CTCLossDescriptor, dataType);
+  status = cudnnSetCTCLossDescriptor(CTCLossDescriptor, dataType);
+#endif
 
   // CUDA: cudnnStatus_t CUDNNWINAPI cudnnDestroyRNNDescriptor(cudnnRNNDescriptor_t rnnDesc);
   // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenDestroyRNNDescriptor(miopenRNNDescriptor_t rnnDesc);
@@ -715,18 +712,6 @@ int main() {
   // CHECK: status = miopenDestroyCTCLossDescriptor(CTCLossDescriptor);
   status = cudnnDestroyCTCLossDescriptor(CTCLossDescriptor);
 
-  // TODO [#837]: Insert int* blank_label_id, bool* apply_softmax_layer in the hipified miopenGetCTCLossDescriptor: will need variable declaration
-  // CUDA: cudnnStatus_t CUDNNWINAPI cudnnGetCTCLossDescriptor(cudnnCTCLossDescriptor_t ctcLossDesc, cudnnDataType_t* compType);
-  // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenGetCTCLossDescriptor(miopenCTCLossDescriptor_t ctcLossDesc, miopenDataType_t* dataType, int* blank_label_id, bool* apply_softmax_layer);
-  // CHECK: status = miopenGetCTCLossDescriptor(CTCLossDescriptor, &dataType);
-  status = cudnnGetCTCLossDescriptor(CTCLossDescriptor, &dataType);
-
-  // TODO [#837]: Insert int blank_label_id, bool apply_softmax_layer in the hipified miopenSetCTCLossDescriptor: will need variable declaration
-  // CUDA: cudnnStatus_t CUDNNWINAPI cudnnSetCTCLossDescriptor(cudnnCTCLossDescriptor_t ctcLossDesc, cudnnDataType_t compType);
-  // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenSetCTCLossDescriptor(miopenCTCLossDescriptor_t ctcLossDesc, miopenDataType_t dataType, const int blank_label_id, bool apply_softmax_layer);
-  // CHECK: status = miopenSetCTCLossDescriptor(CTCLossDescriptor, dataType);
-  status = cudnnSetCTCLossDescriptor(CTCLossDescriptor, dataType);
-
   int labels = 0;
   int labelLengths = 0;
   int inputLengths = 0;