Tweet
Tweet

Slide 1

Slide 1 text

Compiler Fuzzing through Deep Learning https://chriscummins.cc/issta18

Slide 2

Slide 2 text

Chris Cummins Codeplay Software University of Edinburgh Pavlos Petoumenos Alastair Murray Hugh Leather University of Edinburgh University of Edinburgh

Slide 3

Slide 3 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... clang6.0 fuzzing a compiler circa [McKeenan98] a.out error: use of undeclared … Timeout Crash

Slide 4

Slide 4 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... clang6.0 fuzzing a compiler circa [McKeenan98] a.out error: use of undeclared … Timeout Bug! Crash

Slide 5

Slide 5 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... clang6.0 fuzzing a compiler circa [McKeenan98] a.out error: use of undeclared … Timeout Bug?? Bug! Crash

Slide 6

Slide 6 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... gcc5.5 differential testing compilers circa [McKeenan98] a.out clang6.0 clang3.6 a.out a.out

Slide 7

Slide 7 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... gcc5.5 differential testing compilers circa [McKeenan98] a.out clang6.0 clang3.6 a.out a.out $ ./a.out 42 $ ./a.out 42 $ ./a.out -14522312

Slide 8

Slide 8 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... gcc5.5 differential testing compilers circa [McKeenan98] a.out clang6.0 clang3.6 a.out a.out $ ./a.out 42 $ ./a.out 42 $ ./a.out -14522312 Majority rules

Slide 9

Slide 9 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... gcc5.5 differential testing compilers circa [McKeenan98] clang6.0 clang3.6 Also works for build failures

Slide 10

Slide 10 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... gcc5.5 differential testing compilers circa [McKeenan98] clang6.0 clang3.6 error: use of undeclared … error: use of undeclared … a.out Also works for build failures

Slide 11

Slide 11 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... gcc5.5 differential testing compilers circa [McKeenan98] clang6.0 clang3.6 error: use of undeclared … error: use of undeclared … a.out Also works for build failures hard to generate!

Slide 12

Slide 12 text

an ideal fuzzer 1. Cheap
 Easy to implement and extend
 (Languages and features grow quickly) 2. Interpretable Testcases
 Necessary for triage
 (i.e. 45 lines or less [Sun2016]) 3. Plausible Output
 Representative of handwritten code
 (So that bugs gets fixed)

Slide 13

Slide 13 text

state-of-the-art: CLSmith Random grammar enumeration. Extensive static analyses support subset of OpenCL features. Targets compiler middle ends. Incredibly effective!
 100s of bugs to date. https://github.com/ChrisLidbury/CLSmith #include "CLSmith.h" struct S0 { int32_t g_4[4][10]; ... }; kernel void A(global ulong *r) { int i, j, k; struct S0 c_1856; struct S0* p_1855 = &c_1856; c_1856 = c_1857; func_1(p_1855); barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); for (i = 0; i < 4; i++) for (j = 0; j < 10; j++) …>g_4[i][j], "p_1855->g_4[i][j]", print_hash_value); result[get_linear_global_id()] = crc64_context ^ 0xFFFFFFFFFFFFFFFFUL; }

Slide 14

Slide 14 text

state-of-the-art: CLSmith https://github.com/ChrisLidbury/CLSmith 1. Cheap ✖ nope!
 Years to develop! 50k lines of C++.
 Each PL feature engineered by hand. 2. Interpretable Testcases ✖ nope!
 Avg. 1200 lines (excluding headers).
 Requires reduction: ~4 hours / test. 3. Plausible Output ✖ nope!
 Unusual and restricted combinations of PL features.
 87 dials control “shape” of output - hand tuned.

Slide 15

Slide 15 text

Automatic inference of fuzzers from examples. 102x less code than state-of-art. Similar bug finding power, simpler test cases. contributions

Slide 16

Slide 16 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign

Slide 17

Slide 17 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign Mined from • 1k repos • 10k files • 2.0M LOC Filtered by oracle compiler.

Slide 18

Slide 18 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign

Slide 19

Slide 19 text

