download - Usando std::vector en código de dispositivo CUDA




toolkit gpu (5)

En el empuje de la biblioteca de cuda, puede usar thrust::device_vector<classT > para definir un vector en el dispositivo, y la transferencia de datos entre el vector STL del host y el vector del dispositivo es muy simple. Puede consultar este útil enlace: http://docs.nvidia.com/cuda/thrust/index.html para encontrar algunos ejemplos útiles.

La pregunta es: ¿hay alguna manera de usar el "vector" de clase en los granos de Cuda? Cuando lo intento, aparece el siguiente error:

error : calling a host function("std::vector<int, std::allocator<int> > ::push_back") from a __device__/__global__ function not allowed

Entonces, ¿hay alguna manera de usar un vector en la sección global? Recientemente intenté lo siguiente:

  1. crear un nuevo proyecto Cuda
  2. ve a las propiedades del proyecto
  3. abrir Cuda C / C ++
  4. ir al dispositivo
  5. cambie el valor en "Generación de código" para establecerlo en este valor: compute_20, sm_20

........ después de eso, pude usar la función de biblioteca estándar printf en mi kernel Cuda.

¿Hay alguna manera de utilizar el vector clase de biblioteca estándar de la forma en que se admite printf en el código del kernel? Este es un ejemplo del uso de printf en el código del kernel:

// this code only to count the 3s in an array using Cuda
//private_count is an array to hold every thread's result separately 

__global__ void countKernel(int *a, int length, int* private_count) 
{
    printf("%d\n",threadIdx.x);  //it's print the thread id and it's working

    // vector<int> y;
    //y.push_back(0); is there a possibility to do this?

    unsigned int offset  = threadIdx.x * length;
    int i = offset;
    for( ; i < offset + length; i++)
    {
        if(a[i] == 3)
        {
            private_count[threadIdx.x]++;
            printf("%d ",a[i]);
        }
    }   
}

Creo que puede implementar un vector de dispositivo por su cuenta, ya que CUDA admite la asignación de memoria dinámica en los códigos de dispositivo. Operador nuevo / eliminar también son compatibles. Aquí hay un prototipo de vector de dispositivo extremadamente simple en CUDA, pero funciona. No ha sido probado lo suficiente.

template<typename T>
class LocalVector
{
private:
    T* m_begin;
    T* m_end;

    size_t capacity;
    size_t length;
    __device__ void expand() {
        capacity *= 2;
        size_t tempLength = (m_end - m_begin);
        T* tempBegin = new T[capacity];

        memcpy(tempBegin, m_begin, tempLength * sizeof(T));
        delete[] m_begin;
        m_begin = tempBegin;
        m_end = m_begin + tempLength;
        length = static_cast<size_t>(m_end - m_begin);
    }
public:
    __device__  explicit LocalVector() : length(0), capacity(16) {
        m_begin = new T[capacity];
        m_end = m_begin;
    }
    __device__ T& operator[] (unsigned int index) {
        return *(m_begin + index);//*(begin+index)
    }
    __device__ T* begin() {
        return m_begin;
    }
    __device__ T* end() {
        return m_end;
    }
    __device__ ~LocalVector()
    {
        delete[] m_begin;
        m_begin = nullptr;
    }

    __device__ void add(T t) {

        if ((m_end - m_begin) >= capacity) {
            expand();
        }

        new (m_end) T(t);
        m_end++;
        length++;
    }
    __device__ T pop() {
        T endElement = (*m_end);
        delete m_end;
        m_end--;
        return endElement;
    }

    __device__ size_t getSize() {
        return length;
    }
};

no puede usar std::vector en el código del dispositivo, debe usar una matriz en su lugar.


No puede usar el STL en CUDA, pero puede usar la biblioteca Thrust para hacer lo que desee. De lo contrario, solo copie el contenido del vector en el dispositivo y opere normalmente.


Hay dos partes en esa respuesta (la escribí). Una parte es fácil de cuantificar, la otra es más empírica.

Restricciones de hardware:

Esta es la parte fácil de cuantificar. El Apéndice F de la guía de programación actual de CUDA enumera una serie de límites estrictos que limitan la cantidad de hilos por bloque que puede tener un lanzamiento de kernel. Si supera cualquiera de estos, su núcleo nunca se ejecutará. Se pueden resumir aproximadamente como sigue:

  1. Cada bloque no puede tener más de 512/1024 subprocesos en total ( Compute Capability 1.xo 2.x y posterior respectivamente)
  2. Las dimensiones máximas de cada bloque están limitadas a [512,512,64] / [1024,1024,64] (Compute 1.x / 2.x o posterior)
  3. Cada bloque no puede consumir más de 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k de registros totales (Compute 1.0.1.1 / 1.2,1.3 / 2.x- / 3.0 / 3.2 / 3.5-5.2 / 5.3 / 6-6.1 / 6.2 / 7.0)
  4. Cada bloque no puede consumir más de 16kb / 48kb / 96kb de memoria compartida (Compute 1.x / 2.x-6.2 / 7.0)

Si se mantiene dentro de esos límites, cualquier kernel que pueda compilar exitosamente se lanzará sin error.

La optimización del rendimiento:

Esta es la parte empírica. El número de subprocesos por bloque que elija dentro de las restricciones de hardware descritas anteriormente puede afectar el rendimiento del código que se ejecuta en el hardware. Cómo se comporta cada código será diferente y la única manera real de cuantificarlo es mediante una comparación cuidadosa y un perfil. Pero de nuevo, muy resumido:

  1. El número de hilos por bloque debe ser un múltiplo redondo del tamaño de urdimbre, que es 32 en todo el hardware actual.
  2. Cada unidad multiprocesador de transmisión en la GPU debe tener suficientes warps activos para ocultar suficientemente toda la memoria diferente y la latencia de la tubería de instrucciones de la arquitectura y lograr el máximo rendimiento. El enfoque ortodoxo aquí es intentar lograr una ocupación óptima del hardware (a lo que se refiere la respuesta de Roger Dahl ).

El segundo punto es un gran tema que dudo que alguien vaya a tratar de cubrir en una sola respuesta de . Hay personas que escriben tesis de doctorado en torno al análisis cuantitativo de aspectos del problema (ver esta presentación de Vasily Volkov de UC Berkley y este documento de Henry Wong de la Universidad de Toronto para ejemplos de cuán compleja es en realidad la pregunta).

En el nivel de entrada, debes saber que el tamaño de bloque que elijas (dentro del rango de tamaños de bloques legales definidos por las restricciones anteriores) puede tener un impacto en la velocidad de ejecución del código, pero depende del hardware usted tiene y el código que está ejecutando. Al hacer un benchmarking, probablemente descubras que la mayoría del código no trivial tiene un "punto óptimo" en el rango de 128-512 hilos por bloque, pero requerirá un análisis de tu parte para encontrar dónde está. La buena noticia es que, debido a que está trabajando en múltiplos del tamaño de la urdimbre, el espacio de búsqueda es muy finito y la mejor configuración para una pieza determinada de código es relativamente fácil de encontrar.