rc4とcudaの最適化

Question

私は、cudaでrc4暗号アルゴリズムの効率的な実装を試みました。私は、内部の置換状態を保存するために共用メモリを使用しました。これは、ワープ内の並列スレッドへのアクセスで、バンクされたメモリのレイアウトを時間のペナルティに費やしました。また、 'i'インデックスの読み書きアクセスが連続しており、32ビットワードでパックすることができるという事実を利用して、アクセス回数を最小限に抑えようとしました。最後に、定数メモリを使って置換状態を初期化しました。

これらの「巧妙な」トリックにもかかわらず、ホストとGPUの間のブロックされていない通信を使用することを考慮しても、報告された実装（たとえばguapdfクラッカーなど）のスループットの約50％部分的に計算をカバーする。私は理由を理解することができず、新しい改善アイデアや悪い前提についてのコメントを探しています。

私の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 
} 

Answer 1

コードは、少なくとも部分的に再発注バイトを必要とするようです。 FermiクラスのGPUを使用している場合は、Fermiクラスのデバイスのハードウェア命令にマップされる__byte_perm（）組み込み関数を使用して、バイトをより効率的に並べ替えることができます。

他の実装と比較すると、リンゴとリンゴが同じタイプのGPUであると仮定します。このコードは完全に境界を計算しているため、スループットは主にGPUの整数命令のスループットに依存し、パフォーマンススペクトラムは広いです。