如何解决用cuda优化rc4
| 我做了一些尝试,以便在cuda中实现高效的rc4密码算法。我使用共享内存来存储内部排列状态,并注意存储库的内存布局会浪费时间并在扭曲中进行并行线程访问。我还尝试利用带有'i \'索引的读/写访问是连续的并且可以用32位字打包的事实来最大程度地减少访问次数。最后,我利用常量内存来初始化置换状态。 尽管有这些“聪明”的技巧,但即使考虑到主机和gpu之间的畅通通信可以用来部分覆盖,我仍可以预期仅达到报告的最佳实现的大约50%的吞吐量(例如,参见guapdf Cracker)。计算。我不知道为什么,我正在寻找新的改进想法或对我可能做出的错误假设的评论。 这是我的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
}
解决方法
似乎代码至少部分涉及对字节进行重新排序。如果您使用的是Fermi类的GPU,则可以考虑使用__byte_perm()内在函数,该内在函数映射到Fermi类设备上的硬件指令,并允许人们更高效地对字节进行重新排序。
我假设当您与其他实施方案进行比较时,即是相同的,即在相同类型的GPU上。该代码看起来完全不受计算限制,因此吞吐量将在很大程度上取决于GPU的整数指令吞吐量,并且性能范围很广。
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。