whoops

hercules_cg/src/gpu.rs

@@ -1879,8 +1879,10 @@ namespace cg = cooperative_groups;
                     let element_size = self.get_size(*element_type, None);
                     let element_align = self.get_alignment(*element_type);
                     index_ptr.push_str(&format!(
-                        " * (({} + {} - 1 / {} * {}))",
-                        element_size, element_align, element_align, element_align
+                        " * (({} + {}) & ~{})",
+                        element_size,
+                        element_align - 1,
+                        element_align - 1
                     ));
                     type_id = *element_type;
                 }
@@ -1932,8 +1934,10 @@ namespace cg = cooperative_groups;
                     let alignment = self.get_alignment(*type_id);
                     let size = self.get_size(*type_id, None);
                     *dynamic_shared_offset = format!(
-                        "(({} + {} - 1) / {}) * {}",
-                        dynamic_shared_offset, alignment, alignment, alignment
+                        "(({} + {}) & ~{})",
+                        dynamic_shared_offset,
+                        alignment - 1,
+                        alignment - 1,
                     );
                     write!(
                         w,
@@ -1981,8 +1985,10 @@ namespace cg = cooperative_groups;
                     let alignment = self.get_alignment(*type_id);
                     let size = self.get_size(*type_id, None);
                     *dynamic_shared_offset = format!(
-                        "(({} + {} - 1) / {}) * {}",
-                        dynamic_shared_offset, alignment, alignment, alignment
+                        "(({} + {}) & ~{})",
+                        dynamic_shared_offset,
+                        alignment - 1,
+                        alignment - 1
                     );
                     write!(
                         w,
@@ -2023,8 +2029,10 @@ namespace cg = cooperative_groups;
                 let alignment = self.get_alignment(*type_id);
                 let size = self.get_size(*type_id, None);
                 *dynamic_shared_offset = format!(
-                    "(({} + {} - 1) / {}) * {}",
-                    dynamic_shared_offset, alignment, alignment, alignment
+                    "(({} + {}) & ~{})",
+                    dynamic_shared_offset,
+                    alignment - 1,
+                    alignment - 1,
                 );
                 write!(
                     w,
@@ -2055,11 +2063,10 @@ namespace cg = cooperative_groups;
                 let array_size = multiply_dcs(extents);
                 let elem_align = self.get_alignment(type_id);
                 format!(
-                    "(({} + {} - 1) / {} * {}) * {}",
+                    "(({} + {}) & ~{}) * {}",
                     self.get_size(*element_type, None),
-                    elem_align,
-                    elem_align,
-                    elem_align,
+                    elem_align - 1,
+                    elem_align - 1,
                     array_size
                 )
             }
@@ -2073,10 +2080,7 @@ namespace cg = cooperative_groups;
                         if acc == "0" {
                             size
                         } else {
-                            format!(
-                                "({} + {} - 1) / {} * {} + {}",
-                                acc, align, align, align, size
-                            )
+                            format!("(({} + {}) & ~{}) + {}", acc, align - 1, align - 1, size)
                         }
                     })
             }
@@ -2101,8 +2105,10 @@ namespace cg = cooperative_groups;
                     .max()
                     .unwrap_or(0);
                 format!(
-                    "({} + {} - 1) / {} * {}",
-                    max_size, max_alignment, max_alignment, max_alignment
+                    "(({} + {}) & ~{})",
+                    max_size,
+                    max_alignment - 1,
+                    max_alignment - 1,
                 )
             }
             _ => {