[GUI] [mac] Support fast_gui on macOS (#1981)

* [GUI] [mac] Support fast_gui on macOS

* address comments

python/taichi/lang/meta.py

@@ -19,10 +19,20 @@ def tensor_to_ext_arr(tensor: ti.template(), arr: ti.ext_arr()):
 def vector_to_fast_image(img: ti.template(), out: ti.ext_arr()):
     # FIXME: Why is ``for i, j in img:`` slower than:
     for i, j in ti.ndrange(*img.shape):
-        u, v, w = min(255, max(0, int(img[i, img.shape[1] - 1 - j] * 255)))
+        r, g, b = min(255, max(0, int(img[i, img.shape[1] - 1 - j] * 255)))
+        idx = j * img.shape[0] + i
         # We use i32 for |out| since OpenGL and Metal doesn't support u8 types
-        # TODO: treat Cocoa and Big-endian machines, with XOR logic
-        out[j * img.shape[0] + i] = w + (v << 8) + (u << 16)
+        if ti.static(ti.get_os_name() != 'osx'):
+            out[idx] = (r << 16) + (g << 8) + b
+        else:
+            # What's -16777216?
+            #
+            # On Mac, we need to set the alpha channel to 0xff. Since Mac's GUI
+            # is big-endian, the color is stored in ABGR order, and we need to
+            # add 0xff000000, which is -16777216 in I32's legit range. (Albeit
+            # the clarity, adding 0xff000000 doesn't work.)
+            alpha = -16777216
+            out[idx] = (b << 16) + (g << 8) + r + alpha
 
 
 @ti.kernel

taichi/gui/cocoa.cpp

@@ -194,21 +194,27 @@ void updateLayer(id self, SEL _) {
   using namespace taichi;
   auto *gui = gui_from_id[self];
   auto width = gui->width, height = gui->height;
-  auto &img = gui->canvas->img;
-  auto &data = gui->img_data;
-  for (int j = 0; j < height; j++) {
-    for (int i = 0; i < width; i++) {
-      int index = 4 * (i + j * width);
-      auto pixel = img[i][height - j - 1];
-      data[index++] = uint8(clamp(int(pixel[0] * 255.0_f), 0, 255));
-      data[index++] = uint8(clamp(int(pixel[1] * 255.0_f), 0, 255));
-      data[index++] = uint8(clamp(int(pixel[2] * 255.0_f), 0, 255));
-      data[index++] = 255;  // alpha
+  uint8_t *data_ptr = nullptr;
+  if (gui->fast_gui) {
+    data_ptr = reinterpret_cast<uint8_t *>(gui->fast_buf);
+  } else {
+    auto &img = gui->canvas->img;
+    auto &data = gui->img_data;
+    data_ptr = data.data();
+    for (int j = 0; j < height; j++) {
+      for (int i = 0; i < width; i++) {
+        int index = 4 * (i + j * width);
+        auto pixel = img[i][height - j - 1];
+        data[index++] = uint8(clamp(int(pixel[0] * 255.0_f), 0, 255));
+        data[index++] = uint8(clamp(int(pixel[1] * 255.0_f), 0, 255));
+        data[index++] = uint8(clamp(int(pixel[2] * 255.0_f), 0, 255));
+        data[index++] = 255;  // alpha
+      }
     }
   }
 
   CGDataProviderRef provider = CGDataProviderCreateWithData(
-      nullptr, data.data(), gui->img_data_length, nullptr);
+      nullptr, data_ptr, gui->img_data_length, nullptr);
   CGColorSpaceRef colorspace = CGColorSpaceCreateDeviceRGB();
   CGImageRef image =
       CGImageCreate(width, height, 8, 32, width * 4, colorspace,