结构体数组（SoA）与数组结构体（AoS）

最新推荐文章于 2025-02-24 17:02:55 发布
最新推荐文章于 2025-02-24 17:02:55 发布
阅读量8.7k 收藏 10
点赞数 1
分类专栏： 【语言与编译】 【高性能计算】
本文对比了结构体数组（SoA）与数组结构体（AoS）两种数据布局方式，并通过CUDA编程实践展示了它们在并行计算中的性能差异。文中详细介绍了两种布局在CPU与GPU上的实现方式，以及如何利用CUDA进行有效的内存访问。

摘要生成于 C知道 ，由 DeepSeek-R1 满血版支持， 前往体验 >

1.结构体数组（SoA）

/*
 * SoA 结构体数组定义
 */
struct InnerArray
{
    float x[LEN];
    float y[LEN];
};
/*
 * CPU -> SoA 结构体数组的CPU计算形式
 */
void testInnerArrayHost(InnerArray *A, InnerArray *C, const int n)
{
    for (int idx = 0; idx < n; idx++)
    {
        C->x[idx] = A->x[idx] + 10.f;
        C->y[idx] = A->y[idx] + 20.f;
    }

    return;
}
/*
 * GPU -> SoA 结构体数组的CUDA计算模式
 */
__global__ void testInnerArrayDevice(InnerArray *data, InnerArray * result,
                               const int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        float tmpx = data->x[i];
        float tmpy = data->y[i];

        tmpx += 10.f;
        tmpy += 20.f;
        result->x[i] = tmpx;
        result->y[i] = tmpy;
    }
}

2.数组结构体（AoS） 

/*
 * AoS
 */
struct innerStruct
{
    float x;
    float y;
};
/*
 * CPU -> AoS
 */
void testInnerStructHost(innerStruct *A, innerStruct *C, const int n)
{
    for (int idx = 0; idx < n; idx++)
    {
        C[idx].x = A[idx].x + 10.f;
        C[idx].y = A[idx].y + 20.f;
    }

    return;
}
/*
 * GPU -> AoS
 */
__global__ void testInnerStructDevice(innerStruct *data, innerStruct * result,
                                const int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        innerStruct tmp = data[i];
        tmp.x += 10.f;
        tmp.y += 20.f;
        result[i] = tmp;
    }
}

3.结构体数组（SoA）与数组结构体（AoS）二者的区别

许多并行编程范式，尤其是SIMD（单指令多数据）型范式，更倾向于使用SoA。在CUDA C编程中也普遍倾向于SoA，一维数据元素是为全局内存的有效合并访问而预先准备好的，而相同内存操作引用的同字段元素在存储时时彼此相邻的。


4.给出源代码示例，（《CUDA C编程中文翻译版本》，如有侵权，联系删除）源码网址：点击打开链接

SoA

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 * A simple example of using a structore of arrays to store data on the device.
 * This example is used to study the impact on performance of data layout on the
 * GPU.
 *
 * SoA: contiguous reads for x and y
 */

#define LEN 1<<22

struct InnerArray
{
    float x[LEN];
    float y[LEN];
};

// functions for inner array outer struct
void initialInnerArray(InnerArray *ip,  int size)
{
    for (int i = 0; i < size; i++)
    {
        ip->x[i] = (float)( rand() & 0xFF ) / 100.0f;
        ip->y[i] = (float)( rand() & 0xFF ) / 100.0f;
    }

    return;
}

void testInnerArrayHost(InnerArray *A, InnerArray *C, const int n)
{
    for (int idx = 0; idx < n; idx++)
    {
        C->x[idx] = A->x[idx] + 10.f;
        C->y[idx] = A->y[idx] + 20.f;
    }

    return;
}


void printfHostResult(InnerArray *C, const int n)
{
    for (int idx = 0; idx < n; idx++)
    {
        printf("printout idx %d:  x %f y %f\n", idx, C->x[idx], C->y[idx]);
    }

    return;
}

