小编Car*_*íos的帖子

我有两个版本的内核执行相同的任务 - 填充链接单元列表 - 两个内核之间的区别是存储粒子位置的数据类型，第一个使用浮点数组来存储位置（每个粒子有 4 个浮点数）到 128 位读/写），第二个使用 vec3f 结构数组来存储位置（一个包含 3 个浮点数的结构）。

使用 nvprof 做一些测试，我发现第二个内核（使用 vec3f）比第一个运行得更快：

 Time(%)      Time   Calls       Avg       Min       Max  Name
   42.88    37.26s       2    18.63s   23.97us    37.26s  adentu_grid_cuda_filling_kernel(int*, int*, int*, float*, int, _vec3f, _vec3f, _vec3i)
   11.00     3.93s       2     1.97s   25.00us     3.93s  adentu_grid_cuda_filling_kernel(int*, int*, int*, _vec3f*, int, _vec3f, _vec3f, _vec3i)

测试是尝试使用 256 和 512000 个粒子填充链接的单元格列表。

我的问题是，这里发生了什么？我认为 float 数组应该由于合并内存而进行更好的内存访问，而不是使用具有未对齐内存的 vec3f 结构数组。我误会了什么？

这些是内核，第一个内核：

__global__ void adentu_grid_cuda_filling_kernel (int *head,
                                                 int *linked,
                                                 int *cellnAtoms,
                                                 float *pos, 
                                                 int nAtoms, 
                                                 vec3f origin, 
                                                 vec3f h, …