[IGPU]: test new Kernel (V3/V4)

- update low N gemm kernel
- add analyse trace.
- corrected syncthreads

Author: Djip007

README.md

@@ -11,19 +11,30 @@
 
 #include <iostream>
 #include <vector>
+#include <unordered_set>
 
 //#define IGPU_TRACE(...) std::cout << "#> ggml-igpu: " << __VA_ARGS__ << std::endl
 #define IGPU_TRACE(...)
+//#define IGPU_DEV(...) std::cout << "#> ggml-igpu: " << __VA_ARGS__ << std::endl
+#define IGPU_DEV(...)
 
-//#define BLOC_V1
-#define BLOC_V2
+//#define BLOC_V1    // mieux? N=[23-47]  => Voir comment faire "mieux" dans cette bande.
+//#define BLOC_V2
+//#define BLOC_V3
+#define BLOC_V4  // OK N=[1-22] N=[48...]
 
 #ifdef BLOC_V1
 #include "mulmat-bf16bloc_V1.h"
 #endif
 #ifdef BLOC_V2
 #include "mulmat-bf16bloc_V2.h"
 #endif
+#ifdef BLOC_V3
+#include "mulmat-bf16bloc_V3.h"
+#endif
+#ifdef BLOC_V4
+#include "mulmat-bf16bloc_V4.h"
+#endif
 
 /*
 #> version bloc-bf16 V0.
@@ -72,7 +83,7 @@ namespace ggml::backend::igpu {
             // - cas RAM/CPU
             //m_data = new (std::align_val_t(32)) uint8_t[m_size];
             //GGML_ASSERT(m_data);
-            // - cas HHIP/IGPU
+            // - cas HIP/IGPU
             m_host_data = ggml::hip::allocateHost<uint8_t>(m_size);
             m_device_data = ggml::hip::getDeviceMem(m_host_data);
         }
@@ -373,7 +384,7 @@ namespace ggml::backend::igpu {
 
         // TODO: retourner une REF !
         bool caps_host_buffer() override { return true; }
-        buffer_type* get_host_buffer_type() override {
+        ggml::cpp::backend::buffer_type* get_host_buffer_type() override {
             // IGPU_TRACE(" ####################  device[" << m_name << "] get_host_buffer_type!");
             return m_host_buffer_type;
         }
@@ -384,6 +395,21 @@ namespace ggml::backend::igpu {
         // std::vector<buffer_type> get_extra_bufts() override ; ???
 
         bool supports_op(const ggml_tensor & op) override {
+            // histoire de lister toutes les OPs...
+            static std::unordered_set<std::string> list_ops;
+            if (list_ops.count(op.name) == 0) {
+                list_ops.insert(op.name);
+                IGPU_DEV("##>> op("<< op.name<<"<"<<ggml_op_name(op.op) <<">) : "
+                        << ggml_type_name(op.type)<< "[" <<op.ne[0]<<", "<<op.ne[1]<<", "<<op.ne[2]<<", "<<op.ne[3]<<"]");
+                for (int i=0; i<GGML_MAX_SRC; ++i) {
+                    if (op.src[i] != nullptr) {
+                        IGPU_DEV("   {"<<i<<"} " << op.src[i]->name << "<" <<ggml_type_name(op.src[i]->type)<<"> "
+                                << "["<<op.src[i]->ne[0]<<", "<<op.src[i]->ne[1]<<", "<<op.src[i]->ne[2]<<", "<<op.src[i]->ne[3]<<"]");
+                    }
+                }
+            }
+
+
             switch (op.op) {
             case GGML_OP_NONE:
             case GGML_OP_RESHAPE:

ggml/src/ggml-igpu/ggml-igpu.cpp

@@ -215,6 +215,9 @@ namespace ggml::backend::igpu {
                     }
                 }
             }
+
+            __syncthreads(); // OK il faut attendre que tout soit reformaté.
+
         }
 
         // save C (sauf fin)