Introduction

Many kernels are "memory-bound", which means that the kernel's runtime is determined by how long it takes to move data to and from the SM. Changing the way data is laid out in memory can significantly affect a kernel's effective memory throughput, which results in a noticeable impact on the runtime.

So, a natural question people ask when writing CUDA code is: how should I lay out data in memory to work best on the GPU?

If you search the internet for that question, you find a lot of papers and forum discussions (and even AI overviews) that make the same claim: "GPUs work best when you store data in a Struct-of-Array (SoA) layout".

In this document, we'll investigate whether that recommendation is actually true. We will consider 3 of the most common data access patterns found in kernels, and compare the effectiveness of the "Array-of-Structs" and "Struct-of-Arrays" layouts.

Let's quickly review what is meant by these two terms before diving in to the performance experiment.

Array of Structs

A very common pattern in C++ is to define custom POD types and put them in standard containers. For example:

// some custom type
struct vec3 { double x, y, z; };

std::vector< vec3 > p = { vec3{1.0, 2.0, 3.0}, vec3{4.0, 5.0, 6.0}, ... };

If we look at the memory locations for the p.data() buffer we see something like

That is, all of the data from p[0], followed by all of the data from p[1] and so on. We use the term "Array of Structs" to describe this layout, since it arises naturally from putting POD structs in arrays.

Struct of Arrays

The other pattern we'll consider in this document is one where data is laid out in the following way:

This layout groups the data by component, rather than struct. This layout tends to be less common in regular C++, as the language currently lacks a natural way to generate it from a general struct definition (although it is possible with circle, and should be possible with C++26 reflection).

Experiment

We'll be considering two basic kernels for this study.

The first kernel takes a list of vectors and translates each of them by some constant:

__global__ void translate_kernel(vec3 * vertices, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) {
    vertices[i] += vec3{0.2, -0.7, 0.3};
  }
}

This kernel represents memory-bound calculations where data access is strided but predictable.

The second kernel takes a triangle mesh and computes the surface normal for each triangle:

__global__ void calculate_normals_kernel(const vec3 * vertices, const vec3i * triangles, vec3 * normals, int n) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < n) {
    vec3i tri = triangles[i];
    vec3 v0 = vertices[tri[0]];
    vec3 v1 = vertices[tri[1]];
    vec3 v2 = vertices[tri[2]];
    normals[i] = normalize(cross(v1 - v0, v2 - v0));
  }
}

The triangles array holds the indices of which nodes correspond to each triangle. For example, {1, 13, 2} might be one of the vec3i entries in that array for the mesh below:

This kernel has an unstructured data access pattern for the vertex coordinates (which account for most of the global memory accesses).

The kernels above are written in terms of the AoS layouts. The equivalent SoA versions are given below.

__global__ void translate_kernel_soa(double * vertices, int num_vertices) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < num_vertices) {
    vertices[i + 0 * num_vertices] += 0.2;
    vertices[i + 1 * num_vertices] -= 0.7;
    vertices[i + 2 * num_vertices] += 0.3;
  }
}

__global__ void calculate_normals_kernel_soa(
  const double * vertices,
  const int * triangles,
  double * normals,
  int num_vertices,
  int num_triangles
) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < num_triangles) {
    const double * vertex_x = vertices + 0 * num_vertices;
    const double * vertex_y = vertices + 1 * num_vertices;
    const double * vertex_z = vertices + 2 * num_vertices;

    int i0 = triangles[i + 0 * num_triangles];
    int i1 = triangles[i + 1 * num_triangles];
    int i2 = triangles[i + 2 * num_triangles];

    vec3 v0 = vec3{vertex_x[i0], vertex_y[i0], vertex_z[i0]};
    vec3 v1 = vec3{vertex_x[i1], vertex_y[i1], vertex_z[i1]};
    vec3 v2 = vec3{vertex_x[i2], vertex_y[i2], vertex_z[i2]};
    vec3 normal = normalize(cross(v1 - v0, v2 - v0));

    normals[i + 0 * num_triangles] = normal[0];
    normals[i + 1 * num_triangles] = normal[1];
    normals[i + 2 * num_triangles] = normal[2];
  }
}

