fix mac os and CI

FindDefinition · Dec 5, 2024 · 64798e8 · 64798e8
1 parent 09b759e
commit 64798e8
Show file tree

Hide file tree

Showing 3 changed files with 7 additions and 6 deletions.
diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml
@@ -164,7 +164,7 @@ jobs:
       - name: Store the binary wheel
         uses: actions/upload-artifact@v4
         with:
-          name: python-package-distributions
+          name: python-package-distributions-${{ matrix.python-version }}
           path: dist
 
   deploy-macos:
@@ -176,7 +176,7 @@ jobs:
       - name: Download all the dists
         uses: actions/download-artifact@v4
         with:
-          name: python-package-distributions
+          name: python-package-distributions-*
           path: dist/
       - name: Publish a Python distribution to PyPI
         if: github.event_name == 'push' && startsWith(github.ref, 'refs/tags')

diff --git a/include/tensorview/core/printf2.h b/include/tensorview/core/printf2.h
@@ -17,6 +17,7 @@
 #include "defs.h"
 #ifndef TV_PARALLEL_RTC
 #include <cstdio>
+#include <cstdint>
 #endif
 #include "const_string.h"
 namespace tv {

diff --git a/include/tensorview/parallel/ops.h b/include/tensorview/parallel/ops.h
@@ -62,7 +62,7 @@ uint __apple_metal_warp_index [[simdgroup_index_in_threadgroup]];
 uint __apple_metal_lane_index [[thread_index_in_simdgroup]];
 
 uint3 __apple_metal_block_indexes [[threadgroup_position_in_grid]];
-uint3 __apple_metal_block_dims [[threadgroup_size_in_grid]];
+uint3 __apple_metal_block_dims [[threadgroups_per_grid]];
 uint3 __apple_metal_grid_dims [[grid_size]];
 uint3 __apple_metal_thread_indexes [[thread_position_in_threadgroup]];
 
@@ -181,7 +181,7 @@ template <typename T> TV_DEVICE_INLINE T atomicAdd(TV_METAL_DEVICE T *ctr, T val
 
 template <typename T, size_t N> TV_DEVICE_INLINE array<T, N> atomicAdd(TV_METAL_DEVICE array<T, N> *ctr, const TV_METAL_THREAD array<T, N>& val) {
   auto ptr_array = arrayops::create_ptr_arange<N>(reinterpret_cast<TV_METAL_DEVICE T*>(ctr));
-  return arrayops::apply(atomicAdd<float>, reinterpret_cast<tv::array<T*, N>&>(ptr_array), val);
+  return arrayops::apply(atomicAdd<float>, reinterpret_cast<TV_METAL_DEVICE tv::array<TV_METAL_DEVICE T*, N>&>(ptr_array), val);
 }
 
 template <typename T> TV_DEVICE_INLINE T atomicMax(TV_METAL_DEVICE T *ctr, T val) {
@@ -259,12 +259,12 @@ TV_DEVICE_INLINE float atomicMin (device float * addr, float value) {
 
 template <typename T, size_t N> TV_DEVICE_INLINE array<T, N> atomicMax(TV_METAL_DEVICE array<T, N> *ctr, const TV_METAL_THREAD array<T, N>& val) {
   auto ptr_array = arrayops::create_ptr_arange<N>(reinterpret_cast<TV_METAL_DEVICE T*>(ctr));
-  return arrayops::apply(atomicMax<float>, reinterpret_cast<tv::array<T*, N>&>(ptr_array), val);
+  return arrayops::apply(atomicMax<float>, reinterpret_cast<TV_METAL_DEVICE tv::array<TV_METAL_DEVICE T*, N>&>(ptr_array), val);
 }
 
 template <typename T, size_t N> TV_DEVICE_INLINE array<T, N> atomicMin(TV_METAL_DEVICE array<T, N> *ctr, const TV_METAL_THREAD array<T, N>& val) {
   auto ptr_array = arrayops::create_ptr_arange<N>(reinterpret_cast<TV_METAL_DEVICE T*>(ctr));
-  return arrayops::apply(atomicMin<float>, reinterpret_cast<tv::array<T*, N>&>(ptr_array), val);
+  return arrayops::apply(atomicMin<float>, reinterpret_cast<TV_METAL_DEVICE tv::array<TV_METAL_DEVICE T*, N>&>(ptr_array), val);
 }
 
 #ifdef TV_METAL_RTC