Implementación de AES usando CUDA -- ++ campo con performance campo con aes campo con cuda camp codereview Relacionados El problema

Implementation of AES using CUDA


5
vote

problema

Español

Estoy tratando de implementar AES en la GPU utilizando la programación de CUDA. Utilizo 4 Toxes en mi implementación que requiere 4kb de memoria GPU. He usado una matriz de 1KB para un simple texto de 1KB. Primero, todo tipo de texto, se copiaría a la memoria GPU, luego se iniciaría el cifrado usando

  System.Timers.Timer1  

entonces, el kernel global se ejecutará:

  System.Timers.Timer2  

System.Timers.Timer3 realiza una ronda de algoritmo AES en un estado de 1024 bits, luego 32 bytes de la matriz con 32 bits del estado.

  System.Timers.Timer4  

No hay problema en el código o en mi programa, pero el gran problema es su velocidad: solo 1 megabyte por segundo. No sé por qué este programa es tan lento.

  • Yo uso Microsoft Visual C ++ en Windows 8.1 con CUDA Toolkit 7.
  • My GPU es un GeForce GT 720M
    • Capacidad de cómputo: 2.1
    • Número de SMS: 2
    • reloj gráfico (MHz): 625
    • reloj de procesador (MHz): 1250
    • reloj de memoria (MHz): 800
  • Mi procesador es el microprocesador Intel Core (TM) I5-3337U CPU @ 1.80GHz (4 CPUS).

Código fuente del algoritmo se puede encontrar aquí .

Original en ingles

I am trying to implement AES on GPU using CUDA programming. I use 4 TBoxes in my implementation that requires 4kB of GPU Memory. I have used a 1KB array for 1KB plaintext. first all of plaintext would be copied to GPU memory, then encryption would be started using

cudaMemcpyToSymbolAsync(DEV_message, H_Message, 1024, 0, cudaMemcpyHostToDevice); 

Then, the global kernel will run:

AESROUND<<< 8, 16, 16 >>>(1024); 

AESROUND() performs one round of AES algorithm on a 1024-bit state, then 32 bytes of the array will XOR with 32 bits of the state.

__global__ void AESROUND_AD(const int SIZE) {     __shared__ unsigned char dev_rkey[16];     __shared__ unsigned char dev_sh_state[16];      int tid = blockIdx.x * 16 + threadIdx.x;     if (tid < 128)     {         for (long long i = 0; i <SIZE/32; i++)         {                dev_sh_state[threadIdx.x] = dev_state[(tid + 112) % 128];             dev_rkey[threadIdx.x] = dev_state[tid];             __syncthreads();              if (threadIdx.x < 4)             {                 U32 v1 = ((U32*)dev_TE0)[*(dev_sh_state + threadIdx.x * 4)];                 U32 v2 = ((U32*)dev_TE1)[*(dev_sh_state + ((threadIdx.x * 4 + 5) % 16))];                 U32 v3 = ((U32*)dev_TE2)[*(dev_sh_state + ((threadIdx.x * 4 + 10) % 16))];                 U32 v4 = ((U32*)dev_TE3)[*(dev_sh_state + ((threadIdx.x * 4 + 15) % 16))];                  ((U32*)dev_sh_state)[threadIdx.x] = v1 ^ v2 ^ v3 ^ v4 ^ ((U32*)dev_rkey)[threadIdx.x];             }             __syncthreads();              dev_state[tid] = dev_sh_state[threadIdx.x];             __syncthreads();          }     } } 

There is no problem in code or in my program, but the big problem is its speed: just 1 megabyte per second. I don't know why this program is so slow.

  • I use Microsoft Visual C++ in Windows 8.1 with CUDA Toolkit 7.
  • My GPU is a GeForce GT 720M
    • Compute capability: 2.1
    • Number of SMs: 2
    • Graphic clock (MHz): 625
    • Processor clock (MHz): 1250
    • Memory clock (MHz): 800
  • My processor is Intel Core(TM) i5-3337U CPU @ 1.80GHz (4 CPUs) microprocessor.

Source code of the algorithm can be found here.

           

Lista de respuestas

1
 
vote
vote
La mejor respuesta
 

