From 38127d0daf2a1c997f4006dc740dbf48c6e474ab Mon Sep 17 00:00:00 2001 From: Taichi Gardener Date: Tue, 24 Mar 2020 17:51:29 -0400 Subject: [PATCH] [skip ci] enforce code format --- .../metal/shaders/runtime_utils.metal.h | 138 ++++++++++-------- 1 file changed, 76 insertions(+), 62 deletions(-) diff --git a/taichi/platform/metal/shaders/runtime_utils.metal.h b/taichi/platform/metal/shaders/runtime_utils.metal.h index 3c719b0961c238..f19956b195ae3a 100644 --- 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 - int append(device ListManager *list, - thread const T &elem, - device byte *data_addr) { + template 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(&(list->next)), 1, @@ -58,60 +56,76 @@ STR( (i * list->element_stride)); } - [[maybe_unused]] void clear(device ListManager *list) { - atomic_store_explicit( - reinterpret_cast(&(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( - 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( - 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( - 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(&(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( + 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( + 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( + 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