Effizienz der CUDA-Vektor-Typen (float2, float3, float4)

Ich versuche zu verstehen, die integrate_functor im particles_kernel.cu von CUDA-Beispiele:

struct integrate_functor
{
    float deltaTime;    
    //constructor for functor
    //...

    template <typename Tuple>
    __device__
    void operator()(Tuple t)
    {
        volatile float4 posData = thrust::get<2>(t);
        volatile float4 velData = thrust::get<3>(t);

        float3 pos = make_float3(posData.x, posData.y, posData.z);
        float3 vel = make_float3(velData.x, velData.y, velData.z);

        //update position and velocity
        //...

        //store new position and velocity
        thrust::get<0>(t) = make_float4(pos, posData.w);
        thrust::get<1>(t) = make_float4(vel, velData.w);
    }
};

Nennen wir make_float4(pos, age) aber make_float4 definiert ist, in vector_functions.h als

static __inline__ __host__ __device__ float4 make_float4(float x, float y, float z, float w)
{
    float4 t; t.x = x; t.y = y; t.z = z; t.w = w; return t;
} 

Sind CUDA-Vektor-Typen (float3 und float4) effizienter für die GPU und wie weiß der compiler wissen, wie überlastet die Funktion make_float4?

  • Ich glaube, Sie werden feststellen, es gibt eine Anzahl von Instanzen von make_float4 und der, den du gepostet hast ist nicht das einzige Lebewesen in diesem code..
  • Sollten Sie auf der Suche an/der Verwendung von Funktionen aus vector_types.h in die CUDA-include-Verzeichnis. Mit einem geeigneten Vektor-Typs (sagen wir float4), der compiler erstellen können, die Anweisungen, die Last der gesamten Menge in einer einzigen Transaktion. Innerhalb von Grenzen können diese arbeiten rund um die AoS/SoA-problem, für bestimmte Vektor-arrangements. Also, ja, kann es effizienter sein, je nachdem, was Sie sind, zu vergleichen.
  • Also in Bezug auf Speicher-alignment es besser float4 statt float3? In dem Beispiel, das Sie verwenden, float4 für die Lagerung und float3 für Operationen. Sie verwenden nicht die Daten.w
  • Die GPU-hardware bietet load-Anweisungen für 32-bit, 64-bit und 128-bit-Daten, die eine Zuordnung zu den float, float2, und float4 Datentypen (sowie die int, int2, und int4 Arten). Die Daten müssen natürlich ausgerichtet für den load-Anweisungen, um korrekt zu arbeiten und im Allgemeinen größere Lasten bieten höhere maximale Speicherbandbreite. So float4 ist bevorzugt über float3 aus performance-Gründen.
  • Danke, ich fand eine Reihe von Konstruktoren in der Kopfzeile helper_math.h
InformationsquelleAutor ilciavo | 2014-10-31
Schreibe einen Kommentar