Update CK GeMM (#200)

* Upgrade CK to tag `rocm-6.0.0`
* Fix potential throughput bugs
* More robust correctness tests

---------

Co-authored-by: Binyang Li <[email protected]>

Author: chhwang

.azure-pipelines/ut.yml

@@ -44,6 +44,7 @@ jobs:
       script: |
         git submodule foreach --recursive git reset --hard
         git submodule foreach --recursive git clean -fdx
+        git submodule foreach git fetch
         git submodule update --init --recursive
 
   - task: Bash@3
@@ -63,7 +64,7 @@ jobs:
     inputs:
       targetType: 'inline'
       script: |
-        cd build && ARK_ROOT=$PWD ctest --stop-on-failure --verbose --schedule-random
+        cd build && ARK_ROOT=$PWD ARK_IGNORE_BINARY_CACHE=1 ctest --stop-on-failure --verbose --schedule-random
       workingDirectory: '$(System.DefaultWorkingDirectory)'
 
   - task: Bash@3

.github/workflows/ut-cuda.yml

@@ -41,6 +41,7 @@ jobs:
           git config --global --add safe.directory /__w/ark/ark
           git submodule foreach --recursive git reset --hard
           git submodule foreach --recursive git clean -fdx
+          git submodule foreach git fetch
           git submodule update --init --recursive
 
       - name: Build
@@ -51,7 +52,7 @@ jobs:
 
       - name: RunUT
         run: |
-          cd build && ARK_ROOT=$PWD ctest --stop-on-failure --verbose --schedule-random
+          cd build && ARK_ROOT=$PWD ARK_IGNORE_BINARY_CACHE=1 ctest --stop-on-failure --verbose --schedule-random
 
       - name: ReportCoverage
         run: |

.github/workflows/ut-rocm.yml

@@ -33,6 +33,7 @@ jobs:
           git config --global --add safe.directory /__w/ark/ark
           git submodule foreach --recursive git reset --hard
           git submodule foreach --recursive git clean -fdx
+          git submodule foreach git fetch
           git submodule update --init --recursive
 
       - name: Build
@@ -43,7 +44,7 @@ jobs:
 
       - name: RunUT
         run: |
-          cd build && ARK_ROOT=$PWD ctest --stop-on-failure --verbose --schedule-random -E "ops_matmul_test"
+          cd build && ARK_ROOT=$PWD ARK_IGNORE_BINARY_CACHE=1 ctest --stop-on-failure --verbose --schedule-random
 
       - name: ReportCoverage
         run: |

ark/bfloat16_test.cc

@@ -228,8 +228,26 @@ ark::unittest::State test_bfloat16() {
     return ark::unittest::SUCCESS;
 }
 
+ark::unittest::State test_bfloat16_error() {
+    ark::bfloat16_t x(0.1f);
+    ark::bfloat16_t sum(0.0f);
+    int reduce_length = 256;  // should not exceed 2^8
+
+    for (int i = 0; i < reduce_length; ++i) {
+        sum += x * x;
+    }
+
+    // max diff = 2^(-8) * x * 2 * reduce_length = 0.2
+    UNITTEST_LOG(float(sum));
+    UNITTEST_TRUE(float(sum) >= 2.36f);
+    UNITTEST_TRUE(float(sum) <= 2.76f);
+
+    return ark::unittest::SUCCESS;
+}
+
 int main() {
     ark::init();
     UNITTEST(test_bfloat16);
+    UNITTEST(test_bfloat16_error);
     return 0;
 }

ark/gpu/gpu_compile.cc

@@ -27,8 +27,6 @@
 
 #define ARK_DEBUG_KERNEL 0
 
-using namespace std;
-
 namespace ark {
 
 template <typename ItemType>
@@ -151,7 +149,7 @@ const std::string gpu_compile(const std::vector<std::string> &codes,
     items.reserve(codes.size());
     srand();
     for (auto &code : codes) {
-        string hash_str = fnv1a_hash(code);
+        std::string hash_str = fnv1a_hash(code);
         items.emplace_back(code, "/tmp/ark_" + hash_str);
     }
     assert(items.size() == 1);
@@ -187,25 +185,25 @@ const std::string gpu_compile(const std::vector<std::string> &codes,
             LOG(INFO, "Compiling: ", code_file_path);
             LOG(DEBUG, compile_cmd);
             // Run the command.
-            array<char, 4096> buffer;
-            stringstream exec_print;
-            unique_ptr<FILE, decltype(&pclose)> pipe(
+            std::array<char, 4096> buffer;
+            std::stringstream exec_print;
+            std::unique_ptr<FILE, decltype(&pclose)> pipe(
                 popen(compile_cmd.c_str(), "r"), pclose);
             if (!pipe) {
                 ERR(SystemError, "popen() failed");
             }
             while (fgets(buffer.data(), buffer.size(), pipe.get()) != nullptr) {
                 exec_print << buffer.data();
             }
-            string exec_print_str = exec_print.str();
+            std::string exec_print_str = exec_print.str();
             if (exec_print_str.size() > 0) {
                 ERR(ExecutorError, "\n", compile_cmd, "\n", exec_print_str,
                     "\n");
             }
             LOG(INFO, "Compile succeed: ", code_file_path, " (",
                 cpu_timer() - start, " seconds)");
         });
-    string gpubin_file_path = items[0].second + ".cubin";
+    std::string gpubin_file_path = items[0].second + ".cubin";
     return read_file(gpubin_file_path);
 }
 

ark/gpu/gpu_kernel_test.cc

@@ -18,7 +18,7 @@ ark::unittest::State test_gpu_kernel() {
 
 //
 const std::string test_kernel_loop_void =
-    "__device__ void ark_loop_body(int _iter) {\n"
+    "__device__ void ark_loop_body(char *_buf, int _iter) {\n"
     "  // Do nothing. Print iteration counter.\n"
     "  if (threadIdx.x == 0 && blockIdx.x == 0) {\n"
     "    if (_iter % 50 == 49) {\n"

ark/gpu/gpu_loop_kernel.cc

@@ -73,6 +73,7 @@ GpuLoopKernel::GpuLoopKernel(std::shared_ptr<GpuContext> ctx,
         "extern \"C\" __global__ __launch_bounds__(" << block_dim_[0] << ", 1)\n"
         "void " << kernel_name_ << "(int *_it)\n"
         "{\n"
+        "  char *_buf = " ARK_BUF_NAME ";\n"
         "  int *shared_mem = (int *)_ARK_SMEM;\n"
         "  for (int i = threadIdx.x; i < ARK_SMEM_RESERVED_BYTES / sizeof(int); i += blockDim.x) {\n"
         "    shared_mem[i] = 0;\n"
@@ -88,7 +89,7 @@ GpuLoopKernel::GpuLoopKernel(std::shared_ptr<GpuContext> ctx,
         "      return;\n"
         "    }\n"
         "    for (int _i = 0; _i < _ITER; ++_i) {\n"
-        "      ark_loop_body(_i);\n"
+        "      ark_loop_body(_buf, _i);\n"
         "      ark::sync_gpu<" << num_sm << ">(" ARK_LSS_NAME ");\n"
         "    }\n"
         "    if (threadIdx.x == 0 && blockIdx.x == 0) {\n"

ark/half_test.cc

@@ -227,8 +227,26 @@ ark::unittest::State test_half() {
     return ark::unittest::SUCCESS;
 }
 
+ark::unittest::State test_half_error() {
+    ark::half_t x(0.1f);
+    ark::half_t sum(0.0f);
+    int reduce_length = 1024;  // should not exceed 2^11
+
+    for (int i = 0; i < reduce_length; ++i) {
+        sum += x * x;
+    }
+
+    // max diff = 2^(-11) * x * 2 * reduce_length = 0.1
+    UNITTEST_LOG(float(sum));
+    UNITTEST_TRUE(float(sum) >= 10.14f);
+    UNITTEST_TRUE(float(sum) <= 10.34f);
+
+    return ark::unittest::SUCCESS;
+}
+
 int main() {
     ark::init();
     UNITTEST(test_half);
+    UNITTEST(test_half_error);
     return 0;
 }