Estoy tratando de comprender la integrate_functor en particles_kernel.cu de CUDA ejemplos:

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);
    }
};

Llamamos make_float4(pos, age) pero make_float4 se define en vector_functions.h como

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;
} 

Son CUDA los tipos de vectores (float3 y float4) más eficientes para la GPU y cómo el compilador sabe cómo la función de sobrecarga make_float4?

  • Creo que usted encontrará que hay una serie de instancias de make_float4, y en la que se ha publicado no es el que está siendo utilizado en ese código..
  • Usted debe buscar en/uso de las funciones de vector_types.h en el CUDA incluyen directorio. Con un adecuado tipo de vector (es decir, float4), el compilador puede crear instrucciones que se cargue toda la cantidad en una sola transacción. Dentro de los límites, esto puede evitar el am/SoA problema, por cierto vector de arreglos. Así que, sí, puede ser más eficiente, dependiendo de lo que usted está comparando.
  • Así que en términos de memoria la alineación es mejor usar float4 en lugar de float3? En el ejemplo que uso float4 para el almacenamiento y float3 para las operaciones. No utilizar los Datos.w
  • El hardware GPU proporciona instrucciones de carga para 32-bit, 64-bit y 128-bits de datos, que se asigna a la float, float2, y float4 tipos de datos (así como a la int, int2, y int4 tipos). Los datos deben ser naturalmente alineado para las instrucciones de carga para que funcione correctamente y, en general, más amplio cargas de proporcionar una mayor pico de ancho de banda de memoria. Así float4 es preferido sobre float3 por razones de rendimiento.
  • Gracias me encontré con un conjunto de constructores en el encabezado helper_math.h
InformationsquelleAutor ilciavo | 2014-10-31

Dejar respuesta

Please enter your comment!
Please enter your name here