Оптимизация RC4 с CUDA - PullRequest
Новая
       8

Оптимизация RC4 с CUDA

1 голос
/ 22 июня 2011

Я сделал несколько попыток реализовать эффективный алгоритм шифрования rc4 в cuda. Я использовал разделяемую память для хранения состояния внутренней перестановки, заботясь о расположении банковской памяти и уменьшении времени с параллельным доступом к потокам в деформации. Я также попытался свести к минимуму количество обращений, используя тот факт, что доступ для чтения / записи с индексом 'i' является смежным и может быть упакован в 32-разрядные слова. Наконец, я использовал постоянную память для инициализации состояния перестановки.

Несмотря на эти «хитрые» уловки, я могу рассчитывать на достижение только примерно 50% пропускной способности лучших из представленных реализаций (см., Например, guapdf cracker), даже принимая во внимание, что неблокируемая связь между хостом и gpu может использоваться для частичной покрыть вычисление. Я не могу понять, почему, и я ищу новые идеи по улучшению или комментарии о плохих предположениях, которые я мог бы сделать.

Вот игрушечная реализация моего ядра KSA (установка ключа) с ключом, уменьшенным до 4 байтов. 

__constant__ unsigned int c_init[256*32/4];

__global__ void rc4Block(unsigned int *d_out, unsigned int *d_in)
{
__shared__ unsigned int s_data[256*32/4];

int inOffset  = blockDim.x * blockIdx.x;
int in  = inOffset + threadIdx.x;
unsigned int key, u;

// initialization 
key = d_in[in];

for(int i=0; i<(256/4); i++) {  // read from constant memory
  s_data[i*32+threadIdx.x] = c_init[i*32+threadIdx.x];
}
// key mixing
unsigned char j = 0;
unsigned char k0 = key & 0xFF;
unsigned char k1 = (key >> 8) & 0xFF;
unsigned char k2 = (key >> 8) & 0xFF;
unsigned char k3 = (key >> 8) & 0xFF;

for(int i=0; i<256; i+=4) { // unrolled

  unsigned int u, sj, v;
  unsigned int si = s_data[(i/4)*32+threadIdx.x];
  unsigned int shiftj;

  u = si & 0xff;
  j = (j + k0 + u) & 0xFF;
  sj = s_data[(j/4)*32+threadIdx.x];
  shiftj = 8*(j%4);
  v = (sj >> shiftj) & 0xff;
  si = (si & 0xffffff00) | v;
  sj = (sj & ~(0xFFu << (8*(j%4)))) | (u << shiftj);
  s_data[(j/4)*32+threadIdx.x] = sj;

  u = (si >> 8) & 0xff;
  j = (j + k1 + u) & 0xFF;
  sj = s_data[(j/4)*32+threadIdx.x];
  shiftj = 8*(j%4);
  v = (sj >> shiftj) & 0xff;
  si = (si & 0xffff00ff) | (v<<8);
  sj = (sj & ~(0xFFu << (8*(j%4)))) | (u << shiftj);
  s_data[(j/4)*32+threadIdx.x] = sj;

  u = (si >> 16) & 0xff;
  j = (j + k2 +u) & 0xFF;
  sj = s_data[(j/4)*32+threadIdx.x];
  shiftj = 8*(j%4);
  v = (sj >> shiftj) & 0xff;
  si = (si & 0xff00ffff) | (v<<16);
  sj = (sj & ~(0xFFu << (8*(j%4)))) | (u << shiftj);
  s_data[(j/4)*32+threadIdx.x] = sj;

  u = (si >> 24) & 0xff;
  j = (j + k3 + u) & 0xFF;
  sj = s_data[(j/4)*32+threadIdx.x];
  shiftj = 8*(j%4);
  v = (sj >> shiftj) & 0xff;
  si = (si & 0xffffff) | (v<<24);
  sj = (sj & ~(0xFFu << (8*(j%4)))) | (u << shiftj);
  s_data[(j/4)*32+threadIdx.x] = sj;

  s_data[(i/4)*32+threadIdx.x] = si;
}
d_out[in] = s_data[threadIdx.x]; // unrelevant debug output 
}

1 Ответ

1 голос
/ 22 июня 2011

Кажется, код хотя бы частично включает в себя переупорядочение байтов.Если вы используете графический процессор класса Fermi, вы можете использовать встроенную функцию __byte_perm (), которая сопоставляется с аппаратной инструкцией на устройствах класса Fermi и позволяет более эффективно переупорядочивать байты.

Полагаю,если сравнивать с другими реализациями, то это яблоки к яблокам, то есть на графическом процессоре того же типа.Этот код выглядит полностью привязанным к вычислениям, поэтому пропускная способность будет в значительной степени зависеть от пропускной способности целочисленных команд графического процессора, а спектр производительности будет широким.

...