pipline_gemm报错CUDA: misaligned address

[tangpanyu]

========= Invalid __shared__ read of size 16 bytes
=========     at matrixMul(const float *, const float *, float *, int, int, int, float, float)+0x7120 in /home/tpy/gitcode/Hands-on-GEMM/src/cuda/pipeline_gemm.cu:94
=========     by thread (0,4,0) in block (0,0,0)
=========     Address 0x10a8 is misaligned
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x23b3b7]
=========                in /usr/lib/wsl/drivers/nv_dispig.inf_amd64_0afec3f2050014a0/libcuda.so.1.1
=========     Host Frame:libcudart_static_4d8b33a106dceb3c07a56e26de61f2d53bb62a68 [0x12a42]
=========                in /home/tpy/gitcode/Hands-on-GEMM/./bin/benchmark_pipeline
=========     Host Frame:cudaLaunchKernel [0x766ad]
=========                in /home/tpy/gitcode/Hands-on-GEMM/./bin/benchmark_pipeline
=========     Host Frame:sgemm(int, int, int, float*, float*, float*, float, float) [0xd05a]
=========                in /home/tpy/gitcode/Hands-on-GEMM/./bin/benchmark_pipeline
=========     Host Frame:main [0xb1c7]
=========                in /home/tpy/gitcode/Hands-on-GEMM/./bin/benchmark_pipeline
=========     Host Frame: [0x29d8f]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:__libc_start_main [0x29e3f]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0xbe24]
=========                in /home/tpy/gitcode/Hands-on-GEMM/./bin/benchmark_pipeline
........

[tangpanyu]

__shared__ cuda::pipeline_shared_state<cuda::thread_scope::thread_scope_block, stage_count> shared_state;

这一行放在下面，但是计算结果错误

[tangpanyu]

__shared__ cuda::pipeline_shared_state<cuda::thread_scope::thread_scope_block, stage_count> shared_state;

这一行放在下面，但是计算结果错误