Enable fp8 support (#63)

* initial checkin

* rename the fp8 datatype name

* update based on cr comments

* resolve the build issue

* resolve fp8 campability issue

* fix minior bug and catch up to reflex latest develop branch change

* add fp8 + operatior support

* update fp8 header file

* resolve merge issue from develop branch

[ROCm/rccl-tests commit: 21e59fb283]

projects/rccl-tests/src/CMakeLists.txt

@@ -71,6 +71,7 @@ set(COMMON_FILES
   common.h
   common.cu
   nccl1_compat.h
+  rccl_bfloat8.h
   rccl_bfloat16.h
   timer.h
   timer.cc

projects/rccl-tests/src/common.cu

@@ -2,11 +2,13 @@
 /*************************************************************************
  * Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
  * Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
+ * Modifications Copyright (c) Microsoft Corporation. Licensed under the MIT License.
  *
  * See LICENSE.txt for license information
  ************************************************************************/
 
 #include "cuda_runtime.h"
+#include "rccl_bfloat8.h"
 #include "rccl_bfloat16.h"
 #include "common.h"
 #include <pthread.h>
@@ -28,12 +30,18 @@ int test_ncclVersion = 0; // init'd with ncclGetVersion()
   #if RCCL_BFLOAT16 == 1
     , ncclBfloat16
   #endif
+  #if RCCL_FLOAT8 == 1
+    , ncclFp8E4M3, ncclFp8E5M2
+  #endif
   };
   const char *test_typenames[ncclNumTypes] = {
     "int8", "uint8", "int32", "uint32", "int64", "uint64", "half", "float", "double"
   #if RCCL_BFLOAT16 == 1
     , "bfloat16"
   #endif
+  #if RCCL_FLOAT8 == 1
+    , "fp8_e4m3", "fp8_e5m2"
+  #endif
   };
   int test_typenum = -1;
 
@@ -100,13 +108,13 @@ static int enable_out_of_place = 1;
 static double parsesize(const char *value) {
     long long int units;
     double size;
-    char size_lit;
+    char size_lit[2];
 
-    int count = sscanf(value, "%lf %1s", &size, &size_lit);
+    int count = sscanf(value, "%lf %1s", &size, size_lit);
 
     switch (count) {
     case 2:
-      switch (size_lit) {
+      switch (size_lit[0]) {
       case 'G':
       case 'g':
         units = 1024*1024*1024;
@@ -401,6 +409,9 @@ testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
         #if defined(RCCL_BFLOAT16)
         rccl_bfloat16 bf16;
         #endif
+        #if defined(RCCL_FLOAT8)
+        rccl_float8 fp8_e4m3; rccl_bfloat8 fp8_e5m2;
+        #endif
       };
       switch(type) {
       case ncclInt8: i8 = ncclVerifiablePremulScalar<int8_t>(rank); break;
@@ -415,6 +426,11 @@ testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
       #if defined(RCCL_BFLOAT16)
       case ncclBfloat16: bf16 = ncclVerifiablePremulScalar<rccl_bfloat16>(rank); break;
       #endif
+      #if defined(RCCL_FLOAT8)
+      case ncclFp8E4M3: fp8_e4m3 = ncclVerifiablePremulScalar<rccl_float8>(rank); break;
+      case ncclFp8E5M2: fp8_e5m2 = ncclVerifiablePremulScalar<rccl_bfloat8>(rank); break;
+      #endif
+      case ncclNumTypes: break;
       }
       NCCLCHECK(ncclRedOpCreatePreMulSum(&op, &u64, type, ncclScalarHostImmediate, args->comms[i]));
     }
@@ -753,6 +769,10 @@ int main(int argc, char* argv[]) {
       #if defined(RCCL_BFLOAT16)
         test_typenum++; // bfloat16
       #endif
+      #if defined(RCCL_FLOAT8)
+        test_typenum++; // fp8_e4m3
+        test_typenum++; // fp8_e5m2
+      #endif
     }
     if (NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) && test_ncclVersion >= NCCL_VERSION(2,11,0)) {
       test_opnum++; // PreMulSum

projects/rccl-tests/src/common.h

@@ -1,6 +1,7 @@
 /*************************************************************************
  * Copyright (c) 2016-2022, NVIDIA CORPORATION. All rights reserved.
  * Modifications Copyright (c) 2019-2022 Advanced Micro Devices, Inc. All rights reserved.
+ * Modifications Copyright (c) Microsoft Corporation. Licensed under the MIT License.
  *
  * See LICENSE.txt for license information
  ************************************************************************/
@@ -221,6 +222,10 @@ static size_t wordSize(ncclDataType_t type) {
 #if NCCL_MAJOR >= 2
     //case ncclInt8:
     case ncclUint8:
+#if NCCL_MAJOR >= 2 && RCCL_FLOAT8 == 1
+    case ncclFp8E4M3:
+    case ncclFp8E5M2:
+#endif
 #endif
       return 1;
     case ncclHalf: