Skip to content

Commit d4ea720

Browse files
committed
Remove initialization from aligned_array
1 parent 71a4184 commit d4ea720

File tree

3 files changed

+30
-8
lines changed

3 files changed

+30
-8
lines changed

include/kernel_float/base.h

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -22,13 +22,15 @@ struct alignas(Alignment) aligned_array {
2222
return items_;
2323
}
2424

25-
T items_[N] = {};
25+
T items_[N];
2626
};
2727

2828
template<typename T, size_t Alignment>
2929
struct alignas(Alignment) aligned_array<T, 1, Alignment> {
30+
aligned_array() = default;
31+
3032
KERNEL_FLOAT_INLINE
31-
aligned_array(T item = {}) : item_(item) {}
33+
aligned_array(T item) : item_(item) {}
3234

3335
KERNEL_FLOAT_INLINE
3436
operator T() const {
@@ -45,7 +47,7 @@ struct alignas(Alignment) aligned_array<T, 1, Alignment> {
4547
return &item_;
4648
}
4749

48-
T item_ = {};
50+
T item_;
4951
};
5052

5153
template<typename T, size_t Alignment>

single_include/kernel_float.h

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -16,8 +16,8 @@
1616

1717
//================================================================================
1818
// this file has been auto-generated, do not modify its contents!
19-
// date: 2026-03-02 16:31:41.575770
20-
// git hash: a9bc5a6b4eaaa934175759ad44eb4b3b89f7ded2
19+
// date: 2026-03-11 14:54:13.201632
20+
// git hash: 82ce926cbd4572184b5e19b456eda9c03dba1064
2121
//================================================================================
2222

2323
#ifndef KERNEL_FLOAT_MACROS_H
@@ -424,13 +424,15 @@ struct alignas(Alignment) aligned_array {
424424
return items_;
425425
}
426426

427-
T items_[N] = {};
427+
T items_[N];
428428
};
429429

430430
template<typename T, size_t Alignment>
431431
struct alignas(Alignment) aligned_array<T, 1, Alignment> {
432+
aligned_array() = default;
433+
432434
KERNEL_FLOAT_INLINE
433-
aligned_array(T item = {}) : item_(item) {}
435+
aligned_array(T item) : item_(item) {}
434436

435437
KERNEL_FLOAT_INLINE
436438
operator T() const {
@@ -447,7 +449,7 @@ struct alignas(Alignment) aligned_array<T, 1, Alignment> {
447449
return &item_;
448450
}
449451

450-
T item_ = {};
452+
T item_;
451453
};
452454

453455
template<typename T, size_t Alignment>

tests/basics.cu

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -146,3 +146,21 @@ struct creation_tests {
146146
};
147147

148148
REGISTER_TEST_CASE("into_vec and make_vec", creation_tests, int, float)
149+
150+
struct shared_memory_tests {
151+
template<typename T, size_t... I, size_t N = sizeof...(I)>
152+
__device__ void operator()(generator<T> gen, std::index_sequence<I...>) {
153+
// It should be possible to use vector_storage for shared memory.
154+
// Using the regular vector type is not possible as it needs initialization.
155+
__shared__ kernel_float::vector_storage<T, N> items;
156+
}
157+
};
158+
159+
REGISTER_TEST_CASE_GPU(
160+
"shared memory",
161+
shared_memory_tests,
162+
int,
163+
float,
164+
double,
165+
__half,
166+
__nv_bfloat16)

0 commit comments

Comments
 (0)