Python modul PyCUDA: rad s vektorskim tipovima podataka

PyCUDA, kao i CUDA, ima ugrađenu podršku za vektorske tipove: višedimenzionalni podaci s jednom do četiri komponente, adresirane kao .x, .y, .z, .w. Neki od postojećih predefiniranih vektorskih tipova podataka su:

struct uchar1
{
  unsigned char x;
};

struct __align__(4) ushort2
{
  unsigned short x, y;
};

struct uint3
{
  unsigned int x, y, z;
};

struct __align__(16) float4
{
  float x, y, z, w;
};

Uočimo funkciju __align__(), koja se koristi za poravnanje podataka. Naime, GPU može iz memorije u jednoj instrukciji dohvatiti:

• 32 bita, odnosno 4 bajta, za što se koristi __align(4),
• 64 bita, odnosno 8 bajta, za što se koristi __align(8),
• 128 bita, odnosno 16 bajta, za što se koristi __align(16).

Primjeri slučajeva korištenja vektorskih tipova podataka su:

• slike s 8-bitnom dubinom boje (uchar3) (RGB => x, y, z),
• slike s 24-bitnom dubinom boje i transparencijom (uint4) (RGBA => x, y, z, w),
• sustavi fizikalnih objekata u prostoru (float3) (x, y, z: prostorne koordinate),
• sustavi fizikalnih objekata u prostorno-vremenskom kontinuumu (float4) (x, y, z, w: 3 prostorne koordinate i jedna vremenska).

Sa predefiniranim vektorskim tipovima podataka radimo na intuitivan način; zrno je oblika

#include <stdio.h>

__global__ void print_coordinates (float3 point)
{
  printf ("x: %3.2f, y: %3.2f, z: %3.2f\\n", point.x, point.y, point.z);
}

Pridruženi Python kod je oblika

import pycuda.autoinit
import pycuda.driver as drv
import pycuda.gpuarray as ga
import numpy as np
from pycuda.compiler import SourceModule

mod = SourceModule(open("vector_types.cu").read())

print_coordinates = mod.get_function("print_coordinates")

point_in_space = ga.vec.make_float3(3.6, 5.4, 2.25)

print_coordinates(point_in_space, block=(1,1,1), grid=(1,1))

Author: Vedran Miletić, Kristijan Lenković