void checkInnerArray(InnerArray *hostRef, InnerArray *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef->x[i] - gpuRef->x[i]) > epsilon)
        {
            match = 0;
            printf("different on x %dth element: host %f gpu %f\n", i,
                   hostRef->x[i], gpuRef->x[i]);
            break;
        }

        if (abs(hostRef->y[i] - gpuRef->y[i]) > epsilon)
        {
            match = 0;
            printf("different on y %dth element: host %f gpu %f\n", i,
                   hostRef->y[i], gpuRef->y[i]);
            break;
        }
    }

    if (!match)  printf("Arrays do not match.\n\n");
}

__global__ void testInnerArray(InnerArray *data, InnerArray * result,
                               const int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        float tmpx = data->x[i];
        float tmpy = data->y[i];

        tmpx += 10.f;
        tmpy += 20.f;
        result->x[i] = tmpx;
        result->y[i] = tmpy;
    }
}

__global__ void warmup2(InnerArray *data, InnerArray * result, const int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        float tmpx = data->x[i];
        float tmpy = data->y[i];
        tmpx += 10.f;
        tmpy += 20.f;
        result->x[i] = tmpx;
        result->y[i] = tmpy;
    }
}

// test for array of struct
int main(int argc, char **argv)
{
    // set up device
    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("%s test struct of array at ", argv[0]);
    printf("device %d: %s \n", dev, deviceProp.name);
    CHECK(cudaSetDevice(dev));

    // allocate host memory
    int nElem = LEN;
    size_t nBytes = sizeof(InnerArray);
    InnerArray     *h_A = (InnerArray *)malloc(nBytes);
    InnerArray *hostRef = (InnerArray *)malloc(nBytes);
    InnerArray *gpuRef  = (InnerArray *)malloc(nBytes);

    // initialize host array
    initialInnerArray(h_A, nElem);
    testInnerArrayHost(h_A, hostRef, nElem);

    // allocate device memory
    InnerArray *d_A, *d_C;
    CHECK(cudaMalloc((InnerArray**)&d_A, nBytes));
    CHECK(cudaMalloc((InnerArray**)&d_C, nBytes));

    // copy data from host to device
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));

    // set up offset for summary
    int blocksize = 128;

    if (argc > 1) blocksize = atoi(argv[1]);

    // execution configuration
    dim3 block (blocksize, 1);
    dim3 grid  ((nElem + block.x - 1) / block.x, 1);

    // kernel 1:
    double iStart = seconds();
    warmup2<<<grid, block>>>(d_A, d_C, nElem);
    CHECK(cudaDeviceSynchronize());
    double iElaps = seconds() - iStart;
    printf("warmup2      <<< %3d, %3d >>> elapsed %f sec\n", grid.x, block.x,
           iElaps);
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));
    checkInnerArray(hostRef, gpuRef, nElem);
    CHECK(cudaGetLastError());

    iStart = seconds();
    testInnerArray<<<grid, block>>>(d_A, d_C, nElem);
    CHECK(cudaDeviceSynchronize());
    iElaps = seconds() - iStart;
    printf("innerarray   <<< %3d, %3d >>> elapsed %f sec\n", grid.x, block.x,
           iElaps);
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));
    checkInnerArray(hostRef, gpuRef, nElem);
    CHECK(cudaGetLastError());

    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_C));
    free(h_A);
    free(hostRef);
    free(gpuRef);

    // reset device
    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

AoS 

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 * A simple example of using an array of structures to store data on the device.
 * This example is used to study the impact on performance of data layout on the
 * GPU.
 *
 * AoS: one contiguous 64-bit read to get x and y (up to 300 cycles)
 */

#define LEN 1<<22

struct innerStruct
{
    float x;
    float y;
};

struct innerArray
{
    float x[LEN];
    float y[LEN];
};

