Optimize OpenCL Addition

Unify opencl addition for FP16/32
Add possibility to pass nullptr as local work size

**Self-evaluation:**
1. Build test: [X]Passed [ ]Failed [ ]Skipped
2. Run test:   [X]Passed [ ]Failed [ ]Skipped

Signed-off-by: Grzegorz Kisala <[email protected]>

Author: gkisalapl

nntrainer/opencl/opencl_command_queue_manager.cpp

@@ -387,4 +387,22 @@ bool CommandQueueManager::DispatchCommand(
   return true;
 }
 
+bool CommandQueueManager::DispatchCommandAndWait(
+  const cl_kernel kernel, const uint32_t work_dim,
+  const size_t *global_work_size, const size_t *local_work_size) {
+
+  const auto error_code = clEnqueueNDRangeKernel(
+    command_queue_, kernel, work_dim, nullptr, global_work_size,
+    local_work_size, 0, nullptr, nullptr);
+  if (error_code != CL_SUCCESS) {
+    ml_loge("Failed to clEnqueueNDRangeKernel. OpenCL error code: %d",
+            error_code);
+    return false;
+  }
+
+  clFinish(command_queue_);
+
+  return true;
+}
+
 } // namespace nntrainer::opencl

nntrainer/opencl/opencl_command_queue_manager.h

@@ -191,6 +191,19 @@ class CommandQueueManager : public Singleton<CommandQueueManager> {
                        const int (&work_group_size)[3],
                        cl_event *event = nullptr);
 
+  /**
+   * @brief Overloaded function to initiate execution of the command queue.
+   *
+   * @param kernel OpenCL kernel
+   * @param global_work_size Total number of work items that will execute the
+   * kernel function
+   * @param local_work_size Number of work items that make up a work group
+   * @return true if command queue execution is successful or false otherwise
+   */
+  bool DispatchCommandAndWait(const cl_kernel kernel, const uint32_t work_dim,
+                              const size_t *global_work_size,
+                              const size_t *local_work_size);
+
   /**
    * @brief Get the OpenCL Command Queue object
    *

nntrainer/tensor/cl_operations/blas_kernel_interface.cpp

@@ -220,23 +220,18 @@ void add_i_cl(Tensor &result, Tensor const &input) {
        result.channel() == input.channel() &&
        result.height() == input.height() && result.width() == input.width())) {
 
-    if (result.getDataType() == ml::train::TensorDim::DataType::FP32) {
-      float *Y = result.getData();
-      const float *X = input.getData();
+    const unsigned int size_input = input.size();
+    const unsigned int size_res = result.size();
 
-      for (unsigned int i = 0; i < result.batch() / input.batch(); ++i) {
-        axpy_cl(input.size(), 1.0f, X, Y);
-        Y += input.size();
-      }
+    if (result.getDataType() == ml::train::TensorDim::DataType::FP32) {
+      const auto *data_input = input.getData<float>();
+      auto *data_res = result.getData<float>();
+      addition_cl(data_input, data_res, size_input, size_res);
     } else if (result.getDataType() == ml::train::TensorDim::DataType::FP16) {
 #ifdef ENABLE_FP16
-      unsigned int size_res = result.size();
-      unsigned int size_input = input.size();
-      _FP16 *data_res = result.getData<_FP16>();
-      const _FP16 *data_input = input.getData<_FP16>();
-
+      const auto *data_input = input.getData<_FP16>();
+      auto *data_res = result.getData<_FP16>();
       addition_cl(data_input, data_res, size_input, size_res);
-
 #else
       throw std::invalid_argument("Error: enable-fp16 is not enabled");
 #endif

nntrainer/tensor/cl_operations/blas_kernel_strings.cpp

@@ -394,7 +394,6 @@ const std::string &getSgemmClTransABKernel() {
 const std::string &getAdditionClKernel() {
   static const std::string addition_cl_kernel_ =
     R"(__kernel void addition_cl(const __global float* input, __global float* output, unsigned int size_input, unsigned int size_res) {
-        #pragma printf_support
         size_t idx = get_global_id(0);
         if (idx < size_res) {
             output[idx] = output[idx] + input[idx % size_input];

nntrainer/tensor/cl_operations/blas_kernels_templates.h

@@ -280,14 +280,15 @@ addition_cl_internal(ClContext::SharedPtrClKernel kernel, const T *input,
     return;
   }
 
-  result = kernel->SetKernelArguments(0, clbuffInstance.getInBufferA(),
-                                      sizeof(cl_mem));
+  auto bufferInA = clbuffInstance.getInBufferA()->GetBuffer();
+  auto bufferOutA = clbuffInstance.getOutBufferA()->GetBuffer();
+
+  result = kernel->SetKernelArguments(0, &bufferInA, sizeof(cl_mem));
   if (!result) {
     return;
   }
 
-  result = kernel->SetKernelArguments(1, clbuffInstance.getOutBufferA(),
-                                      sizeof(cl_mem));
+  result = kernel->SetKernelArguments(1, &bufferOutA, sizeof(cl_mem));
   if (!result) {
     return;
   }
@@ -302,11 +303,11 @@ addition_cl_internal(ClContext::SharedPtrClKernel kernel, const T *input,
     return;
   }
 
-  const int work_groups_count[3] = {(int)size_res, 1, 1};
-  /// @todo: create a group size by device & input
-  const int work_group_size[3] = {1, 1, 1}; // test-value
-  result = blas_cc->command_queue_inst_.DispatchCommand(
-    kernel, work_groups_count, work_group_size);
+  std::array<size_t, 3> global_work_size = {size_res, 1, 1};
+
+  result = blas_cc->command_queue_inst_.DispatchCommandAndWait(
+    kernel->GetKernel(), global_work_size.size(), global_work_size.data(),
+    nullptr);
   if (!result) {
     return;
   }

test/unittest/unittest_blas_kernels_cl.cpp

@@ -441,8 +441,8 @@ TEST(blas_kernels, dot_gemm_50_768_2048_transAB) {
 TEST(blas_kernels, addition_i) {
   const int batch = 12;
   const int channel = 1;
-  const int height = 26;
-  const int width = 26;
+  const int height = 2048;
+  const int width = 2048;
 
   const int batch_b = 1;
 
@@ -474,8 +474,20 @@ TEST(blas_kernels, addition_i) {
                             MOD) *
                              alpha);
 
+  auto t1 = std::chrono::high_resolution_clock::now();
   A_fp32.add_i(B_fp32);
+  auto t2 = std::chrono::high_resolution_clock::now();
+  auto dt_cpu = std::chrono::duration_cast<std::chrono::microseconds>(t2 - t1);
+
+  auto t3 = std::chrono::high_resolution_clock::now();
   add_i_cl(C_fp32, D_fp32);
+  auto t4 = std::chrono::high_resolution_clock::now();
+  auto dt_gpu = std::chrono::duration_cast<std::chrono::microseconds>(t4 - t3);
+
+  std::cout << "FP32 ADD : N: " << batch << " C: " << channel
+            << " H: " << height << " W: " << width << std::endl;
+  std::cout << " - time : CPU = " << dt_cpu.count() << " us" << std::endl;
+  std::cout << " - time : GPU = " << dt_gpu.count() << " us" << std::endl;
 
   float mseError =
     mse<float>(A_fp32.getData<float>(), C_fp32.getData<float>(), A_fp32.size());