c++ - mundo - programacion cuda linux

Implementación de operadores de rotación de bits utilizando SIMD en CUDA

Question

Sé que StackOverflow no está destinado a preguntar el código a otras personas, pero déjame hablar.

Estoy tratando de implementar algunas funciones AES en el código del dispositivo CUDA C ++. Mientras trataba de implementar el operador de giro de byte izquierdo, me desconcertó ver que no había intrínseco SIMD nativo para eso. Así que comencé una implementación ingenua, pero .... es enorme, y aunque aún no lo he probado, simplemente no será rápido debido a los costosos desembalaje / embalaje ... Entonces, ¿hay una manera de hacerlo? una operación de rotación de bit por byte que sea al menos algo eficiente?

Aquí está el código si quieres echarle un vistazo.

__inline__ __device__ uint32_t per_byte_bit_left_rotate(uint32_t input, uint8_t amount) {
return ((((input & 0xFF) >> 0) << amount) | (((input & 0xFF) >> 0) >> 7) & ~0x100) << 0 |
     ((((input & 0xFF00) >> 8) << amount) | ((input & 0xFF00 >> 8) >> 7) & ~0x100) << 8 |
     ((((input & 0xFF0000) >> 16) << amount) | ((input & 0xFF0000 >> 16) >> 7) & ~0x100) << 16 |
     ((((input & 0xFF000000) >> 24) << amount) | ((input & 0xFF000000 >> 24) >> 7) & ~0x100) << 24; } // The XORs are for clearing the old 7th bit who is getting pushed to the next byte of the intermediate int

Answer 1

CUDA tiene un __byte_perm() intrínseco que se asigna directamente a la instrucción PRMT en el nivel de código de máquina (SASS), que es una instrucción de permutación en bytes. Se puede usar para extraer y combinar bytes de manera eficiente. Para afectar una rotación izquierda en bytes, podemos doblar cada byte, cambiar los pares de bytes por la cantidad deseada, luego extraer y fusionar los cuatro bytes altos de los pares de bytes.

Para la rotación por bytes, solo necesitamos los tres bits más bajos de la cantidad de cambio, ya que una rotación por s es igual a una rotación por s mod 8 . Para mayor eficiencia, es mejor evitar los tipos enteros que comprenden menos de 32 bits, ya que la semántica de C ++ requiere que los tipos enteros más angostos que int se amplíen a int antes de su uso en las expresiones. Esto puede y genera una sobrecarga de conversión en muchas arquitecturas, incluidas las GPU.

El rendimiento de la instrucción PRMT depende de la arquitectura, por lo que el uso de __byte_perm() puede conducir a un código más rápido o más lento que el uso del método clásico SIMD-en-un-registro demostrado en otra respuesta , así que asegúrese de compararlo el contexto de su caso de uso antes de la implementación.

#include <stdio.h>
#include <stdint.h>
#include <stdlib.h>

__device__ uint32_t per_byte_bit_left_rotate (uint32_t input, uint32_t amount)
{
     uint32_t l = __byte_perm (input, 0, 0x1100) << (amount & 7);
     uint32_t h = __byte_perm (input, 0, 0x3322) << (amount & 7);
     return __byte_perm (l, h, 0x7531);
}

__global__ void rotl_kernel (uint32_t input, uint32_t amount, uint32_t *res)
{
    *res = per_byte_bit_left_rotate (input, amount);
}

uint32_t ref_per_byte_bit_left_rotate (uint32_t input, uint32_t amount)
{
   int s = amount & 7;
   uint8_t b0 = (input >>  0) & 0xff;
   uint8_t b1 = (input >>  8) & 0xff;
   uint8_t b2 = (input >> 16) & 0xff;
   uint8_t b3 = (input >> 24) & 0xff;
   b0 = s ? ((b0 << s) | (b0 >> (8 - s))) : b0;
   b1 = s ? ((b1 << s) | (b1 >> (8 - s))) : b1;
   b2 = s ? ((b2 << s) | (b2 >> (8 - s))) : b2;
   b3 = s ? ((b3 << s) | (b3 >> (8 - s))) : b3;
   return (b3 << 24) | (b2 << 16) | (b1 << 8) | (b0 << 0);
}

// Fixes via: Greg Rose, KISS: A Bit Too Simple. http://eprint.iacr.org/2011/007
static unsigned int z=362436069,w=521288629,jsr=362436069,jcong=123456789;
#define znew (z=36969*(z&0xffff)+(z>>16))
#define wnew (w=18000*(w&0xffff)+(w>>16))
#define MWC  ((znew<<16)+wnew)
#define SHR3 (jsr^=(jsr<<13),jsr^=(jsr>>17),jsr^=(jsr<<5)) /* 2^32-1 */
#define CONG (jcong=69069*jcong+13579)                     /* 2^32 */
#define KISS ((MWC^CONG)+SHR3)

// Macro to catch CUDA errors in CUDA runtime calls
#define CUDA_SAFE_CALL(call)                                          /
do {                                                                  /
    cudaError_t err = call;                                           /
    if (cudaSuccess != err) {                                         /
        fprintf (stderr, "Cuda error in file ''%s'' in line %i : %s./n",/
                 __FILE__, __LINE__, cudaGetErrorString(err) );       /
        exit(EXIT_FAILURE);                                           /
    }                                                                 /
} while (0)