vocabulary encoding #define MY_CONST 3.14 // A very simple kernel. kernel void Foo(global float* input, const float x) { input[get_global_id(0)] *= MY_CONST + x; } kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 1. Preprocess. 2. Enforce code style. 3. Rename variables. 4. Rename functions.

Slide 20

Slide 20 text

vocabulary encoding kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } Vocab: Encoded:

Slide 21

Slide 21 text

vocabulary encoding Token Index kernel 0 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 Vocab: Encoded:

Slide 22

Slide 22 text

vocabulary encoding Token Index kernel 0 [space] 1 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 Vocab: Encoded:

Slide 23

Slide 23 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 Vocab: Encoded:

Slide 24

Slide 24 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 Vocab: Encoded: 1

Slide 25

Slide 25 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 A 3 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 1 3 Vocab: Encoded:

Slide 26

Slide 26 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 A 3 ( 4 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 1 3 4 Vocab: Encoded:

Slide 27

Slide 27 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 A 3 ( 4 global 5 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 1 3 4 5 Vocab: Encoded:

Slide 28

Slide 28 text

vocabulary encoding kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 1 3 4 5 1 Vocab: Encoded: Token Index kernel 0 [space] 1 void 2 A 3 ( 4 global 5

Slide 29

Slide 29 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 A 3 ( 4 global 5 float 6 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 1 3 4 5 1 6 Vocab: Encoded:

Slide 30

Slide 30 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 A 3 ( 4 global 5 float 6 * 7 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 1 3 4 5 1 6 7 Vocab: Encoded:

Slide 31

Slide 31 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 A 3 ( 4 global 5 float 6 * 7 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 1 3 4 5 1 6 7 Vocab: Encoded: 1

Slide 32

Slide 32 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 A 3 ( 4 global 5 float 6 * 7 a 8 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 1 3 4 5 1 6 7 1 8 Vocab: Encoded:

Slide 33

Slide 33 text

vocabulary encoding Token Index kernel 0 [space] 1 void 2 A 3 ( 4 global 5 float 6 * 7 a 8 Token Index ] 18 = 19 3 20 . 21 1 22 4 23 + 24 ; 25 Token Index , 9 const 10 b 11 ) 12 { 13 \n 14 [ 15 get_global_id 16 0 17 kernel void A(global float* a, const float b) { a[get_global_id(0)] *= 3.14 + b; } 0 1 2 1 3 4 5 1 6 7 1 8 ... Vocab: Encoded:

Slide 34

Slide 34 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign

Slide 35

Slide 35 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign

Slide 36

Slide 36 text

neural network 181 tokens 181 tokens Input: 30M token corpus
 Learns probability distribution over corpus. < 500 lines of code, 12 hours training on GPU. 512x2 LSTM 0 1 2 ...

Slide 37

Slide 37 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign

Slide 38

Slide 38 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign

Slide 39

Slide 39 text

synthesizer + harness 0 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void … Input:

Slide 40

Slide 40 text

synthesizer + harness 0 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void … Input:

Slide 41

Slide 41 text

synthesizer + harness 0 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A … Input:

Slide 42

Slide 42 text

synthesizer + harness 0 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A … Input: 3

Slide 43

Slide 43 text

synthesizer + harness 0 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A … Input: 3

Slide 44

Slide 44 text

synthesizer + harness 0 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A( … Input: 3

Slide 45

Slide 45 text

synthesizer + harness 0 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A( … Input: 3 4

Slide 46

Slide 46 text

synthesizer + harness 0 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A( … Input: 3 4

Slide 47

Slide 47 text

synthesizer + harness 0 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global … Input: 3 4 kernel void A(global 0 1 kernel ' ' void A ( global int double float

Slide 48

Slide 48 text

synthesizer + harness 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global … Input: 3 4 … 5

Slide 49

Slide 49 text

synthesizer + harness 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global … Input: 3 4 … 5

Slide 50

Slide 50 text

synthesizer + harness 1 2 1 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global … Input: 3 4 … 5

Slide 51

Slide 51 text

synthesizer + harness 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global … Input: … 1 2 1 3 4 5

Slide 52

Slide 52 text

synthesizer + harness 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global … Input: … 1 2 1 3 4 5 1

Slide 53

Slide 53 text

synthesizer + harness 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global … Input: … 1 2 1 3 4 5 1

Slide 54

Slide 54 text

synthesizer + harness 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global int … Input: … 1 2 1 3 4 5 1

Slide 55

Slide 55 text

synthesizer + harness 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global int … Input: … 1 2 1 3 4 5 6 3 4 5 1

Slide 56

Slide 56 text

synthesizer + harness 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. Decoded: Output: 0 1 kernel ' ' void A ( global int double float kernel void A(global int* a) { /* snip */ } … Input: … 1 2 1 3 4 5 1 1 2 1 3 4 5 6 3 4 5 1

Slide 57

Slide 57 text

synthesizer + harness 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. 3. Can we parse signature? Yes: Generate input data, compile and run it. No: Compile it but don’t run it. Decoded: kernel void A(global int* a) { /* snip */ }

Slide 58

Slide 58 text

synthesizer + harness 1. Seed the model with the start of a program. 2. Predict tokens until { } brackets balance. 3. Can we parse signature? Yes: Generate input data, compile and run it. No: Compile it but don’t run it. Decoded: kernel void A(global int* a) { /* snip */ }

Slide 59

Slide 59 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign

Slide 60

Slide 60 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign Standard majority voting. False-positive filtering of runtime behavior: Combination of off-the-shelf tools and ad-hoc filters. Took ~1 dev-day to develop.

Slide 61

Slide 61 text

int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Training Corpus Vocabulary Encoding Neural Network Synthesizer our approach Harness int main( int argc, char** argv) { ... int main( int argc, char** argv) { ... Generated Test Cases Voting Heuristics Testing Campaign

Slide 62

Slide 62 text

how well does it work?

Slide 63

Slide 63 text

48 hours per testbed testing campaign 10 OpenCL compilers 3 GPUs, 5 CPUs, Xeon Phi, Emulator Test with optimizations on / off Treat as separate testbeds

Slide 64

Slide 64 text

Errors in ever compiler! Num results (log) 1 100 10,000 C om piler crash B uild Tim eout B uild Failure Program C rash W rong O utput 69 252 51 860 7,040 results overview

Slide 65

Slide 65 text

… crashes during parsing / compilation 67 bug reports to date… void A() {void* a; uint4 b=0; b=(b>b)?a:a } Affects: Intel OpenCL SDK 1.2.0.25 kernel void A(global int* a) { int b = get_global_id(0); a[b] = (6 * 32) + 4 * (32 / 32) + a; } Affects: Beignet 1.3 “Bad code” finds bugs in error handling

Slide 66

Slide 66 text

kernel void A() { __builtin_astype(d, uint4); } … crashes during type checking 67 bug reports to date… Affects: 6 / 10 compilers we tested Unexpected outcome: Learning from handwritten code leads to bugs found in compiler builtins!

Slide 67

Slide 67 text

kernel void A(global double* a, global double* b, global double* c, int d, int e) { double f;
 int g = get_global_id(0); if (g < e - d - 1) c[g] = (((e) / d) % 5) % (e + d); } 67 bug reports to date… Affects: Intel OpenCL SDK 1.2.0.25 … errors in optimizers CLSmith doesn’t allow
 thread-dependent control flow.

Slide 68

Slide 68 text

$ docker run chriscummins/opencl_fuzz https://chriscummins.cc/issta18 runs in docker try it for yourself! code and paper on GitHub

Slide 69

Slide 69 text

comparison to CLSmith LLVM robustness experiments extending to a second language
 + more! read the paper for … https://chriscummins.cc/issta18 code and paper on GitHub

Slide 70

Slide 70 text

Compiler Fuzzing through Deep Learning https://chriscummins.cc/issta18 Problem: Compilers inputs are hard to generate Use DL to infer PL generator from examples 102x less code than state-of-art, 3.03x faster Lots of bugs!