Skip to content
Snippets Groups Projects
Commit 5cee37ad authored by Russel Arbore's avatar Russel Arbore
Browse files

Revert "whoops" 

This reverts commit 9a58de9e.
parent 9a58de9e
No related branches found
No related tags found
1 merge request!199Fix type layout
Pipeline #201866 failed
Stage: test
This commit is part of merge request !199. Comments created here will be created in the context of that merge request.
Hide whitespace changes
Inline Side-by-side
Showing
with 18 additions and 24 deletions
hercules_cg/src/gpu.rs
+ 18
− 24
View file @ 5cee37ad
...@@ -1879,10 +1879,8 @@ namespace cg = cooperative_groups; ...@@ -1879,10 +1879,8 @@ namespace cg = cooperative_groups;
let element_size = self.get_size(*element_type, None); let element_size = self.get_size(*element_type, None);
let element_align = self.get_alignment(*element_type); let element_align = self.get_alignment(*element_type);
index_ptr.push_str(&format!( index_ptr.push_str(&format!(
" * (({} + {}) & ~{})", " * (({} + {} - 1) / {} * {})",
element_size, element_size, element_align, element_align, element_align
element_align - 1,
element_align - 1
)); ));
type_id = *element_type; type_id = *element_type;
} }
...@@ -1934,10 +1932,8 @@ namespace cg = cooperative_groups; ...@@ -1934,10 +1932,8 @@ namespace cg = cooperative_groups;
let alignment = self.get_alignment(*type_id); let alignment = self.get_alignment(*type_id);
let size = self.get_size(*type_id, None); let size = self.get_size(*type_id, None);
*dynamic_shared_offset = format!( *dynamic_shared_offset = format!(
"(({} + {}) & ~{})", "(({} + {} - 1) / {}) * {}",
dynamic_shared_offset, dynamic_shared_offset, alignment, alignment, alignment
alignment - 1,
alignment - 1,
); );
write!( write!(
w, w,
...@@ -1985,10 +1981,8 @@ namespace cg = cooperative_groups; ...@@ -1985,10 +1981,8 @@ namespace cg = cooperative_groups;
let alignment = self.get_alignment(*type_id); let alignment = self.get_alignment(*type_id);
let size = self.get_size(*type_id, None); let size = self.get_size(*type_id, None);
*dynamic_shared_offset = format!( *dynamic_shared_offset = format!(
"(({} + {}) & ~{})", "(({} + {} - 1) / {}) * {}",
dynamic_shared_offset, dynamic_shared_offset, alignment, alignment, alignment
alignment - 1,
alignment - 1
); );
write!( write!(
w, w,
...@@ -2029,10 +2023,8 @@ namespace cg = cooperative_groups; ...@@ -2029,10 +2023,8 @@ namespace cg = cooperative_groups;
let alignment = self.get_alignment(*type_id); let alignment = self.get_alignment(*type_id);
let size = self.get_size(*type_id, None); let size = self.get_size(*type_id, None);
*dynamic_shared_offset = format!( *dynamic_shared_offset = format!(
"(({} + {}) & ~{})", "(({} + {} - 1) / {}) * {}",
dynamic_shared_offset, dynamic_shared_offset, alignment, alignment, alignment
alignment - 1,
alignment - 1,
); );
write!( write!(
w, w,
...@@ -2063,10 +2055,11 @@ namespace cg = cooperative_groups; ...@@ -2063,10 +2055,11 @@ namespace cg = cooperative_groups;
let array_size = multiply_dcs(extents); let array_size = multiply_dcs(extents);
let elem_align = self.get_alignment(type_id); let elem_align = self.get_alignment(type_id);
format!( format!(
"(({} + {}) & ~{}) * {}", "(({} + {} - 1) / {} * {}) * {}",
self.get_size(*element_type, None), self.get_size(*element_type, None),
elem_align - 1, elem_align,
elem_align - 1, elem_align,
elem_align,
array_size array_size
) )
} }
...@@ -2080,7 +2073,10 @@ namespace cg = cooperative_groups; ...@@ -2080,7 +2073,10 @@ namespace cg = cooperative_groups;
if acc == "0" { if acc == "0" {
size size
} else { } else {
format!("(({} + {}) & ~{}) + {}", acc, align - 1, align - 1, size) format!(
"({} + {} - 1) / {} * {} + {}",
acc, align, align, align, size
)
} }
}) })
} }
...@@ -2105,10 +2101,8 @@ namespace cg = cooperative_groups; ...@@ -2105,10 +2101,8 @@ namespace cg = cooperative_groups;
.max() .max()
.unwrap_or(0); .unwrap_or(0);
format!( format!(
"(({} + {}) & ~{})", "({} + {} - 1) / {} * {}",
max_size, max_size, max_alignment, max_alignment, max_alignment
max_alignment - 1,
max_alignment - 1,
) )
} }
_ => { _ => {
......
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment