Fix alignment in dyn_smem example

NVIDIA · Sep 25, 2021 · e2240f6 · e2240f6
1 parent 1bc8794
commit e2240f6
Showing 1 changed file with 7 additions and 1 deletion.
diff --git a/examples/block/example_block_reduce_dyn_smem.cu b/examples/block/example_block_reduce_dyn_smem.cu
@@ -98,8 +98,14 @@ __global__ void BlockReduceKernel(
     using BlockReduceT = cub::BlockReduce<int, BLOCK_THREADS>;
     using TempStorageT = typename BlockReduceT::TempStorage;
 
+    union ShmemLayout
+    {
+      TempStorageT reduce;
+      int aggregate;
+    };
+
     // shared memory byte-array
-    extern __shared__ char smem[];
+    extern __shared__ __align__(alignof(ShmemLayout)) char smem[];
 
     // cast to lvalue reference of expected type
     auto& temp_storage = reinterpret_cast<TempStorageT&>(smem);