Hay algunos problemas que veo en su código. No estoy seguro de si todo lo que importa, pero intentaré enumerarlos de todos modos.

  1. Usted dijo que llame AESROUND<<< 8, 16, 16 >>>(1024); , pero usted presente __global__ void AESROUND_AD(const int SIZE) . Son AESROUND() y AESROUND_AD() los mismos kernels? ¿Es eso un error tipográfico?

  2. Llamar AESROUND<<< 8, 16, 16 >>>(1024); significa (entre otros) que asigna dinámicamente 16 (tercer parámetro del código <<< >>> ) de un segmento de memoria compartida de tipo a-be -Definene dentro del kernel (declarado extern __shared__ whatever mySegment[]; ). Así que asigna ese segmento en su llamada de kernel, pero en realidad no se declare en el cuerpo del kernel. Sin embargo, otro error tipográfico?

  3. En el propio kernel, se hace referencia dev_state que no se declara en ningún lugar. ¿Es eso una variable global? ¿No debería ser mejor poner un parámetro de entrada del kernel?

  4. El cálculo de tid supone explícitamente que el tamaño del bloque es de 16, que es cierto aquí. Sin embargo, si esto cambia al llamar al kernel, entonces la fórmula se volvería falsa. Así que si esto no se supone que cambia, use una macro para el tamaño del bloque (como #define BLOCKSIZE 16 ) y úselo en su definición de __global__ void AESROUND_AD(const int SIZE)0 y la llamada a __global__ void AESROUND_AD(const int SIZE)1 . Pero si este tamaño puede cambiar en la llamada del kernel, use __global__ void AESROUND_AD(const int SIZE)2 . Y como un comentario más general, trate de evitar "números mágicos" en su código, o al menos colóquelos en macros o __global__ void AESROUND_AD(const int SIZE)3 variables, para que le dé un nombre sensible y, por lo tanto, un significado. < / p>

Ahora, la pregunta inicial fue sobre el rendimiento:
Bueno, supongo que todo esto se reduce a un problema simple: no maximiza ni su ocupación principal, ni su ancho de banda de memoria. De hecho, solo solicita 16 hilos por bloque, lo que no es suficiente para una deformación. Por lo tanto, independientemente de su arquitectura de GPU real, en cualquier momento, a lo sumo, solo la mitad de sus núcleos se utilicen. Además, una vez que los datos iniciales cargados en la memoria compartida, solo calcula usando 4 hilos, por lo que solo utiliza 1/4 de los núcleos disponibles como máximo.

Considerando lo que puedo ver (y una especie de entender) de su kernel, tengo la sensación de que es mejor que use un solo bloque de 128 hilos por mensaje para calcular, y para cargarlo toda la misma en su memoria compartida. Luego, para calcular como lo hace en su núcleo, pero con ahora 4 veces más hilos involucrados. Esto tendría el doble de beneficio de permitirle evitar tener que almacenarlo resultados parciales en la memoria global en cada ronda, y para aumentar su ocupación.

y para utilizar completamente su GPU, deberá calcular varios mensajes en paralelo, uno por bloque.

 

There are a few issues that I see in your code. I'm not sure if all matter, but I'll try to list them anyway.

  1. You said you call AESROUND<<< 8, 16, 16 >>>(1024);, but you present __global__ void AESROUND_AD(const int SIZE). Are AESROUND() and AESROUND_AD() the same kernels? Is that a typo?

  2. Calling AESROUND<<< 8, 16, 16 >>>(1024); means (amongst other) that you dynamically allocate 16 (third parameter of the <<< >>>) elements of a shared memory segment of type to-be-defined inside the kernel (declared extern __shared__ whatever mySegment[];). So you allocate that segment in your kernel call, but do not actually declare in the kernel body. Yet another typo?

  3. In the kernel itself, you reference dev_state which isn't declared anywhere. Is that a global variable? Shouldn't that be better put a input parameter of the kernel?

  4. The computation of tid explicitly assumes that the block size is 16, which is true here. However, if this changes when calling the kernel, then the formula would become false. So if this isn't supposed to change, then use a macro for the block size (such as #define BLOCKSIZE 16) and use it in your definition of tid, and the call to AESROUND(). But if this size might change in the call of the kernel, then use tid = blockIdx.x * blockDim.x + threadIdx.x;. And as a more general comment, try to avoid "magic numbers" in your code, or at least put them into macros or const variables, so that you give them a sensible name and therefore a meaning.

Now, the initial question was about performance:
Well, I guess that this all boils down to one simple issue: you do not maximise neither your core occupancy, nor your memory bandwidth. Indeed, you only request 16 threads per block, which isn't even enough for a warp. Therefore, irrespective of your actual GPU architecture, at any moment, at most only half of your cores are utilised. Moreover, once the initial data loaded in shared memory, you only compute using 4 threads, so only utilising 1/4 of the available cores at most.

Considering what I can see (and sort of understand) of your kernel, I have the feeling that you'd better be using one single block of 128 threads per message to compute, and to load all of it at once in your shared memory. Then to compute as you do in your kernel, but with now 4 times more threads involved. This would have the double benefit of permitting you to avoid having to store you partial results in global memory at each round, and to increase your occupancy.

And to fully utilise your GPU, you would need to compute several messages in parallel, one per block.

 
 
 
 

Relacionados problema

16  Generando números primos usando tamiz de eratóstenes con CUDA  ( Generating prime numbers using sieve of eratosthenes with cuda ) 
Estoy aprendiendo a Cuda y escribí un pequeño programa que genera números primos utilizando el tamiz de Eratóstenes. (Conozco las limitaciones de CUDA, especi...

2  Kernel de CUDA para comparar las entradas de matriz, ponderadas con un patrón  ( Cuda kernel to compare matrix entries weighted with a pattern ) 
Me pregunto si es posible optimizar este código en CUDA. ¿Podría obtener alguna sugerencia cómo? El algoritmo equivalente funciona más rápido en Matlab para m...

3  Multiplicación compleja e integración con CUDA  ( Complex multiplication and integration with cuda ) 
Quiero realizar la multiplicación en dos vectores e integrarlo en un vector llamado public class UpdateAltitude : View { public float AltitudeFactor ...

3  Multiplicación compleja e integración con CUDA  ( Complex multiplication and integration with cuda ) 
Quiero realizar la multiplicación en dos vectores e integrarlo en un vector llamado public class UpdateAltitude : View { public float AltitudeFactor ...

2  Convoluciones con memoria compartida en CUDA  ( Convolutions with shared memory in cuda ) 
Estoy escribiendo un kernel de convolución de desenfoque de ruido en CUDA __global__ void noiseReduction(float *im, float *NR, int height, int width) { ...

3  Una primitiva para dormir para el código del lado del dispositivo CUDA  ( A sleep primitive for cuda device side code ) 
Hasta ahora he sido ahorrado la necesidad de desperdiciar cualquier ciclo en la GPU, pero parece que esto podría cambiar. Después de obtener información, me f...

3  Programa CUDA que emula un tipo de autómata celular más lento en la GPU que en la CPU  ( Cuda program that emulates a kind of cellular automata slower on the gpu than on ) 
Soy un muy nuevo en Cuda, así que todavía estoy tratando de entender cómo hacer el mejor uso de la GPU. He portado a un algoritmo C a él. Este algoritmo funci...

6  CLASE DE MATRIX CUDA  ( Cuda matrix class ) 
Comencé a trabajar en esta clase CUDA C MATRIX para aprender tanto la programación orientada a objetos en C ++ como para aprender CUDA. El objetivo inicial de...

8  Vector 3D CUDA KERNEL  ( 3d vector cuda kernel ) 
Diseñé este kernel de CUDA para calcular una función en un dominio 3D: p1 y Ap son los vectores en 3D que realmente se implementan como una sola matriz ...

1  Número de espejos a bordes de intervalo, hasta que se encuentre en el intervalo  ( Mirrors number at borders of interval untill it lays in the interval ) 
Tengo la siguiente función simple: __device__ mirror(int index , int lB, int uB) { while(index < lB || index >= uB) { if(index < lB) { ...




© 2022 respuesta.top Reservados todos los derechos. Centro de preguntas y respuestas reservados todos los derechos