搜索此博客

2016年12月23日星期五

Curious Case of Rowhammer: Flipping Secret Exponent Bits Using Timing Analysis

本邮件内容由第三方提供，如果您不想继续收到该邮件，可 点此退订 。
Curious Case of Rowhammer: Flipping Secret Exponent Bits Using Timing Analysis  阅读原文»

论文下载

Sarani Bhattachary,Indian Institute of Technology,CHES’16

Combing Timing Analysis and Rowhammer

攻击模型

  • decryption oracle:RSA，对输入的密文解密，然后返回解密结果
  • Adversary：普通用户，已知明密文，可以和decryption oralce交互。目标是通过rowhammer让解密指数出错，因此需要知道secret在内存中的bank地址
  • spy：adversary引入的进程，监视解密过程，使得解密进程访问数据都是从内存中访问，并利用时间旁路确定secret所处的bank
  • k cores,k slices,c cache sets/slice, m way associative

Fig

确定evivtion set

  • Prime+Probe整体思想：由于cache组相连，攻击者可以占满一个cache set的所有组（Prime），当victim有cache进入这个set时，攻击者原有的cache被挤出，攻击者可以再次访问自己的数据看到时间差（Probe）。。

确定secret映射到的LLC slice

  • Prime+Probe可以确定secret所在的cacheset，攻击者可以找出m*k个映射到同一个set的地址，不断在这些地址上Prime+Probe，就可以把secret从缓存中逐出，保证decryption oracle每次都是从内存中取值

确定secret映射到的DRAM bank

  • row buffer collison的时间旁路（Drammer中也是利用这个来确定rowsize）
  • 当同时访问处于同一个bank中但在不同row的内存地址时，row buffer会清空，访问时间被比其他情况下的多 （两个进程的Collision只是理想中能发生的)

在bank中进行rowhammer

  • 在确定了secret所在的bank后，由于无法进一步获知secret所处的row，所以作者在其中进行随机的hammer，如果解密结果不对说明flip影响到了解密进程

Experimental Validation for Inducing Bit Flips On secret

  • 物理地址与cache slice的关系（RAID15）
  • 4核，12路组相连，4个cache slice，每个有2048个cache sets。
  • 攻击者mmap大块内存，通过pagemap得到物理地址，利用映射关系算出cache slice和cache set

Fig

Fig

Countermeasures and Discussion

  • hardware
  • software：ANVIL，统计cache miss
  • fault countermeasure: OpenSSL,计算两次，如果可以稳定的造出两次错误，这样的防护是无效的

黑客技术官网地址：http://www.hackdig.com/

NVIDIA TX-1 的 Zero Copy 和 Pinned Memory  阅读原文»

公司发的圣诞礼物被一个自以为是的胖女人拿走了，不开心

众所周知GPGPU的性能瓶颈为PCI-E传输速度，数据传输时会导致运算资源闲置。因此NVIDIA发明了一个很牛逼的技术Zero Copy，它把主机内存直接映射到GPU内存上，在GPU需要数据时直接从主机内存寻找，隐式的传输到GPU中。还有另一个技术叫Pinned Memory，会在产生一个不会被分页的内存，这块内存不会被交换到磁盘的虚拟内存上，内存地址也不会被重新定位，因此，相比普通的Pageable Memory有更高的速度。
Pageanbe 和 Pinned 的区别

对于普通的GPU使用Zero Copy以后，读取的数据速度限制为PCI-E的速度，所以不适用与频繁读取数据的程序，直到 NVIDIA TX-1（TK-1）的出现。TX-1不区分Host Memory和GPU Memory，同一块内存CPU/GPU共享，在这种情况下，Zero Copy会完全节省掉内存传输时间，特别对于流媒体的应用效果显著。

使用Pinned Memory是一定会提高性能的（只要你不分配太多把内存挤爆了就行）；可是Zero Copy在这个帖子中有人提到会导致某种情况的效率低下。

0x00 CUDA

1.标准的CUDA Pipeline：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
// Host Arrays
float* h_in  = new float;
float* h_out = new float;
//Process h_in
// Device arrays
float *d_out, *d_in;
// Allocate memory on the device
cudaMalloc((void *
) &d_in,  sizeIn ));
cudaMalloc((void **) &d_out, sizeOut));
// Copy array contents of input from the host (CPU) to the device (GPU)
cudaMemcpy(d_in, h_in, sizeX * sizeY * sizeof(float), cudaMemcpyHostToDevice);
// Launch the GPU kernel
kernel<<<blocks, threads>>>(d_out, d_in);
// Copy result back
cudaMemcpy(h_out, d_out, sizeOut, cudaMemcpyDeviceToHost);
// Continue processing on host using h_out

零拷贝的CUDA pipeline：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
// Set flag to enable zero copy access
cudaSetDeviceFlags(cudaDeviceMapHost);
// Host Arrays
float* h_in  = NULL;
float* h_out = NULL;
// Process h_in
// Allocate host memory using CUDA allocation calls
cudaHostAlloc((void **)&h_in,  sizeIn,  cudaHostAllocMapped);
cudaHostAlloc((void **)&h_out, sizeOut, cudaHostAllocMapped);
// Device arrays
float *d_out, *d_in;
// Get device pointer from host memory. No allocation or memcpy
cudaHostGetDevicePointer((void **)&d_in,  (void *) h_in , 0);
cudaHostGetDevicePointer((void **)&d_out, (void *) h_out, 0);
// Launch the GPU kernel
kernel<<<blocks, threads>>>(d_out, d_in);
// No need to copy d_out back
// Continue processing on host using h_out

0x01 OpenCV

1. OpenCV 3

OpenCV3可以使用cv::cuda::HostMem来使用ZeroCopy和Pinned Memory，

  • PAGE_LOCKED: sets a page locked memory type used commonly for fast and asynchronous uploading/downloading data from/to GPU.
  • SHARED: specifies a zero copy memory allocation that enables mapping the host memory to GPU address space, if supported.
  • WRITE_COMBINED: sets the write combined buffer that is not cached by CPU. Such buffers are used to supply GPU with data when GPU only reads it. The advantage is a better CPU cache utilization.

详细参见：cv::cuda::HostMem Class Reference

2. OpenCV 2

Regular cv::gpu::GpuMat
cv::gpu::CudaMem with ALLOC_ZEROCOPY

从这抄了一段代码https://github.com/Error323/gpumat-tk1

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
阅读更多内容
    
  








发帖者



时间：










































没有评论:











发表评论






        
      








订阅：
博文评论 (Atom)