diff --git a/README.md b/README.md
index d1b4947..d1388ab 100644
--- a/README.md
+++ b/README.md
@@ -1 +1,3 @@
-# Blocks
\ No newline at end of file
+# Blocks
+
+This is a collection of manipulation of concetp of computer engineering.
diff --git a/artificial_intelligence/high_performance_operators/conv.cpp b/artificial_intelligence/high_performance_operators/conv.cpp
new file mode 100644
index 0000000..137df01
--- /dev/null
+++ b/artificial_intelligence/high_performance_operators/conv.cpp
@@ -0,0 +1,28 @@
+#include <iostream>
+#include <vector>
+
+std::vector<std::vector<float> >& convolution_2d(std::vector<std::vector<float> > input, std::vector<std::vector<float> > weight) {
+    int h_input = input.size();
+    int w_input = input[0].size();
+    int h_weight = weight.size();
+    int w_weight = weight[0].size();
+    int w_pad = (w_weight / 2);
+    int h_pad = (h_weight / 2)
+    int h_output = h_input - 2 * h_pad;
+    int w_output = w_input - 2 * w_pad;
+
+    std::vector<std::vector<float> > output(h_output, std::vector<float>(w_output, 0));
+    for (int i(0); i < h_output; ++i) {
+        for (int j(0); j < w_output; ++j) {
+            // filter
+            for (int wi(0); wi < h_weight; ++wi) {
+                for (int wj(0); wj < w_weight; ++wj) {
+                    output[i][j] += weight[wi][wj] * input[wi + i][wj + j];
+                }
+            }
+        }
+    }
+}
+
+int main() {
+}
diff --git a/artificial_intelligence/high_performance_operators/gemm/a.out b/artificial_intelligence/high_performance_operators/gemm/a.out
new file mode 100755
index 0000000..0fb089b
Binary files /dev/null and b/artificial_intelligence/high_performance_operators/gemm/a.out differ
diff --git a/artificial_intelligence/high_performance_operators/gemm/gemm b/artificial_intelligence/high_performance_operators/gemm/gemm
new file mode 100755
index 0000000..005530d
Binary files /dev/null and b/artificial_intelligence/high_performance_operators/gemm/gemm differ
diff --git a/artificial_intelligence/high_performance_operators/gemm/gemm.cpp b/artificial_intelligence/high_performance_operators/gemm/gemm.cpp
new file mode 100644
index 0000000..3476600
--- /dev/null
+++ b/artificial_intelligence/high_performance_operators/gemm/gemm.cpp
@@ -0,0 +1,31 @@
+#include <iostream>
+#include <vector>
+#include <sycl/sycl.hpp>
+
+template<class type>
+struct matrix_2d {
+    matrix_2d(std::vector<std::vector<type>> &data) {
+        this->data = data;
+        h = data.size();
+        w = data[0].size();
+    }
+    int h;
+    int w;
+    std::vector<std::vector<type>> data;
+};
+
+template<class type>
+type & inner_product() {
+}
+
+template<class type>
+matrix_2d<type> & gemm2d(matrix_2d<type> & a, matrix_2d<type> & b, matrix_2d<type> & result) {
+
+};
+
+int main() {
+    std::vector<std::vector<float>> tmp(2, std::vector<float>{1, 1});
+    auto a = matrix_2d<float>(tmp);
+    std::cout << a.h << std::endl;
+    return 0;
+}
diff --git a/artificial_intelligence/high_performance_operators/gemm/test_lz.cpp b/artificial_intelligence/high_performance_operators/gemm/test_lz.cpp
new file mode 100644
index 0000000..0ddf2ba
--- /dev/null
+++ b/artificial_intelligence/high_performance_operators/gemm/test_lz.cpp
@@ -0,0 +1 @@
+i
diff --git a/artificial_intelligence/high_performance_operators/gemm/test_sy.cpp b/artificial_intelligence/high_performance_operators/gemm/test_sy.cpp
new file mode 100644
index 0000000..d23abb9
--- /dev/null
+++ b/artificial_intelligence/high_performance_operators/gemm/test_sy.cpp
@@ -0,0 +1,36 @@
+#include <iostream>
+#include <sycl/sycl.hpp>
+using namespace sycl; // (optional) avoids need for "sycl::" before SYCL names
+
+int main() {
+  int data[1024]; // Allocate data to be worked on
+
+  // Create a default queue to enqueue work to the default device
+  queue myQueue;
+
+  // By wrapping all the SYCL work in a {} block, we ensure
+  // all SYCL tasks must complete before exiting the block,
+  // because the destructor of resultBuf will wait
+  {
+    // Wrap our data variable in a buffer
+    buffer<int, 1> resultBuf { data, range<1> { 1024 } };
+
+    // Create a command group to issue commands to the queue
+    myQueue.submit([&](handler& cgh) {
+      // Request write access to the buffer without initialization
+      accessor writeResult { resultBuf, cgh, write_only, no_init };
+
+      // Enqueue a parallel_for task with 1024 work-items
+      cgh.parallel_for(1024, [=](id<1> idx) {
+        // Initialize each buffer element with its own rank number starting at 0
+        writeResult[idx] = idx;
+      }); // End of the kernel function
+    });   // End of our commands for this queue
+  }       // End of scope, so we wait for work producing resultBuf to complete
+
+  // Print result
+  for (int i = 0; i < 1024; i++)
+    std::cout << "data[" << i << "] = " << data[i] << std::endl;
+
+  return 0;
+}
diff --git a/artificial_intelligence/high_performance_operators/im2col.cpp b/artificial_intelligence/high_performance_operators/im2col.cpp
new file mode 100644
index 0000000..e69de29
diff --git a/artificial_intelligence/high_performance_operators/softmax.cpp b/artificial_intelligence/high_performance_operators/softmax.cpp
new file mode 100644
index 0000000..e69de29
diff --git a/artificial_intelligence/high_performance_operators/t_low.cpp b/artificial_intelligence/high_performance_operators/t_low.cpp
new file mode 100644
index 0000000..e69de29
diff --git a/artificial_intelligence/high_performance_operators/transform.cpp b/artificial_intelligence/high_performance_operators/transform.cpp
new file mode 100644
index 0000000..e69de29