Dynamic buffer overflow detection for GPGPUs [CGO 17]

이제 읽어야하는 논문이다.

http://www.computermachines.org/joe/publications/pdfs/cgo2017_clarmor.pdf

저자가 누군지 당연히 모르지만 다 AMD출신이다. AMD에서 GPU도 생산하는지 몰랐다. 

Buffer overflow는 너무나 유명한 security error인데 Dynamic이 왜 붙어있는지 모르겠다.

Abstract

GPU가 buffer overflow를 발생시킨다고 한다. 여기서는 dynamic을 쓰지 않았다.

CPU와 GPU는 다를 memory space를 쓰지 않아 문제가 되지 않았지만 최근 GPU는 CPU와 공유하는 일이 잦고 이에따라 overflow가 문제가 된다.

cpu에서 쓰던 tool을 쓰면 문제가 생김

canary를 쓸 예정

*OpenCL API가 뭔지 모르겠다. 근데 이게 detection이 된다.

당연하지만 optimize 신경을 썼다.

over head가 14퍼센트

1. Introduction

buffer overflow가 뭔지 알지만 정확히 어떻게 attack을 당하는지 모르기에 7번 참고문헌을 언젠가 일어야겠다.

7번 - J.P. Anderson. Computer security technology planning study. technical report ESD-TR-73-51, U.S. Air Force electronic systems division Oct 1972 (공군이라니....)

유명한 security attack - Moris worm, code red, slammer 뭔지 모르겠다. 이게 buffer overflow일으키는 듯

10,12,29,42,45,58,59,60,65,68,74가 모두 overflow막는 방법들이다. 다 못읽을 듯...

GPU 버그가 무시된 이유 1. 다른 memory region, 2. pointer나 function call을 쓰지 않아 발생하기 어려움 3. 밀도가 낫기(?)때문에 필요한 데이터가 침범당할 확률이 적음

PCIe에서 CPU memory와 상호작용함. HSA에서는 virtual memory공유

2. Background 에서 buffer overflow를 설명해준다. GPU의 memory에대해서도 설명해주고 OPEN CL이라는 것도 알려준다.

buffer overflow는 software error. buffer overflow에 대한 자세한 설명은 나와있지 않다.

Miele라는게 나오는데 사람이름인지 회사인지 모르겠다.

heterogeneous : 이기종의 라는데 한국말도 처음듣고 영어도 처음 듣는다.

Margiolas와 O'Boyle가 CPU와 GPU의 DMA transfer를 줄이기 위해 무언가 했다는데 뭔지 모르겠다.

openCL에 대한 설명을 기대했지만 없는것 같아 구글링을 해보니, openCL이란 CPU와 GPU를 이어주는 역할을 하는 인터페이스이다.

http://hoororyn.tistory.com/3

위의 블로그의 도움을 받았다.

openCL에서의 kernel은 그 커널이 아닌, GPU에서 동작하는 함수를 이르는 말이라고 한다.

oepnCL에서 쓰이는 Buffer의 종류도 설명해놨는데 st ack values, local memory, global cl_mem buffers, global cl_mem images, sub-buffers, coarse-grained svm, find-grained svm, fine-grained system svm이 있다. 하나씩 살펴보도록 하자.

stack variables를 메모리에 저장시키지 않는다고 하는데 stack variables란 지역변수를 의미하는 것 같다. 대신 reg로 저장한다고 한다. 이러한 변수를 분석하는 것은 opencl compiler를  변형시켜야한다는데, which is often a proprietary part of a vendor's software stack이 의미하는 건 모르겠다. 그래서 이 연구에서는 stack values의 overflow를 찾아내는데는 쓰지 않는다.

local memory는 kernel이 수행될 때, onchip scratchpad memory에 할당된다고 하는데 무슨말인지 모름. Dram buffer랑은 공유되지 않기에 외부에서 접근하는 것은 GPU kernel을 망가뜨리기 때문에 고려하지 않는다.

global cl_mem buffers 는 host에 의해 clcreatebuffer라는 함수를 이용하여 만들어진다. default상태일 때는 GPU memory에 할당되는데 pointer는 포함하지 않는다. 이 부분이 논문의 목적.

global cl_mem images 이미지는 고차원 buffer인데 clCreateImage2D나 clCreateImage를 이용하여 만든다. 고차원이기에 overflow가 일어날 수 있고 canary를 이용하여 overflow를 찾아낸다.

sub-buffers clcreateSubvuffer를 이용하여 만들어지는데 cl-mem buffer를 참조하며 return값이 cl_mem buffer의 중간을 가르키는 pointer이다. 이 논문에서는 못다루지만 shadow copy라는 것을 만든다고 한다. 이후에 설명된다고 함

coarse-grained SVM  SVM이 ML에 나오는 SVM이 아니라 shared virtual memory라고 한다. 다른 buffer를 가르키는 pointer를 가지는 buffer인데 slSVMAlloc function을 통해 만들어진다. 얘는 cpu memory에 존재하는데, CPU에 접근하고 data를 복사하기 위해서라고 한다. 모든 buffer에 대한 pointer를 가지고 있는듯 하다. 이것도 이 논문의 해당영역이다.