void initialInnerStruct(innerStruct *ip,  int size)
{
    for (int i = 0; i < size; i++)
    {
        ip[i].x = (float)(rand() & 0xFF) / 100.0f;
        ip[i].y = (float)(rand() & 0xFF) / 100.0f;
    }

    return;
}

void testInnerStructHost(innerStruct *A, innerStruct *C, const int n)
{
    for (int idx = 0; idx < n; idx++)
    {
        C[idx].x = A[idx].x + 10.f;
        C[idx].y = A[idx].y + 20.f;
    }

    return;
}

void checkInnerStruct(innerStruct *hostRef, innerStruct *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i].x - gpuRef[i].x) > epsilon)
        {
            match = 0;
            printf("different on %dth element: host %f gpu %f\n", i,
                    hostRef[i].x, gpuRef[i].x);
            break;
        }

        if (abs(hostRef[i].y - gpuRef[i].y) > epsilon)
        {
            match = 0;
            printf("different on %dth element: host %f gpu %f\n", i,
                    hostRef[i].y, gpuRef[i].y);
            break;
        }
    }

    if (!match)  printf("Arrays do not match.\n\n");
}

__global__ void testInnerStruct(innerStruct *data, innerStruct * result,
                                const int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        innerStruct tmp = data[i];
        tmp.x += 10.f;
        tmp.y += 20.f;
        result[i] = tmp;
    }
}

__global__ void warmup(innerStruct *data, innerStruct * result, const int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        innerStruct tmp = data[i];
        tmp.x += 10.f;
        tmp.y += 20.f;
        result[i] = tmp;
    }
}

int main(int argc, char **argv)
{
    // set up device
    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("%s test struct of array at ", argv[0]);
    printf("device %d: %s \n", dev, deviceProp.name);
    CHECK(cudaSetDevice(dev));

    // allocate host memory
    int nElem = LEN;
    size_t nBytes = nElem * sizeof(innerStruct);
    innerStruct     *h_A = (innerStruct *)malloc(nBytes);
    innerStruct *hostRef = (innerStruct *)malloc(nBytes);
    innerStruct *gpuRef  = (innerStruct *)malloc(nBytes);

    // initialize host array
    initialInnerStruct(h_A, nElem);
    testInnerStructHost(h_A, hostRef, nElem);

    // allocate device memory
    innerStruct *d_A, *d_C;
    CHECK(cudaMalloc((innerStruct**)&d_A, nBytes));
    CHECK(cudaMalloc((innerStruct**)&d_C, nBytes));

    // copy data from host to device
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));

    // set up offset for summaryAU: It is blocksize not offset. Thanks.CZ
    int blocksize = 128;

    if (argc > 1) blocksize = atoi(argv[1]);

    // execution configuration
    dim3 block (blocksize, 1);
    dim3 grid  ((nElem + block.x - 1) / block.x, 1);

    // kernel 1: warmup
    double iStart = seconds();
    warmup<<<grid, block>>>(d_A, d_C, nElem);
    CHECK(cudaDeviceSynchronize());
    double iElaps = seconds() - iStart;
    printf("warmup      <<< %3d, %3d >>> elapsed %f sec\n", grid.x, block.x,
           iElaps);
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));
    checkInnerStruct(hostRef, gpuRef, nElem);
    CHECK(cudaGetLastError());

    // kernel 2: testInnerStruct
    iStart = seconds();
    testInnerStruct<<<grid, block>>>(d_A, d_C, nElem);
    CHECK(cudaDeviceSynchronize());
    iElaps = seconds() - iStart;
    printf("innerstruct <<< %3d, %3d >>> elapsed %f sec\n", grid.x, block.x,
           iElaps);
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));
    checkInnerStruct(hostRef, gpuRef, nElem);
    CHECK(cudaGetLastError());

    // free memories both host and device
    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_C));
    free(h_A);
    free(hostRef);
    free(gpuRef);

    // reset device
    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}


