merge amd-staging into amd-mainline

HIPIFY 6.4.0 beta - features complete Change-Id: Iababa494ccf5c919dc466ca613dbee7cf7922f2d
ROCm · Jan 8, 2025 · ae15539 · ae15539
2 parents dca2ac6 + 64b9f26
commit ae15539
Show file tree

Hide file tree

Showing 36 changed files with 1,435 additions and 1,088 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -3,6 +3,27 @@
 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.4.0
+
+### Added
+
+* CUDA 12.6.3 support
+* cuDNN 9.6.0 support
+* cuTENSOR 2.0.2.1 support
+* LLVM 19.1.6 support
+* Full support for direct hipification of `cuRAND` into `rocRAND` under the `--roc` option
+* [#1617](https://github.com/ROCm/HIPIFY/issues/1617) Support for `fp8` math device/host API
+
+### Resolved issues
+
+* `MIOpen` support in hipify-perl under the `-miopen` option
+* [#1769](https://github.com/ROCm/HIPIFY/issues/1769) Support for `fp16` device/host API
+* [#1800](https://github.com/ROCm/HIPIFY/issues/1800) Fix instructions on building LLVM for HIPIFY on Linux
+
+### Known issues
+
+* [#833](https://github.com/ROCm/HIPIFY/issues/833) `hipify-clang` build failure against LLVM 15-18 on `Ubuntu`, `CentOS`, and `Fedora`
+
 ## HIPIFY for ROCm 6.3.1
 
 ### Added
@@ -27,15 +48,15 @@ Documentation for HIPIFY is available at
 * `rocBLAS` 64-bit APIs support
 * Initial support for direct hipification of `cuDNN` into `MIOpen` under the `--roc` option
 * Initial support for direct hipification of `cuRAND` into `rocRAND` under the `--roc` option
-* [#1650] Added a filtering ability for the supplementary hipification scripts
+* [#1650](https://github.com/ROCm/HIPIFY/pull/1650) Added a filtering ability for the supplementary hipification scripts
 
 ### Resolved issues
 
 * Correct `roc` header files support
 
 ### Known issues
 
-* [#1617] Support for `fp8` data types
+* [#1617](https://github.com/ROCm/HIPIFY/issues/1617) Support for `fp8` data types
 
 ## HIPIFY for ROCm 6.2.4
 
@@ -122,8 +143,8 @@ Documentation for HIPIFY is available at
 
 ### Known issues
 
-* [#837] Added a new function to call transformation type "additional non-const arg"
-* [#1014] Added a new function to call transformation type "replace argument with a const"
+* [#837](https://github.com/ROCm/HIPIFY/issues/837) Added a new function to call transformation type "additional non-const arg"
+* [#1014](https://github.com/ROCm/HIPIFY/issues/1014) Added a new function to call transformation type "replace argument with a const"
 
 ## HIPIFY for ROCm 5.7.0
 
@@ -143,8 +164,8 @@ Documentation for HIPIFY is available at
 
 ### Known issues
 
-* [#822] Added a new function to call transformation type "additional const by value arg"
-* [#830] Added a new function to call transformation type "move arg from place X to place Y"
+* [#822](https://github.com/ROCm/HIPIFY/issues/822) Added a new function to call transformation type "additional const by value arg"
+* [#830](https://github.com/ROCm/HIPIFY/issues/830) Added a new function to call transformation type "move arg from place X to place Y"
 
 ## HIPIFY for ROCm 5.6.0
 

diff --git a/LICENSE.txt b/LICENSE.txt
@@ -1,4 +1,4 @@
-Copyright (c) 2024 Advanced Micro Devices, Inc.
+Copyright (c) 2025 Advanced Micro Devices, Inc.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal

diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -1398,15 +1398,28 @@ my %removed_funcs = (
 );
 
 my %experimental_funcs = (
+    "cudaEventRecordWithFlags" => "6.4.0",
+    "cudaErrorInvalidTexture" => "6.4.0",
+    "cudaErrorInvalidChannelDescriptor" => "6.4.0",
     "cuStreamBatchMemOp_v2" => "6.4.0",
     "cuStreamBatchMemOp" => "6.4.0",
     "cuGraphExecBatchMemOpNodeSetParams" => "6.4.0",
     "cuGraphBatchMemOpNodeSetParams" => "6.4.0",
     "cuGraphBatchMemOpNodeGetParams" => "6.4.0",
     "cuGraphAddBatchMemOpNode" => "6.4.0",
+    "cuEventRecordWithFlags" => "6.4.0",
+    "CUstreamBatchMemOpType_enum" => "6.4.0",
+    "CUstreamBatchMemOpType" => "6.4.0",
     "CUstreamBatchMemOpParams_v1" => "6.4.0",
     "CUstreamBatchMemOpParams_union" => "6.4.0",
     "CUstreamBatchMemOpParams" => "6.4.0",
+    "CU_STREAM_MEM_OP_WRITE_VALUE_64" => "6.4.0",
+    "CU_STREAM_MEM_OP_WRITE_VALUE_32" => "6.4.0",
+    "CU_STREAM_MEM_OP_WAIT_VALUE_64" => "6.4.0",
+    "CU_STREAM_MEM_OP_WAIT_VALUE_32" => "6.4.0",
+    "CU_STREAM_MEM_OP_FLUSH_REMOTE_WRITES" => "6.4.0",
+    "CU_STREAM_MEM_OP_BARRIER" => "6.4.0",
+    "CU_GRAPH_NODE_TYPE_BATCH_MEM_OP" => "6.4.0",
     "CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st" => "6.4.0",
     "CUDA_BATCH_MEM_OP_NODE_PARAMS_v2" => "6.4.0",
     "CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st" => "6.4.0",
@@ -1550,6 +1563,8 @@ sub subst {
 }
 
 sub experimentalSubstitutions {
+    subst("cuEventRecordWithFlags", "hipEventRecordWithFlags", "event");
+    subst("cudaEventRecordWithFlags", "hipEventRecordWithFlags", "event");
     subst("cuStreamBatchMemOp", "hipStreamBatchMemOp", "stream_memory");
     subst("cuStreamBatchMemOp_v2", "hipStreamBatchMemOp", "stream_memory");
     subst("cuGraphAddBatchMemOpNode", "hipGraphAddBatchMemOpNode", "graph");
@@ -1565,6 +1580,17 @@ sub experimentalSubstitutions {
     subst("CUstreamBatchMemOpParams", "hipStreamBatchMemOpParams", "type");
     subst("CUstreamBatchMemOpParams_union", "hipStreamBatchMemOpParams_union", "type");
     subst("CUstreamBatchMemOpParams_v1", "hipStreamBatchMemOpParams", "type");
+    subst("CUstreamBatchMemOpType", "hipStreamBatchMemOpType", "type");
+    subst("CUstreamBatchMemOpType_enum", "hipStreamBatchMemOpType", "type");
+    subst("CU_GRAPH_NODE_TYPE_BATCH_MEM_OP", "hipGraphNodeTypeBatchMemOp", "numeric_literal");
+    subst("CU_STREAM_MEM_OP_BARRIER", "hipStreamMemOpBarrier", "numeric_literal");
+    subst("CU_STREAM_MEM_OP_FLUSH_REMOTE_WRITES", "hipStreamMemOpFlushRemoteWrites", "numeric_literal");
+    subst("CU_STREAM_MEM_OP_WAIT_VALUE_32", "hipStreamMemOpWaitValue32", "numeric_literal");
+    subst("CU_STREAM_MEM_OP_WAIT_VALUE_64", "hipStreamMemOpWaitValue64", "numeric_literal");
+    subst("CU_STREAM_MEM_OP_WRITE_VALUE_32", "hipStreamMemOpWriteValue32", "numeric_literal");
+    subst("CU_STREAM_MEM_OP_WRITE_VALUE_64", "hipStreamMemOpWriteValue64", "numeric_literal");
+    subst("cudaErrorInvalidChannelDescriptor", "hipErrorInvalidChannelDescriptor", "numeric_literal");
+    subst("cudaErrorInvalidTexture", "hipErrorInvalidTexture", "numeric_literal");
 }
 
 sub rocSubstitutions {
@@ -3538,6 +3564,7 @@ sub MIOpenSubstitutions {
     subst("CUDNN_BACKEND_OPERATION_REDUCTION_DESCRIPTOR", "MIOPEN_BACKEND_OPERATION_REDUCTION_DESCRIPTOR", "numeric_literal");
     subst("CUDNN_BACKEND_OPERATION_RESAMPLE_BWD_DESCRIPTOR", "MIOPEN_BACKEND_OPERATION_RESAMPLE_BWD_DESCRIPTOR", "numeric_literal");
     subst("CUDNN_BACKEND_OPERATION_RESAMPLE_FWD_DESCRIPTOR", "MIOPEN_BACKEND_OPERATION_RESAMPLE_FWD_DESCRIPTOR", "numeric_literal");
+    subst("CUDNN_BACKEND_OPERATION_RESHAPE_DESCRIPTOR", "MIOPEN_BACKEND_OPERATION_RESHAPE_DESCRIPTOR", "numeric_literal");
     subst("CUDNN_BACKEND_OPERATION_RNG_DESCRIPTOR", "MIOPEN_BACKEND_OPERATION_RNG_DESCRIPTOR", "numeric_literal");
     subst("CUDNN_BACKEND_OPERATION_SIGNAL_DESCRIPTOR", "MIOPEN_BACKEND_OPERATION_SIGNAL_DESCRIPTOR", "numeric_literal");
     subst("CUDNN_BACKEND_POINTWISE_DESCRIPTOR", "MIOPEN_BACKEND_POINTWISE_DESCRIPTOR", "numeric_literal");
@@ -6771,6 +6798,7 @@ sub simpleSubstitutions {
     subst("curand_poisson.h", "hiprand\/hiprand_kernel.h", "include");
     subst("curand_precalc.h", "hiprand\/hiprand_kernel.h", "include");
     subst("curand_uniform.h", "hiprand\/hiprand_kernel.h", "include");
+    subst("cutensorMg.h", "hiptensor.h", "include");
     subst("device_functions.h", "hip\/device_functions.h", "include");
     subst("device_launch_parameters.h", "", "include");
     subst("driver_types.h", "hip\/driver_types.h", "include");
@@ -7384,6 +7412,7 @@ sub simpleSubstitutions {
     subst("cusparseSparseToDenseAlg_t", "hipsparseSparseToDenseAlg_t", "type");
     subst("cusparseStatus_t", "hipsparseStatus_t", "type");
     subst("cutensorAlgo_t", "hiptensorAlgo_t", "type");
+    subst("cutensorComputeType_t", "hiptensorComputeType_t", "type");
     subst("cutensorContractionPlan_t", "hiptensorContractionPlan_t", "type");
     subst("cutensorDataType_t", "hiptensorComputeType_t", "type");
     subst("cutensorHandle_t", "hiptensorHandle_t", "type");
@@ -8694,7 +8723,6 @@ sub simpleSubstitutions {
     subst("cudaSuccess", "hipSuccess", "numeric_literal");
     subst("cudaUserObjectNoDestructorSync", "hipUserObjectNoDestructorSync", "numeric_literal");
     subst("cusolver_int_t", "int", "numeric_literal");
-    subst("cutensorComputeType_t", "hiptensorComputeType_t", "numeric_literal");
     subst("CUB_MAX", "CUB_MAX", "define");
     subst("CUB_MIN", "CUB_MIN", "define");
     subst("CUB_NAMESPACE_BEGIN", "BEGIN_HIPCUB_NAMESPACE", "define");
@@ -9907,6 +9935,42 @@ sub warnHipOnlyUnsupportedFunctions {
         "cutensorOperationDescriptorSetAttribute",
         "cutensorOperationDescriptorGetAttribute",
         "cutensorOperationDescriptorAttribute_t",
+        "cutensorMgTensorDescriptor_t",
+        "cutensorMgTensorDescriptor_s",
+        "cutensorMgHostDevice_t",
+        "cutensorMgHandle_t",
+        "cutensorMgHandle_s",
+        "cutensorMgDestroyTensorDescriptor",
+        "cutensorMgDestroyCopyPlan",
+        "cutensorMgDestroyCopyDescriptor",
+        "cutensorMgDestroyContractionPlan",
+        "cutensorMgDestroyContractionFind",
+        "cutensorMgDestroyContractionDescriptor",
+        "cutensorMgDestroy",
+        "cutensorMgCreateTensorDescriptor",
+        "cutensorMgCreateCopyPlan",
+        "cutensorMgCreateCopyDescriptor",
+        "cutensorMgCreateContractionPlan",
+        "cutensorMgCreateContractionFind",
+        "cutensorMgCreateContractionDescriptor",
+        "cutensorMgCreate",
+        "cutensorMgCopyPlan_t",
+        "cutensorMgCopyPlan_s",
+        "cutensorMgCopyGetWorkspace",
+        "cutensorMgCopyDescriptor_t",
+        "cutensorMgCopyDescriptor_s",
+        "cutensorMgCopy",
+        "cutensorMgContractionPlan_t",
+        "cutensorMgContractionPlan_s",
+        "cutensorMgContractionGetWorkspace",
+        "cutensorMgContractionFind_t",
+        "cutensorMgContractionFind_s",
+        "cutensorMgContractionFindSetAttribute",
+        "cutensorMgContractionFindAttribute_t",
+        "cutensorMgContractionDescriptor_t",
+        "cutensorMgContractionDescriptor_s",
+        "cutensorMgContraction",
+        "cutensorMgAlgo_t",
         "cutensorJitMode_t",
         "cutensorHandleWritePlanCacheToFile",
         "cutensorHandleResizePlanCache",
@@ -10246,8 +10310,10 @@ sub warnHipOnlyUnsupportedFunctions {
         "CUTENSOR_STATUS_CUBLAS_ERROR",
         "CUTENSOR_R_MIN_TF32",
         "CUTENSOR_R_MIN_8U",
+        "CUTENSOR_R_MIN_8I",
         "CUTENSOR_R_MIN_64F",
         "CUTENSOR_R_MIN_32U",
+        "CUTENSOR_R_MIN_32I",
         "CUTENSOR_R_MIN_32F",
         "CUTENSOR_R_MIN_16F",
         "CUTENSOR_R_MIN_16BF",
@@ -10297,6 +10363,8 @@ sub warnHipOnlyUnsupportedFunctions {
         "CUTENSOR_OPERATION_DESCRIPTOR_PADDING_LEFT",
         "CUTENSOR_OPERATION_DESCRIPTOR_MOVED_BYTES",
         "CUTENSOR_OPERATION_DESCRIPTOR_FLOPS",
+        "CUTENSOR_MG_DEVICE_HOST_PINNED",
+        "CUTENSOR_MG_DEVICE_HOST",
         "CUTENSOR_JIT_MODE_NONE",
         "CUTENSOR_JIT_MODE_DEFAULT",
         "CUTENSOR_C_MIN_TF32",
@@ -10316,13 +10384,16 @@ sub warnHipOnlyUnsupportedFunctions {
         "CUTENSOR_C_16F",
         "CUTENSOR_C_16BF",
         "CUTENSOR_COMPUTE_TF32",
+        "CUTENSOR_COMPUTE_3XTF32",
         "CUTENSOR_CACHE_MODE_PEDANTIC",
         "CUTENSOR_CACHE_MODE_NONE",
         "CUTENSOR_AUTOTUNE_MODE_NONE",
         "CUTENSOR_AUTOTUNE_MODE_INCREMENTAL",
         "CUTENSOR_ALGO_TTGT",
         "CUTENSOR_ALGO_TGETT",
         "CUTENSOR_ALGO_GETT",
+        "CUTENSORMG_CONTRACTION_FIND_ATTRIBUTE_MAX",
+        "CUTENSORMG_ALGO_DEFAULT",
         "CUSPARSE_SPSV_UPDATE_GENERAL",
         "CUSPARSE_SPSV_UPDATE_DIAGONAL",
         "CUSPARSE_SPSM_UPDATE_GENERAL",
@@ -11269,6 +11340,42 @@ sub warnRocOnlyUnsupportedFunctions {
         "cutensorOperationDescriptorSetAttribute",
         "cutensorOperationDescriptorGetAttribute",
         "cutensorOperationDescriptorAttribute_t",
+        "cutensorMgTensorDescriptor_t",
+        "cutensorMgTensorDescriptor_s",
+        "cutensorMgHostDevice_t",
+        "cutensorMgHandle_t",
+        "cutensorMgHandle_s",
+        "cutensorMgDestroyTensorDescriptor",
+        "cutensorMgDestroyCopyPlan",
+        "cutensorMgDestroyCopyDescriptor",
+        "cutensorMgDestroyContractionPlan",
+        "cutensorMgDestroyContractionFind",
+        "cutensorMgDestroyContractionDescriptor",
+        "cutensorMgDestroy",
+        "cutensorMgCreateTensorDescriptor",
+        "cutensorMgCreateCopyPlan",
+        "cutensorMgCreateCopyDescriptor",
+        "cutensorMgCreateContractionPlan",
+        "cutensorMgCreateContractionFind",
+        "cutensorMgCreateContractionDescriptor",
+        "cutensorMgCreate",
+        "cutensorMgCopyPlan_t",
+        "cutensorMgCopyPlan_s",
+        "cutensorMgCopyGetWorkspace",
+        "cutensorMgCopyDescriptor_t",
+        "cutensorMgCopyDescriptor_s",
+        "cutensorMgCopy",
+        "cutensorMgContractionPlan_t",
+        "cutensorMgContractionPlan_s",
+        "cutensorMgContractionGetWorkspace",
+        "cutensorMgContractionFind_t",
+        "cutensorMgContractionFind_s",
+        "cutensorMgContractionFindSetAttribute",
+        "cutensorMgContractionFindAttribute_t",
+        "cutensorMgContractionDescriptor_t",
+        "cutensorMgContractionDescriptor_s",
+        "cutensorMgContraction",
+        "cutensorMgAlgo_t",
         "cutensorJitMode_t",
         "cutensorHandleWritePlanCacheToFile",
         "cutensorHandleResizePlanCache",
@@ -11698,8 +11805,10 @@ sub warnRocOnlyUnsupportedFunctions {
         "CUTENSOR_STATUS_CUBLAS_ERROR",
         "CUTENSOR_R_MIN_TF32",
         "CUTENSOR_R_MIN_8U",
+        "CUTENSOR_R_MIN_8I",
         "CUTENSOR_R_MIN_64F",
         "CUTENSOR_R_MIN_32U",
+        "CUTENSOR_R_MIN_32I",
         "CUTENSOR_R_MIN_32F",
         "CUTENSOR_R_MIN_16F",
         "CUTENSOR_R_MIN_16BF",
@@ -11749,6 +11858,8 @@ sub warnRocOnlyUnsupportedFunctions {
         "CUTENSOR_OPERATION_DESCRIPTOR_PADDING_LEFT",
         "CUTENSOR_OPERATION_DESCRIPTOR_MOVED_BYTES",
         "CUTENSOR_OPERATION_DESCRIPTOR_FLOPS",
+        "CUTENSOR_MG_DEVICE_HOST_PINNED",
+        "CUTENSOR_MG_DEVICE_HOST",
         "CUTENSOR_JIT_MODE_NONE",
         "CUTENSOR_JIT_MODE_DEFAULT",
         "CUTENSOR_C_MIN_TF32",
@@ -11768,13 +11879,16 @@ sub warnRocOnlyUnsupportedFunctions {
         "CUTENSOR_C_16F",
         "CUTENSOR_C_16BF",
         "CUTENSOR_COMPUTE_TF32",
+        "CUTENSOR_COMPUTE_3XTF32",
         "CUTENSOR_CACHE_MODE_PEDANTIC",
         "CUTENSOR_CACHE_MODE_NONE",
         "CUTENSOR_AUTOTUNE_MODE_NONE",
         "CUTENSOR_AUTOTUNE_MODE_INCREMENTAL",
         "CUTENSOR_ALGO_TTGT",
         "CUTENSOR_ALGO_TGETT",
         "CUTENSOR_ALGO_GETT",
+        "CUTENSORMG_CONTRACTION_FIND_ATTRIBUTE_MAX",
+        "CUTENSORMG_ALGO_DEFAULT",
         "CUSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED",
         "CUSPARSE_STATUS_MAPPING_ERROR",
         "CUSPARSE_STATUS_INSUFFICIENT_RESOURCES",
@@ -13374,7 +13488,6 @@ sub warnMIOpenOnlyUnsupportedFunctions {
         "CUDNN_BATCHNORM_OPS_BN_ADD_ACTIVATION",
         "CUDNN_BATCHNORM_OPS_BN_ACTIVATION",
         "CUDNN_BATCHNORM_OPS_BN",
-        "CUDNN_BACKEND_OPERATION_RESHAPE_DESCRIPTOR",
         "CUDNN_BACKEND_OPERATION_PAGED_CACHE_LOAD_DESCRIPTOR",
         "CUDNN_BACKEND_OPERATION_BN_FINALIZE_STATISTICS_DESCRIPTOR",
         "CUDNN_BACKEND_OPERATION_BN_BWD_WEIGHTS_DESCRIPTOR",