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 wirfloat4
), 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
, undfloat4
Datentypen (sowie dieint
,int2
, undint4
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. Sofloat4
ist bevorzugt überfloat3
aus performance-Gründen. - Danke, ich fand eine Reihe von Konstruktoren in der Kopfzeile
helper_math.h