【CUDA】内存访问
GG_Bruse的博客
02-07 872
GPU对片外DRAM的访问延迟大，带宽低，导致其成为很多应用的性能瓶颈。因此对DRAM访问的进一步优化可以有效改善程序性能。优化之前，首先看一下GPU对内存的访问模式如上图所示，DRAM内存的读写在物理上是从 片外DRAM -> 片上Cache -> 寄存器。其中，片外DRAM到片上Cache是主要性能瓶颈。DRAM 到 Cache 之间的一次传输(transaction)设计为32、64或者128字节，并且内存地址按照32、64或者128字节的间隔对齐。
数组与结构体
周文慧--廊坊师范学院提高班15期
04-21 3841
数组与结构体的区别与举例： 一、定义 数组：数组是具有相同数据类型的项的有序集合，要访问数组中的某个项，需要同时使用数组名称及该项与数组起点之间的偏移量，在C#中，声明和使用数组的方法与Java有一些重要区别。 结构体：结构体的关键字 struct,struct类型是一种值类型通常用来封装小型变量组，如果冲结构创建一个对象并将对象赋给某个变量 则变量包含结构该变量放入结构的全部值，如果...
【cuda学习日记】4.3 结构体数组与数组结构体
最新发布
anbncn1234的博客
02-24 295
如图说明了AoS和SoA方法的内存布局，用AoS模式在GPU上存储示例数据并执行一个只有x字段的应用程序，将导致50%的带宽损失。4.3 数组结构体（AoS）和结构体数组（SoA）·对齐及合并内存访问，以减少带宽的浪费。·足够的并发内存操作，以隐藏内存延迟。4.3.1 简单示例AoS。4.3.1 简单示例SoA。SoA方法来存储数据。
C语言算法笔记1：结构体、结构体数组实战讲解
NULL指向我的博客
06-30 2064
本文回顾复习结构体的概念以及一些基本操作，本文回顾复习结构体的概念以及一些基本操作，每个知识点都有代码实践演示，可以复制测试查看！争取一篇文章的搜索量，就解决大家对结构体的疑问，教会大家结构体的一些操作手法！
Aos&Soa;和cacheblocking
02-13
代码优化中的两个常用技巧，对编写优质代码十分有帮助，分享给大家
结构体与数组
qq_54075859的博客
08-08 453
结构体与数组；结构体数组；结构体的成员为数组类型；
SOA AOS 面向对象与面向数据
不甘平凡的小鸟的专栏
08-22 1375
转载自https://blog.csdn.net/debugconsole/article/details/17788979 SOA 数组的结构与AOS结构的数组，是面向数据和面向对象设计的区别之一。 在需要高频率（如渲染循环中）访问数据的时候，一般情况下SOA的效率高于AOS，因为将需要频繁访问的数据连续存放会大大提高访问速度。虽然AOS的结构可能更适合面向对象设计，但是在高度依赖效率的地...
SoA vs AoS
unityofficial的博客
02-06 607
SoA vs AoS今天我们聊一点内存中数据组织的话题，题目中提到了两个名词array of structures (AoS)和structure of arrays (SoA)，这分别代表着两种不同的数组组织形式。比如下面一段代码：struct AoSData { public int a; public int b; public int c; public i...
soa2aos（IN，LCHK）：将数组的结构转换为结构的列数组-matlab开发
05-29
句法： OUT = soa2aos(IN) OUT = soa2aos(IN, LCHK) 参数： IN 标量结构，其字段是类的数组'cell'、'numeric' 或 'struct' 并且可以定向为列、行或多维。 LCHK 标量逻辑，可选输入。 true（默认）=> 每个字段数组的...
MATLAB开发：soa2aos函数实现数组结构转换详解
函数的调用格式为OUT = soa2aos(IN)，其中IN为输入的结构体数组。在某些情况下，还可以添加一个可选参数LCHK，其作用是用来校验输入结构体数组的字段长度是否一致。如果设置LCHK为true（默认值），则要求输入结构体...
优化技巧：Aos&Soa与Cacheblocking在访存优化中的应用
"本文将探讨两种重要的代码优化技术——AoS与SoA布局以及Cache blocking，并解释它们在提升代码性能方面的作用。这两种方法主要针对内存访问优化，特别是针对大规模数据处理时的效率提升。了解并合理运用这些技术，...
第8讲-结构体与结构体数组
02-15
讲解区分结构体与结构体数组：结构体的引出、结构体变量的定义和引用、结构体数组
CUDA内存访问模式
A11en3的博客
05-02 1265
大多数设备端数据访问都是从全局内存开始的，并且多数GPU应用程序容易受内存带宽的限制。因此，最大限度的利用全局内存带宽是调控核函数性能的基本。
ECS架构之内存布局
liran2019的博客
12-07 1513
ECS架构是面向数据编程的代表，面向数据编程的意义就是让我们关注数据编程，从而提高程序的性能。数据的处理一般分为数据的读取以及数据的计算。 数据读取 我们需要尽可能的使用高速缓存来提高运行效率，不仅仅针对内存，磁盘的读取在游戏开发中也是非常重要的。 数据的计算 cpu端SIMD计算，以及gpu的CUDA计算都是为了提高数据的计算能力。 无论是提高缓存的命中率还是SIMD及CUDA计算，SOA式的内存布局都要比AOS更加友好。但这也不是绝对，准确的说应该是我们尽量保证数据的局部性。 说到局部性，无
学习随笔 一 SOA AOS 面向对象与面向数据
debugconsole的专栏
01-03 5238
SOA 数组的结构与AOS结构的数组，是面向数据和面向对象设计的区别之一。 在需要高频率（如渲染循环中）访问数据的时候，一般情况下SOA的效率高于AOS，因为将需要频繁访问的数据连续存放会大大提高访问速度。虽然AOS的结构可能更适合面向对象设计，但是在高度依赖效率的地方应该使用SOA。 引用大家都在说的一句话：一开始写程序用struct (c方式的编程,面向数据的编程) 后面开始用class
MPI技术测试-结构体类型的创建（AOS&SOA组合）
qq_33562636的博客
01-25 267
结构体类型的创建（AOS&SOA组合） #include "mpi.h" #include <stdio.h> #include <stdlib.h> #define DOUBLE_NUMS 2 struct Value{ int a; double b[DOUBLE_NUMS]; }; void initArray(struct Value *array, int size,int bSize) { for (int i =
结构体数组
三三想说
08-15 256
/*  * 2.cpp  *  *  Created on: 2011-8-15  *      Author: san  */ # include # include struct person { char name[20]; int
OpenGL vao aos soa
qq_26328385的博客
12-11 477
教程链接 合理的组织数据结构，减少状态切换，可以提高渲染性能。 AOS (Array Of Structure): 把多个模型的数据，以结构体数组的形式，存储在一个 Array Buffer中。 SOA (Structure Of Arrays) ：把多个模型的数据，组合成数组组成的结构体，分别存储在相应的Buffer中。 可以用 glDrawElementsBaseVertex、glDrawElementsInstancedBaseVertex等命令进行渲染。 ...
结构体与结构体数组
m0_69135841的博客
04-10 3084
结构体跟数组的功能有些相似，都是变量的集合。不同点在于，数组中的变量的类型都是一样的，而结构体中的变量的类型是可以自己根据需求来定的。在使用方面，数组能直接在主函数中定义、访问，而结构体再做这些操作前需要在主函数之前定义出结构体的模板，之后才能对结构体进行定义、赋值、访问等等。 struct Player { char name[32]; char sex[32]; char postion[32]; float height; float weight; }; 这是我定义的一个结构体模板。
评论 4
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值