用cuda優化rc4

Question

我做了一些嘗試在cuda中實現rc4密碼算法的高效。我使用共享內存來存儲內部排列狀態，照顧分區內存佈局，以便在變形中使用並行線程訪問來處理時間。我還試圖通過使用'i'索引的讀/寫訪問是連續的並且可以打包成32位字的事實來最小化訪問次數。最後，我利用常量內存來初始化排列狀態。即使考慮到主機和GPU之間的暢通通信可以用於實現最佳報告實現（例如參見guapdf cracker）的吞吐量，我預計只能獲得大約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

看來碼至少部分地涉及到重新排序字節。如果您使用的是費米級GPU，則可以使用映射到費米級設備上的硬件指令的__byte_perm（）內部函數，並允許更有效地重新排序字節。

我認爲當你與其他實現比較時，它是蘋果對蘋果，即在相同類型的GPU上。該代碼看起來完全受計算限制，所以吞吐量將很大程度上取決於GPU的整數指令吞吐量，並且性能頻譜很寬。