The #1 Way to Pass float4 to PyCUDA Kernels (2025)

data = np.random.rand(N, 4).astype(np.float32)
kernel(cuda.In(data), ...)

CUDA Kernel Side:

__global__ void my_kernel(float* data) {
  int idx = threadIdx.x + blockIdx.x * blockDim.x;
  
  // Manual access
  float x = data[idx * 4 + 0];
  float y = data[idx * 4 + 1];
  float z = data[idx * 4 + 2];
  float w = data[idx * 4 + 3];
  
  // ... do work
}

This is better—at least you only have one memory transfer. But it’s still not ideal. The code is less readable, the manual indexing `(idx * 4 + offset)` is prone to off-by-one errors, and you’re not explicitly telling the CUDA compiler that you’re working with a vector type. While the memory layout is contiguous, you're missing out on the clarity and potential compiler optimizations that come from using the native float4 type directly.

The #1 Method: Using NumPy Structured Arrays

Here it is, the moment you've been waiting for. The best way to pass float4 data to a PyCUDA kernel is by using NumPy Structured Arrays.

A structured array allows you to define a custom dtype that mimics a C-style struct. You can define fields by name and type. By creating a dtype that exactly matches CUDA's float4, you create a memory layout in Python that is a perfect, 1:1 match for what the GPU expects. No casting, no manual indexing, no fuss.

Step-by-Step Guide: From NumPy to CUDA Kernel

Let's walk through a complete, practical example.

Step 1: Define the Structured dtype in Python

First, we define our custom data type. A float4 is just four 32-bit floats. In NumPy, the type code for a 32-bit float is 'f4'.

import numpy as np

# Define a dtype that matches the CUDA float4 struct
# We can name the fields anything, but x,y,z,w is conventional.
float4_dtype = np.dtype([('x', np.float32), 
                         ('y', np.float32), 
                         ('z', np.float32), 
                         ('w', np.float32)])

This `float4_dtype` object now describes a memory block of 16 bytes (4 floats * 4 bytes/float), with named fields. NumPy will handle all the memory alignment for us, ensuring it's packed just like a C struct.

Step 2: Create and Populate the Array

Now, we create our host array using this new dtype. Notice how we can access the fields by name, which makes the code incredibly readable.

import pycuda.autoinit
import pycuda.driver as cuda
from pycuda.compiler import SourceModule

N = 1024 # Number of float4 vectors

# Create the host array with our custom dtype
host_data = np.zeros(N, dtype=float4_dtype)

# Populate the data using named fields
host_data['x'] = np.arange(N, dtype=np.float32)
host_data['y'] = np.arange(N, dtype=np.float32) * 2
host_data['z'] = 1.0
host_data['w'] = 0.0

# Allocate memory on the GPU and copy the data
# PyCUDA understands the .nbytes attribute of any NumPy array.
device_data = cuda.mem_alloc(host_data.nbytes)
cuda.memcpy_htod(device_data, host_data)

Step 3: Write the CUDA Kernel

Here’s the beautiful part. The kernel becomes much cleaner. We can accept a pointer to float4 directly. CUDA C has built-in support for this type.

kernel_code = """
__global__ void my_vector_kernel(float4* data, float4* out_data) {
  int idx = threadIdx.x + blockIdx.x * blockDim.x;

  // Access the entire struct at once - this is a coalesced read!
  float4 vec = data[idx];

  // Operate on the vector components naturally
  vec.x *= 2.0f;
  vec.y += vec.z;

  // Write the result back - this is a coalesced write!
  out_data[idx] = vec;
}
"""

Notice the lack of manual indexing. We just grab `data[idx]` and get a full `float4` struct. It's clean, safe, and exactly what the GPU hardware is optimized for.

Step 4: Launch the Kernel

Launching the kernel is straightforward. Just pass the GPU pointer as an argument.

# Compile the kernel
mod = SourceModule(kernel_code)
kernel = mod.get_function("my_vector_kernel")

# Prepare output buffer
host_result = np.empty_like(host_data)
device_result = cuda.mem_alloc(host_result.nbytes)

# Launch!
block_size = 256
grid_size = (N + block_size - 1) // block_size
kernel(device_data, device_result, block=(block_size, 1, 1), grid=(grid_size, 1))

# Copy result back to host
cuda.memcpy_dtoh(host_result, device_result)

print("First result vector:", host_result[0])

Performance Showdown: A Clear Winner Emerges

So, does this method actually make a difference? Absolutely. While a full benchmark is beyond the scope of a single post, the conceptual differences lead to clear performance results.

The structured array method wins on all fronts. It provides the optimal memory layout for the GPU, leading to the fastest possible kernel execution, while also making the Python and C++ code significantly cleaner and easier to understand.

A Quick Note on Alignment and pycuda.driver.Struct

For more complex structs with mixed types, C/C++ compilers can introduce padding bytes for alignment purposes, which can be tricky to replicate perfectly with NumPy dtypes. In those advanced scenarios, PyCUDA provides its own utility, pycuda.driver.Struct. You can define a struct in a C-like string and PyCUDA will compute the correct size and offsets.

However, for standard vector types like float4, int2, etc., which contain only one data type, NumPy's structured arrays are perfectly aligned by default and are generally easier and more "Pythonic" to work with.

Conclusion: Your New Default for Vector Data

Passing data between Python and CUDA doesn't have to be a black art. By leveraging NumPy's structured arrays, you can create a seamless, high-performance pipeline for vector data types like float4.

The key takeaway is simple: match your NumPy dtype to your CUDA C struct. This approach provides:

• Peak Performance: Ensures perfectly coalesced memory access in your kernels.
• Code Readability: Named fields in Python and native struct access in CUDA make your code self-documenting.
• Maintainability: Reduces the risk of bugs from manual pointer arithmetic and memory management.

So, the next time you need to pass vector data to a PyCUDA kernel, don't reach for the old, clunky methods. Make NumPy structured arrays your default choice. Stop fighting with memory layouts and start writing the clean, fast, and professional GPU code you've always wanted to.