Fine-grained SVM CPU와 GPU buffer의 pointer를 가지고 있다. mapped와 copies가 필요 없다고 하는데 그럼 coarse grained svm은 필요하다는 것인가...amd의 opencl은 gpu가 cpu memory에 접근가능하게 함으로서 가능하다. clSVMAlloc을 함으로써 overflow를 dectect할 수 있다.

Fine-grained system SVM 이름은 전것과 너무나 흡사하다. CPU에 있는 buffer인데 pointer를 저장, 이 논문에서는 다루지 않는다.

electric fence 같은 거는 fine-grained system SVM에 효과가 있다. 이 부분에서 일어나는 GPU overflow는 GPU page fault를 일으키는데 이 논문이 필요하지 않다.

3. Design of a Buffer Overflow Detector

드디어 디자인부분이다. canary value를 이용한다. stackguard,controlpolice, electricfence와 비슷하다고 하는데 하나는 읽어봐야할 것 같다. 주기적으로 canary values가 체크된다. stackGuard와 ContraPolice는 application이나 system lib를 recompile해야한다. electricfence는 virtual memory system을 이용하여 보호하는데 canary region에 덧씌워지면 page fault가 발생한다. malloc function과 연결되어 있어서인데 recompile이 필요하지 않다. 이 논문도 이와 같이 recompile이 필요하지 않다. wrapper가 뭔지 궁금하다....  GPU kernel을 만드는 function call을 실행시킨 후에 canary를 체크한다. kernel이 완료된 후에 체크를 한다. 이게 한계점이 될 수 있다. 왜냐하면 체크되기 전에 어택받은 값이 사용될 수 있기 때문이다. GPU의 control을 가져감으로써 attacker가 우리의 check를 피할 수 있다.(그 값을 다시 써놓는 방식이나 체크하기전에 control을 빼앗음으로서) 그럼에도 유용하다. 

3.1 buffer creation APIs

clSMVAlloc,clCreateBuffer,clCreateImage함수를 이용하여 buffer를 만든다. canary가 있기 때문에 8KB만큼 buffer의 크기가 증가한다. CL_MEM_USE_HOST_PTR은 CPU에서 사용되던 memory를 OpenCL에서 사용하게 해준다. 모든 것을 업데이트 할 수 없기 때문에 resize시키지 못한다. 이를 해결하기 위하여 shadow copy를 만든다. 이 shadow copy도 canary가 있다. kernel이 완료될 때마다 shadow buffer에 있는 data는 original host(CPU) memory에도 저장이 된다. 이 부분을 cache라고 표현한다. 

3.2 setting kernel Arguments

clSetKernelArg와 clSetKernelArgSVMPointer를 이용하여 call을 찾아낸다. wrap이 cl~~~를 실행시킨다 의미인가....

이 list에는 check 해야할 buffer가 있는데 아마 pointer를 이용한 것 같다. list에는(list라는게 결국 array인것인가?)buffer size,canary values, pointer가 있다. pointer가 맞았다.ㅎ optimization을 위해 두개 중 하나는 canary values만 check 한다고 나오는데 나머지는 accessiblity도 check하는 것 같다. accessibility -> cananry values냐 canary values check하냐의 차이인듯

3.3 kernel Enqueues

clEnqueueNDRangeKernel이라는 함수를 쓴다. 이름 참 길다 ㄷㄷ 실제 overflow detection이 일어나는 부분. kernel의 argument를 먼저 분석한다고 하는데 왜인지는 모르겠다. If however there are cl_mem buffers passed to ther kernel, we must verify that these vuffers' canary refions were not perturbed by kernel이 무슨 뜻인지 모르겠는데 3.1에 나온 shadow copy를 만들경우를 고려하는 것 같은데 확실히 모르겠다.(cl_mem buffer가 kernel에 arg로 들어가면 or 사용되면 인것 같다.) kernel이 수행되면 queue에 넣는데 kernel이 끝나고 나면 즉시 접근가능한 canary를 check한다. SVM buffer는 kernel arg만 보고 accessiblity를 알 수 없기에 더 복잡하다. SVMbuffer는 다른 buffer에 대한 pointer를 가지고 있기 때문에 SVM의 canary는 꼭 확인되어야한다.

image buffer를 check 하는 것은 더 복잡하다. 전체적인 overflow가 없더라고 차원내의 간섭이 있을 수 있기 때문에 각 차원마다 canary를 넣어주어야한다. cananry의 size가 dimension과 관련이 있는것은 당연하다. canary를 일차원 array로 읽게된다. 이것은 각 이미지의 끝부분을 가르킨다.

overflow를 감지하면 중단되고 메세지를 출력하도록 했다. kernel name, argument name, 언제 corruption이 일어났는지도출력해준다. argument의 이름은 clGetKernelArginfo를 통해 알아낼 수 있다. 

3.4 API Checking

API가 어떤 것인지 몰라서 무슨 뜻인지 모르겠다. 

4. Accelerating Buffer Overflow Detection