[HIPIFY][ROCm#1617][fp8][tests] Support for fp8 math - Part 3 - Tests

+ Added missing header files mapping + Updated the regenerated `hipify-perl` + [MISC] Minor fixes in the rest of synthetic tests
emankov · Nov 22, 2024 · 76c63d9 · 76c63d9
1 parent ddc3b10
commit 76c63d9
Show file tree

Hide file tree

Showing 23 changed files with 83 additions and 18 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -6161,6 +6161,7 @@ sub simpleSubstitutions {
     subst("cublasLt.h", "hipblaslt.h", "include");
     subst("cublas_api.h", "hipblas.h", "include");
     subst("cuda_fp16.h", "hip\/hip_fp16.h", "include");
+    subst("cuda_fp8.h", "hip\/hip_fp8.h", "include");
     subst("cuda_profiler_api.h", "hip\/hip_runtime_api.h", "include");
     subst("cuda_runtime_api.h", "hip\/hip_runtime_api.h", "include");
     subst("cuda_texture_types.h", "hip\/hip_texture_types.h", "include");

diff --git a/src/CUDA2HIP.cpp b/src/CUDA2HIP.cpp
@@ -33,6 +33,7 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP {
   {"device_functions.h",                                    {"hip/device_functions.h",                                "", CONV_INCLUDE,                API_RUNTIME, 0}},
   {"driver_types.h",                                        {"hip/driver_types.h",                                    "", CONV_INCLUDE,                API_RUNTIME, 0}},
   {"cuda_fp16.h",                                           {"hip/hip_fp16.h",                                        "", CONV_INCLUDE,                API_RUNTIME, 0}},
+  {"cuda_fp8.h",                                            {"hip/hip_fp8.h",                                         "", CONV_INCLUDE,                API_RUNTIME, 0}},
   {"cuda_texture_types.h",                                  {"hip/hip_texture_types.h",                               "", CONV_INCLUDE,                API_RUNTIME, 0}},
   {"texture_fetch_functions.h",                             {"",                                                      "", CONV_INCLUDE,                API_RUNTIME, 0}},
   {"vector_types.h",                                        {"hip/hip_vector_types.h",                                "", CONV_INCLUDE,                API_RUNTIME, 0}},

diff --git a/tests/lit.cfg b/tests/lit.cfg
@@ -127,6 +127,9 @@ if config.cuda_version_major < 11 or (config.cuda_version_major == 11 and config
 if config.cuda_version_major >= 11:
     config.excludes.append('cusparse2rocsparse_before_11000.cu')
 
+if config.cuda_version_major < 11 or (config.cuda_version_major == 11 and config.cuda_version_minor < 8):
+    config.excludes.append('cudevice2hipdevice.cu')
+
 if config.cuda_version_major < 12:
     config.excludes.append('headers_test_06_12000.cu')
     config.excludes.append('headers_test_07_12000.cu')

diff --git a/tests/unit_tests/synthetic/libraries/cublaslt2hipblaslt.cu b/tests/unit_tests/synthetic/libraries/cublaslt2hipblaslt.cu
@@ -13,7 +13,7 @@
 #endif
 
 int main() {
-  printf("20. cuBLASLt API to hipBLASLt API synthetic test\n");
+  printf("21. cuBLASLt API to hipBLASLt API synthetic test\n");
 
   // CHECK: hipblasLtHandle_t blasLtHandle;
   cublasLtHandle_t blasLtHandle;

diff --git a/tests/unit_tests/synthetic/libraries/cublaslt2hipblaslt_10010_10020.cu b/tests/unit_tests/synthetic/libraries/cublaslt2hipblaslt_10010_10020.cu
@@ -8,7 +8,7 @@
 // CHECK-NOT: #include "hipblaslt.h"
 
 int main() {
-  printf("20.1. cuBLASLt API to hipBLASLt API synthetic test\n");
+  printf("21.10010.10020. cuBLASLt API to hipBLASLt API synthetic test\n");
 
 #if CUDA_VERSION >= 10010 && CUDA_VERSION <= 10020
   // CHECK: hipblasLtMatrixLayoutOpaque_t blasLtMatrixLayoutOpaque;

diff --git a/tests/unit_tests/synthetic/libraries/cudevice2hipdevice.cu b/tests/unit_tests/synthetic/libraries/cudevice2hipdevice.cu
@@ -0,0 +1,58 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args 3 --amap --skip-excluded-preprocessor-conditional-blocks --experimental %clang_args -ferror-limit=500
+
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda_runtime.h>
+#include <stdio.h>
+// CHECK: #include "hip/hip_fp8.h"
+#include "cuda_fp8.h"
+// CHECK-NOT: #include "hip/hip_fp8.h"
+// CHECK-NOT: #include "cuda_fp8.h"
+
+int main() {
+  printf("24. CUDA Device API to HIP Device API synthetic test\n");
+
+#if CUDA_VERSION >= 11080
+  // CHECK: __hip_fp8_storage_t fp8_storage_t;
+  __nv_fp8_storage_t fp8_storage_t;
+
+  // CHECK: __hip_fp8x2_storage_t fp8x2_storage_t;
+  __nv_fp8x2_storage_t fp8x2_storage_t;
+
+  // CHECK: __hip_fp8x4_storage_t fp8x4_storage_t;
+  __nv_fp8x4_storage_t fp8x4_storage_t;
+
+  // CHECK: __hip_fp8_e5m2_fnuz fp8_e5m2;
+  __nv_fp8_e5m2 fp8_e5m2;
+
+  // CHECK: __hip_fp8x2_e5m2_fnuz fp8x2_e5m2;
+  __nv_fp8x2_e5m2 fp8x2_e5m2;
+
+  // CHECK: __hip_fp8_e4m3_fnuz fp8_e4m3;
+  __nv_fp8_e4m3 fp8_e4m3;
+
+  // CHECK: __hip_fp8x2_e4m3_fnuz fp8x2_e4m3;
+  __nv_fp8x2_e4m3 fp8x2_e4m3;
+
+  // CHECK: __hip_fp8x4_e4m3_fnuz fp8x4_e4m3;
+  __nv_fp8x4_e4m3 fp8x4_e4m3;
+
+  // CHECK: __hip_saturation_t saturation_t;
+  // CHECK-NEXT: __hip_saturation_t NOSAT = __HIP_NOSAT;
+  // CHECK-NEXT: __hip_saturation_t SATFINITE = __HIP_SATFINITE;
+  __nv_saturation_t saturation_t;
+  __nv_saturation_t NOSAT = __NV_NOSAT;
+  __nv_saturation_t SATFINITE = __NV_SATFINITE;
+
+  // CHECK: __hip_fp8_interpretation_t fp8_interpretation_t;
+  // CHECK-NEXT: __hip_fp8_interpretation_t E4M3 = __HIP_E4M3_FNUZ;
+  // CHECK-NEXT: __hip_fp8_interpretation_t E5M2 = __HIP_E5M2_FNUZ;
+  __nv_fp8_interpretation_t fp8_interpretation_t;
+  __nv_fp8_interpretation_t E4M3 = __NV_E4M3;
+  __nv_fp8_interpretation_t E5M2 = __NV_E5M2;
+
+  // CHECK: __hip_fp8x4_e5m2_fnuz fp8x4_e5m2;
+  __nv_fp8x4_e5m2 fp8x4_e5m2;
+#endif
+
+  return 0;
+}
diff --git a/tests/unit_tests/synthetic/libraries/cudnn2miopen_before_9000.cu b/tests/unit_tests/synthetic/libraries/cudnn2miopen_before_9000.cu
@@ -7,7 +7,7 @@
 #include "cudnn.h"
 
 int main() {
-  printf("15.1. cuDNN API to MIOpen API synthetic test\n");
+  printf("15.before.9000. cuDNN API to MIOpen API synthetic test\n");
 
   // CHECK: miopenStatus_t status;
   cudnnStatus_t status;

diff --git a/tests/unit_tests/synthetic/libraries/curand2hiprand.cu b/tests/unit_tests/synthetic/libraries/curand2hiprand.cu
@@ -16,7 +16,7 @@
 #endif
 
 int main() {
-  printf("21. cuRAND API to hipRAND API synthetic test\n");
+  printf("22. cuRAND API to hipRAND API synthetic test\n");
 
   unsigned int *outputPtr = nullptr;
   unsigned int *constants = nullptr;

diff --git a/tests/unit_tests/synthetic/libraries/curand2rocrand.cu b/tests/unit_tests/synthetic/libraries/curand2rocrand.cu
@@ -11,7 +11,7 @@
 // CHECK-NOT: #include "rocrand/rocrand_kernel.h"
 
 int main() {
-  printf("21.1. cuRAND API to rocRAND API synthetic test\n");
+  printf("23. cuRAND API to rocRAND API synthetic test\n");
 
   unsigned int *outputPtr = nullptr;
   unsigned int *constants = nullptr;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2hipsparse_11030_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2hipsparse_11030_12000.cu
@@ -15,7 +15,7 @@
 #endif
 
 int main() {
-  printf("17.1. cuSPARSE API to hipSPARSE API synthetic test\n");
+  printf("17.11030.12000. cuSPARSE API to hipSPARSE API synthetic test\n");
 
   // CHECK: hipsparseStatus_t status_t;
   cusparseStatus_t status_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2hipsparse_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2hipsparse_12000.cu
@@ -15,7 +15,7 @@
 #endif
 
 int main() {
-  printf("17.1. cuSPARSE API to hipSPARSE API synthetic test\n");
+  printf("17.12000. cuSPARSE API to hipSPARSE API synthetic test\n");
 
   // CHECK: hipsparseStatus_t status_t;
   cusparseStatus_t status_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_10000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_10000.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.10000. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: _rocsparse_handle *handle = nullptr;
   // CHECK-NEXT: rocsparse_handle handle_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_10010.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_10010.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.10010. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: _rocsparse_handle *handle = nullptr;
   // CHECK-NEXT: rocsparse_handle handle_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_10010_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_10010_12000.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.10010.12000. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: _rocsparse_handle *handle = nullptr;
   // CHECK-NEXT: rocsparse_handle handle_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_11010_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_11010_12000.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.11010.12000. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: rocsparse_status status_t;
   cusparseStatus_t status_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_11030_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_11030_12000.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.11030.12000. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: rocsparse_status status_t;
   cusparseStatus_t status_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_12000.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.12000. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: rocsparse_status status_t;
   cusparseStatus_t status_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_7050.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_7050.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.7050. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: _rocsparse_handle *handle = nullptr;
   // CHECK-NEXT: rocsparse_handle handle_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_9020.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_9020.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.9020. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: _rocsparse_handle *handle = nullptr;
   // CHECK-NEXT: rocsparse_handle handle_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_9020_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_9020_12000.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.9020.12000. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: _rocsparse_handle *handle = nullptr;
   // CHECK-NEXT: rocsparse_handle handle_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_before_11000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_before_11000.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.before.11000. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: rocsparse_status status_t;
   cusparseStatus_t status_t;

diff --git a/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_before_12000.cu b/tests/unit_tests/synthetic/libraries/cusparse2rocsparse_before_12000.cu
@@ -10,7 +10,7 @@
 // CHECK-NOT: #include "rocsparse.h"
 
 int main() {
-  printf("18.1. cuSPARSE API to rocSPARSE API synthetic test\n");
+  printf("18.before.12000. cuSPARSE API to rocSPARSE API synthetic test\n");
 
   // CHECK: rocsparse_status status_t;
   cusparseStatus_t status_t;

diff --git a/tests/unit_tests/synthetic/libraries/cutensor2hiptensor.cu b/tests/unit_tests/synthetic/libraries/cutensor2hiptensor.cu
@@ -79,4 +79,6 @@ int main() {
   // CHECK: hiptensorAlgo_t TENSOR_ALGO_DEFAULT_PATIENT = HIPTENSOR_ALGO_DEFAULT_PATIENT;
   cutensorAlgo_t TENSOR_ALGO_DEFAULT_PATIENT = CUTENSOR_ALGO_DEFAULT_PATIENT;
 #endif
+
+  return 0;
 }