apply pr comments

Author: Lyxot

mlx/backend/cuda/device/hadamard.cuh

@@ -10,7 +10,9 @@ __device__ __forceinline__ void hadamard_radix_m(float* x);
 
 template <int N>
 struct Pow2Log2 {
-  static_assert(N > 0 && (N % 2 == 0), "N must be a power-of-two > 1.");
+  static_assert(
+      (N > 0) && ((N & (N - 1)) == 0),
+      "N must be a positive power of two.");
   static constexpr int value = 1 + Pow2Log2<N / 2>::value;
 };
 

mlx/backend/cuda/hadamard.cu

@@ -159,8 +159,8 @@ void hadamard_mn_contiguous(
     args.append(num_transforms);
 
     auto kernel = mod.get_kernel(n1_kernel_name);
-    encoder.add_kernel_node(
-        kernel, num_blocks, n1 / max_radix_1, 0, args.args());
+    encoder.add_kernel_node_raw(
+        kernel, num_blocks, n1 / max_radix_1, {}, 0, args.args());
   }
 
   {
@@ -179,8 +179,8 @@ void hadamard_mn_contiguous(
     args.append(num_transforms);
 
     auto kernel = mod.get_kernel(n2_kernel_name);
-    encoder.add_kernel_node(
-        kernel, num_blocks, n2 / max_radix_2, 0, args.args());
+    encoder.add_kernel_node_raw(
+        kernel, num_blocks, n2 / max_radix_2, {}, 0, args.args());
   }
 
   if (m > 1) {
@@ -200,7 +200,8 @@ void hadamard_mn_contiguous(
     args.append(num_tasks);
 
     auto kernel = mod.get_kernel(m_kernel_name);
-    encoder.add_kernel_node(kernel, num_blocks, block_dim, 0, args.args());
+    encoder.add_kernel_node_raw(
+        kernel, num_blocks, block_dim, {}, 0, args.args());
   }
 }