[skip ci] enforce code format

taichi-dev · Mar 24, 2020 · 38127d0 · 38127d0
1 parent 2352670
commit 38127d0
Showing 1 changed file with 76 additions and 62 deletions.
diff --git a/taichi/platform/metal/shaders/runtime_utils.metal.h b/taichi/platform/metal/shaders/runtime_utils.metal.h
@@ -29,15 +29,13 @@
 #endif  // TI_INSIDE_METAL_CODEGEN
 
 METAL_BEGIN_RUNTIME_UTILS_DEF
-STR(
-    [[maybe_unused]] int num_active(device const ListManager *list) {
-      return list->next;
-    }
+STR([[maybe_unused]] int num_active(device const ListManager *list) {
+  return list->next;
+}
 
-    template <typename T>
-    int append(device ListManager *list,
-               thread const T &elem,
-               device byte *data_addr) {
+    template <typename T> int append(device ListManager *list,
+                                     thread const T &elem,
+                                     device byte *data_addr) {
       thread char *elem_ptr = (thread char *)(&elem);
       int me = atomic_fetch_add_explicit(
           reinterpret_cast<device atomic_int *>(&(list->next)), 1,
@@ -58,60 +56,76 @@ STR(
                                                  (i * list->element_stride));
     }
 
-    [[maybe_unused]] void clear(device ListManager *list) {
-      atomic_store_explicit(
-          reinterpret_cast<device atomic_int *>(&(list->next)), 0,
-          metal::memory_order_relaxed);
-    }
-
-    [[maybe_unused]] int is_active(device byte *addr, SNodeMeta meta, int i) {
-      if (meta.type == SNodeMeta::Root || meta.type == SNodeMeta::Dense) {
-        return true;
-      }
-      device auto *ptr =
-          reinterpret_cast<device atomic_uint *>(
-              addr + ((meta.num_slots - i) * meta.element_stride)) +
-          (i / (sizeof(uint32_t) * 8));
-      uint32_t bits = atomic_load_explicit(ptr, metal::memory_order_relaxed);
-      return ((bits >> (i % (sizeof(uint32_t) * 8))) & 1);
-    }
-
-    [[maybe_unused]] void activate(device byte *addr, SNodeMeta meta, int i) {
-      if (meta.type == SNodeMeta::Root || meta.type == SNodeMeta::Dense) {
-        return;
-      }
-      device auto *ptr =
-          reinterpret_cast<device atomic_uint *>(
-              addr + ((meta.num_slots - i) * meta.element_stride)) +
-          (i / (sizeof(uint32_t) * 8));
-      const uint32_t mask = (1 << (i % (sizeof(uint32_t) * 8)));
-      atomic_fetch_or_explicit(ptr, mask, metal::memory_order_relaxed);
-    }
-
-    [[maybe_unused]] void deactivate(device byte *addr, SNodeMeta meta, int i) {
-      if (meta.type == SNodeMeta::Root || meta.type == SNodeMeta::Dense) {
-        return;
-      }
-      device auto *ptr =
-          reinterpret_cast<device atomic_uint *>(
-              addr + ((meta.num_slots - i) * meta.element_stride)) +
-          (i / (sizeof(uint32_t) * 8));
-      const uint32_t mask = ~(1 << (i % (sizeof(uint32_t) * 8)));
-      atomic_fetch_and_explicit(ptr, mask, metal::memory_order_relaxed);
-    }
-
-    [[maybe_unused]] void refine_coordinates(
-        thread const ListgenElement &parent_elem,
-        device const SNodeExtractors &child_extrators,
-        int l,
-        thread ListgenElement *child_elem) {
-      for (int i = 0; i < kTaichiMaxNumIndices; ++i) {
-        device const auto &ex = child_extrators.extractors[i];
-        const int mask = ((1 << ex.num_bits) - 1);
-        const int addition = (((l >> ex.acc_offset) & mask) << ex.start);
-        child_elem->coords[i] = (parent_elem.coords[i] | addition);
-      }
-    })
+        [[maybe_unused]] void clear(device ListManager *list) {
+          atomic_store_explicit(
+              reinterpret_cast<device atomic_int *>(&(list->next)), 0,
+              metal::memory_order_relaxed);
+        }
+
+            [[maybe_unused]] int is_active(device byte *addr,
+                                           SNodeMeta meta,
+                                           int i) {
+              if (meta.type == SNodeMeta::Root ||
+                  meta.type == SNodeMeta::Dense) {
+                return true;
+              }
+              device auto *ptr =
+                  reinterpret_cast<device atomic_uint *>(
+                      addr + ((meta.num_slots - i) * meta.element_stride)) +
+                  (i / (sizeof(uint32_t) * 8));
+              uint32_t bits =
+                  atomic_load_explicit(ptr, metal::memory_order_relaxed);
+              return ((bits >> (i % (sizeof(uint32_t) * 8))) & 1);
+            }
+
+                [[maybe_unused]] void activate(device byte *addr,
+                                               SNodeMeta meta,
+                                               int i) {
+                  if (meta.type == SNodeMeta::Root ||
+                      meta.type == SNodeMeta::Dense) {
+                    return;
+                  }
+                  device auto *ptr =
+                      reinterpret_cast<device atomic_uint *>(
+                          addr + ((meta.num_slots - i) * meta.element_stride)) +
+                      (i / (sizeof(uint32_t) * 8));
+                  const uint32_t mask = (1 << (i % (sizeof(uint32_t) * 8)));
+                  atomic_fetch_or_explicit(ptr, mask,
+                                           metal::memory_order_relaxed);
+                }
+
+                    [[maybe_unused]] void deactivate(device byte *addr,
+                                                     SNodeMeta meta,
+                                                     int i) {
+                      if (meta.type == SNodeMeta::Root ||
+                          meta.type == SNodeMeta::Dense) {
+                        return;
+                      }
+                      device auto *ptr = reinterpret_cast<device atomic_uint *>(
+                                             addr + ((meta.num_slots - i) *
+                                                     meta.element_stride)) +
+                                         (i / (sizeof(uint32_t) * 8));
+                      const uint32_t mask =
+                          ~(1 << (i % (sizeof(uint32_t) * 8)));
+                      atomic_fetch_and_explicit(ptr, mask,
+                                                metal::memory_order_relaxed);
+                    }
+
+                        [[maybe_unused]] void refine_coordinates(
+                            thread const ListgenElement &parent_elem,
+                            device const SNodeExtractors &child_extrators,
+                            int l,
+                            thread ListgenElement *child_elem) {
+                          for (int i = 0; i < kTaichiMaxNumIndices; ++i) {
+                            device const auto &ex =
+                                child_extrators.extractors[i];
+                            const int mask = ((1 << ex.num_bits) - 1);
+                            const int addition =
+                                (((l >> ex.acc_offset) & mask) << ex.start);
+                            child_elem->coords[i] =
+                                (parent_elem.coords[i] | addition);
+                          }
+                        })
 METAL_END_RUNTIME_UTILS_DEF
 
 #undef METAL_BEGIN_RUNTIME_UTILS_DEF