Note: the SoA implementations are more complicated, although some of that complexity can be mitigated with appropriate abstractions.

Results

First, let's look at the translate kernel. The timings below are measurements from an Intel 13900K and an A100 for a mesh with 2621442 vertices and 5242880 triangles:

(cpu, AoS) translate calculation took 3.8359 ms
(cpu, SoA) translate calculation took 2.61891 ms
(gpu, AoS) translate calculation took 0.098816 ms
(gpu, SoA) translate calculation took 0.100147 ms

Transposing the data layout from AoS to SoA did not lead to any improvement on the GPU kernel! Both CUDA kernels achieve around 80% of theoretical memory throughput.

Interestingly: AoS, which is often assumed to be the "CPU" layout, performs worse on the CPU.

Next, let's look at the surface normal calculation.

(cpu, AoS) normal calculation took 14.7559 ms
(cpu, SoA) normal calculation took 17.4566 ms
(gpu, AoS) normal calculation took 0.19866 ms
(gpu, SoA) normal calculation took 0.19968 ms

Once again, we see no benefit when using the SoA layout on the GPU and a small performance regression when using SoA instead of AoS on the CPU.

Explanation

The experimental results seem to contradict the overwhelming guidance that SoA is better on GPU. What's going on?

It's important to remember that:

1. memory addresses are grouped into contiguous 128-byte chunks called cache lines.

2. reading data the first time from global memory is expensive, but reading from L1$ is relatively cheap.

Translation Kernel

With those in mind, let's think about the first kernel. The vector increment

vertices[i] += vec3{0.2, -0.7, 0.3};

compiles to the following SASS code

xxxxxxxxxx
00000100 052a5590 LDG.E.64 R4, [R2.64]
00000100 052a55a0 LDG.E.64 R6, [R2.64+0x8]
00000100 052a55b0 LDG.E.64 R8, [R2.64+0x10]
00000100 052a55c0 DADD R4, R4, c[0x2][0x0]
00000100 052a55d0 DADD R6, R6, c[0x2][0x8]
00000100 052a55e0 STG.E.64 [R2.64], R4
00000100 052a55f0 DADD R8, R8, c[0x2][0x10]
00000100 052a5600 STG.E.64 [R2.64+0x8], R6
00000100 052a5610 STG.E.64 [R2.64+0x10], R8

That is: the kernel loads the x, y, and z components of the vector separately, increments them, and then writes the new values back out to global memory, component-by-component. If we visualize the data access pattern for the first load (each square is 4 bytes, ⬜ = data is not in cache, 🟨 = data is in cache, 🟩 = data being accessed)

we see that the access pattern is not coalesced (i.e. adjacent threads in the warp do not access adjacent locations in memory) and not in cache, resulting in reading 6 cache lines from global memory. However, the memory transactions for the next two loads (the y- and z-components of each vec3) are already in cache from the first load.

This means that for this stride, the AoS layout does not result in redundant reads from global memory. Because of this, its performance ends up being comparable to the SoA data access pattern:

Surface Normal Kernel

The defining feature of kernel 2 is its irregular data access pattern that comes from the triangle to node connectivity array. Unstructured data access patterns are often impossible to coalesce perfectly, but that doesn't mean we can't achieve high performance.

When the mesh nodes and elements are numbered in a way such that

then a given cache line is likely to still be accessed by multiple threads within a warp (or in a block). This means many of the vertex data accesses hit in L1$ or L2$.

There are many equivalent ways to encode the connectivity of a mesh or graph. Some mathematically equivalent representations of the same graph perform much better than others. For example, here is one representation of the mesh shown before, where the nodes are numbered randomly.

As we can see, the random numbering produces triangles that need to access very different parts of the vertex position array. These sorts of accesses are unlikely to be on the same cache lines. In contrast, the mesh below is ordered in a way that respects spatial locality.

The triangles in this mesh have node ids that are "close together", which helps improve cache hit rate.

Followup Questions

Let's explore the follow-up questions posed in the previous section:

1. What would the performance of the translation kernel look like if the vector had 4 components? 8 components? 16 components?

   ---

   Let's update our kernel to accept a template parameter:

   xxxxxxxxxx
   template <uint32_t dim>
   __global__ void translate_kernel(vec<dim> * vertices, int n) {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
       vec<dim> v = vertices[i];
   
       if (dim >  0) v[ 0] += 0.2;
       if (dim >  1) v[ 1] -= 0.1;
       if (dim >  2) v[ 2] += 0.2;
       if (dim >  3) v[ 3] -= 0.3;
       if (dim >  4) v[ 4] += 0.2;
       if (dim >  5) v[ 5] -= 0.8;
       if (dim >  6) v[ 6] -= 0.2;
       if (dim >  7) v[ 7] += 0.9;
       if (dim >  8) v[ 8] -= 0.1;
       if (dim >  9) v[ 9] -= 0.4;
       if (dim > 10) v[10] += 0.3;
       if (dim > 11) v[11] -= 0.1;
       if (dim > 12) v[12] -= 0.2;
       if (dim > 13) v[13] += 0.7;
       if (dim > 14) v[14] += 0.5;
       if (dim > 15) v[15] -= 0.2;
   
       vertices[i] = v;
     }
   }
   

   and plot the performance of this kernel (relative to the theoretical maximum) as a function of dim:

   

   SoA layout has the most consistent performance across structs of different sizes. Unaligned AoS is competitive for small vectors, but slows down as the vector size increases. Adding the following alignment statement to the vec class improves the AoS performance considerably for even sizes:

   xxxxxxxxxx
   template < uint32_t dim, typename T = double >
   struct alignas(sizeof(T) * (dim & -dim)) vec {
     ...
   };
   

   If one needs to work with a vector with an odd size like 13, AoS with alignment and minimal padding is a simple way to get competitive performance with only a small amount of wasted space.

   ---

2. How would this surface normal kernel perform for the mesh with randomly-numbered nodes? How would the size of the mesh affect the performance of SoA vs AoS?

   ---

   The plot below shows the runtimes associated with AoS and SoA implementations on meshes of different sizes and node numbering conventions:

   

   There are several trends worth noting here:

   1. Randomizing the mesh's numbering significantly increases the runtime (by about 10x)

   2. There is no appreciable benefit to SoA over AoS for the original meshes

   3. On the randomized meshes, SoA was slower than AoS, and the discrepancy is more pronounced once the problem size exceeds the L2$ capacity.

      • For random data access patterns, AoS wins because the different components in the struct are contiguous in memory (so at least some of the accesses hit in L1$). On the other hand, grouping data by components means that SoA layouts practically never hit in L1$ for random access patterns.

Summary

There is a lot of misinformation surrounding struct-of-array vs. array-of-struct layouts. Some of the commonly held beliefs like "CPUs prefer AoS" or "GPUs prefer SoA" are not generally true. Although the SoA layout is frequently useful, it is not a requirement to achieve high performance on a GPU. In fact, refactoring a code to adopt an SoA layout can take a considerable amount of effort and frequently does not help performance.

A more useful rule of thumb is: the appropriate data layout depends on the data access pattern.

• for simple strided and coalesced access patterns, both AoS and SoA can work well (on CPU and GPU).

  • On CPUs, SoA often performs better since it is easier for the compiler to vectorize!

  • Naive AoS often underperforms, but alignment and minimal padding make it competitive again.

• for random data access patterns, AoS performs better than SoA.

  • AoS layout ensures that at least some of the struct data members hit in L1$.

• When calculations only need a small subset of the member variables from a struct, SoA wins.

  • However, if many calculations only need a subset of a struct's member variables, that is likely a hint that the struct is not the natural way to group data. In this case, splitting the struct up into smaller, more natural groupings is a better solution.

In summation, many of the dogmatic statements about SoA and AoS that one finds on the internet and in the literature are outdated. In 2010, GPU programming was very different than it is today. Modern GPU hardware has large L2 caches, and tends to be much more forgiving of strided and random data access patterns than the old GPUs.

So, as of 2026 if you find yourself considering how to choose a data layout for your application, please make the decision based on how the data is used and how it interacts with the memory hierarchy.

The sources used in this document are available here.