From 5cee37adcca4c3de33c025f4975d381cf113f2bf Mon Sep 17 00:00:00 2001
From: Russel Arbore <russel.jma@gmail.com>
Date: Tue, 25 Feb 2025 14:52:16 -0600
Subject: [PATCH] Revert "whoops"
This reverts commit 9a58de9e055004d05d0df66d062987c94819b790.
---
hercules_cg/src/gpu.rs | 42 ++++++++++++++++++------------------------
1 file changed, 18 insertions(+), 24 deletions(-)
diff --git a/hercules_cg/src/gpu.rs b/hercules_cg/src/gpu.rs
index d857fdda..e3dcc0c3 100644
--- a/hercules_cg/src/gpu.rs
+++ b/hercules_cg/src/gpu.rs
@@ -1879,10 +1879,8 @@ 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!(
- " * (({} + {}) & ~{})",
- element_size,
- element_align - 1,
- element_align - 1
+ " * (({} + {} - 1) / {} * {})",
+ element_size, element_align, element_align, element_align
));
type_id = *element_type;
}
@@ -1934,10 +1932,8 @@ namespace cg = cooperative_groups;
let alignment = self.get_alignment(*type_id);
let size = self.get_size(*type_id, None);
*dynamic_shared_offset = format!(
- "(({} + {}) & ~{})",
- dynamic_shared_offset,
- alignment - 1,
- alignment - 1,
+ "(({} + {} - 1) / {}) * {}",
+ dynamic_shared_offset, alignment, alignment, alignment
);
write!(
w,
@@ -1985,10 +1981,8 @@ namespace cg = cooperative_groups;
let alignment = self.get_alignment(*type_id);
let size = self.get_size(*type_id, None);
*dynamic_shared_offset = format!(
- "(({} + {}) & ~{})",
- dynamic_shared_offset,
- alignment - 1,
- alignment - 1
+ "(({} + {} - 1) / {}) * {}",
+ dynamic_shared_offset, alignment, alignment, alignment
);
write!(
w,
@@ -2029,10 +2023,8 @@ namespace cg = cooperative_groups;
let alignment = self.get_alignment(*type_id);
let size = self.get_size(*type_id, None);
*dynamic_shared_offset = format!(
- "(({} + {}) & ~{})",
- dynamic_shared_offset,
- alignment - 1,
- alignment - 1,
+ "(({} + {} - 1) / {}) * {}",
+ dynamic_shared_offset, alignment, alignment, alignment
);
write!(
w,
@@ -2063,10 +2055,11 @@ 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 - 1,
- elem_align - 1,
+ elem_align,
+ elem_align,
+ elem_align,
array_size
)
}
@@ -2080,7 +2073,10 @@ namespace cg = cooperative_groups;
if acc == "0" {
size
} else {
- format!("(({} + {}) & ~{}) + {}", acc, align - 1, align - 1, size)
+ format!(
+ "({} + {} - 1) / {} * {} + {}",
+ acc, align, align, align, size
+ )
}
})
}
@@ -2105,10 +2101,8 @@ namespace cg = cooperative_groups;
.max()
.unwrap_or(0);
format!(
- "(({} + {}) & ~{})",
- max_size,
- max_alignment - 1,
- max_alignment - 1,
+ "({} + {} - 1) / {} * {}",
+ max_size, max_alignment, max_alignment, max_alignment
)
}
_ => {
--
GitLab