// Macro to catch CUDA errors in kernel launches
#define CHECK_LAUNCH_ERROR()                                          /
do {                                                                  /
    /* Check synchronous errors, i.e. pre-launch */                   /
    cudaError_t err = cudaGetLastError();                             /
    if (cudaSuccess != err) {                                         /
        fprintf (stderr, "Cuda error in file ''%s'' in line %i : %s./n",/
                 __FILE__, __LINE__, cudaGetErrorString(err) );       /
        exit(EXIT_FAILURE);                                           /
    }                                                                 /
    /* Check asynchronous errors, i.e. kernel failed (ULF) */         /
    err = cudaThreadSynchronize();                                    /
    if (cudaSuccess != err) {                                         /
        fprintf (stderr, "Cuda error in file ''%s'' in line %i : %s./n",/
                 __FILE__, __LINE__, cudaGetErrorString( err) );      /
        exit(EXIT_FAILURE);                                           /
    }                                                                 /
} while (0)

int main (void)
{
    uint32_t arg, ref, res = 0, *res_d = 0;
    uint32_t shft;

    CUDA_SAFE_CALL (cudaMalloc ((void**)&res_d, sizeof(*res_d)));
    for (int i = 0; i < 100000; i++) {
        arg  = KISS;
        shft = KISS;
        ref = ref_per_byte_bit_left_rotate (arg, shft);
        rotl_kernel <<<1,1>>>(arg, shft, res_d);
        CHECK_LAUNCH_ERROR();
        CUDA_SAFE_CALL (cudaMemcpy (&res, res_d, sizeof (res), 
                                    cudaMemcpyDeviceToHost));
        if (res != ref) {
            printf ("!!!! arg=%08x shft=%d  res=%08x  ref=%08x/n", 
                    arg, shft, res, ref);
        }
    }
    CUDA_SAFE_CALL (cudaFree (res_d));
    CUDA_SAFE_CALL (cudaDeviceSynchronize());
    return EXIT_SUCCESS;
}

Answer 2

El conteo de rotación es el mismo para todos los elementos, ¿verdad?

Desplace toda la entrada hacia la izquierda y hacia la derecha, y luego AND con las máscaras que ponen a cero todos los bits que cruzaron un límite de bytes , para los 4 bytes en un AND. Creo que la amount siempre es una constante de tiempo de compilación en AES, por lo que no tiene que preocuparse por el costo del tiempo de ejecución de generar las máscaras sobre la marcha. Solo deja que el compilador lo haga. (IDK CUDA, pero este parece ser el mismo problema que escribir un SWAR bit-hack con enteros de 32 bits para C ++ normal)

Esto se basa en lo usual (x << count) | (x >> (32-count)) (x << count) | (x >> (32-count)) rotar idioma , con enmascaramiento y un recuento diferente de desplazamiento a la derecha para convertirlo en rotaciones separadas de 8 bits.

inline
uint32_t per_byte_bit_left_rotate(uint32_t input, unsigned amount)
{
    // With constant amount, the left/right masks are constants
    uint32_t rmask = 0xFF >> ((8 - amount) & 7);
    rmask = (rmask<<24 | rmask<<16 | rmask<<8 | rmask);
    uint32_t lmask = ~rmask;

    uint32_t lshift = input << amount;
    lshift &= lmask;
    if (amount == 1) {  // special case left-shift by 1 using an in-lane add instead of shift&mask
        lshift = __vadd4(input, input);
    }
    uint32_t rshift = input >> ((8 - amount) & 7);
    rshift &= rmask;

    uint32_t rotated = lshift | rshift;
    return rotated;
}

Puede ser aún más eficaz enmascarar la entrada de una manera antes del desplazamiento y enmascarar la salida después del cambio ( (in&lmask)<<amount | ((in>>(8-amount))&rmask) , con una máscara diferente. El hardware de NVidia está en orden superescalar y los turnos tienen un rendimiento limitado . Hacerlo de esa manera sería más probable que se ejecute como dos pares de mayúsculas y mayúsculas independientes.

(Esto no intenta evitar C ++ UB con una cantidad> = 32. Consulte las mejores prácticas para las operaciones de cambio circular (rotar) en C ++ . En este caso, creo que cambiar a lshift = input << (amount & 7) haría truco.

Para probar que esto se compila de manera eficiente, miré la salida clang -O3 asm para x86-64 con una amount constante. El explorador del compilador Godbolt tiene compiladores para varias arquitecturas (aunque no para CUDA), así que haga clic en ese enlace y cambie a ARM, MIPS o PowerPC si puede leer esos lenguajes asm más fácilmente que x86.

uint32_t rol7(uint32_t a) {
    return per_byte_bit_left_rotate(a, 7);
}
    mov     eax, edi
    shl     eax, 7
    shr     edi
    and     eax, -2139062144   # 0x80808080
    and     edi, 2139062143    # 0x7F7F7F7F
    lea     eax, [rdi + rax]   # ADD = OR when no bits intersect
    ret

Perfecto, exactamente lo que esperaba.

Un par de casos de prueba:

uint32_t test_rol() {
    return per_byte_bit_left_rotate(0x02ffff04, 0);
}
    // yup, returns the input with count=0
    // return 0x2FFFF04


uint32_t test2_rol() {
    return per_byte_bit_left_rotate(0x02f73804, 4);
}
    // yup, swaps nibbles
    // return 0x207F8340

Este es el mismo tipo de cosas que debe hacer para los cambios de 8 bits con x86 SSE2 / AVX2, ya que la granularidad de cambio más pequeña que admite el hardware es de 16 bits.