Geekly Articles each Day

पासवर्ड हैकिंग के साथ ओपनसीएल की खोज

प्रविष्टि


हाल ही में, GPGPU के बारे में विभिन्न लेखों और प्रस्तुतियों को पढ़ने के बाद, मैंने अपने लिए भी वीडियो कार्ड के लिए प्रोग्रामिंग का प्रयास करने का निर्णय लिया। वास्तव में, इस क्षेत्र में प्रौद्योगिकियों का विकल्प महान नहीं है - केवल सीयूडीए (एनवीडिया मालिकाना मानक) और ओपनसीएल (मुफ्त मानक, एटीआई, एनवीडिया से जीपीयू पर काम करता है, साथ ही केंद्रीय प्रोसेसर पर) अभी भी जीवित और विकासशील है। इस तथ्य के कारण कि मेरे लैपटॉप में एक एटीआई ग्राफिक्स कार्ड (मोबिलिटी राडर्डन 5650 एचडी) है, यह विकल्प पूरी तरह से एक विकल्प में कम हो गया था - ओपनसीएल। यह लेख स्क्रैच से ओपनसीएल सीखने की प्रक्रिया पर चर्चा करेगा, साथ ही इसके बारे में क्या आया।

OpenCL और PyOpenCl का अवलोकन


पहली नज़र में, सब कुछ मेरे लिए बहुत ही भ्रामक लग रहा था, दोनों सी में नियंत्रण कोड और तथाकथित गुठली के कोड - गुठली। प्रदान की गई सी एपीआई में, यहां तक ​​कि एक साधारण कार्यक्रम का शुभारंभ भी बड़ी संख्या में लाइनें लेता है, खासकर कम से कम कुछ त्रुटियों के प्रसंस्करण के साथ, इसलिए मैं कुछ अधिक सुविधाजनक और मानवीय खोजना चाहता था। पसंद PyOpenCL पुस्तकालय पर गिर गया, जिसके नाम से यह स्पष्ट है कि नियंत्रण कोड पायथन में लिखा गया है। इसमें सब कुछ अधिक समझ में आता है, यहां तक ​​कि पहली बार ओपनसीएल कोड को देखने वाले किसी व्यक्ति के लिए (बेशक, यह केवल सरल उदाहरणों पर लागू होता है)। हालांकि, खुद गुठली का कोड अभी भी थोड़ा संशोधित सी में लिखा गया है, इसलिए आपको अभी भी इसका अध्ययन करना होगा। इस पर पूर्ण प्रलेखन मानक डेवलपर की वेबसाइट ( क्रोनोस ) पर प्राप्त किया जा सकता है, और विशिष्ट कार्यान्वयन पर जानकारी - क्रमशः एटीआई और एनवीडिया वेबसाइटों पर।
आप सरल उदाहरण (दो सरणियों को जोड़कर) का उपयोग करके भाषा की पहली छाप प्राप्त कर सकते हैं:

__kernel void sum(__global const float *a, __global const float *b, __global float *c) { int gid = get_global_id(0); c[gid] = a[gid] + b[gid]; }


और यहाँ इस उदाहरण को चलाने और सत्यापित करने के लिए पूर्ण आवश्यक कोड है (PyOpenCL प्रलेखन से लिया गया है):

छिपा हुआ पाठ
 import pyopencl as cl import numpy import numpy.linalg as la a = numpy.random.rand(50000).astype(numpy.float32) b = numpy.random.rand(50000).astype(numpy.float32) ctx = cl.create_some_context() queue = cl.CommandQueue(ctx) mf = cl.mem_flags a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a) b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b) dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes) prg = cl.Program(ctx, """ __kernel void sum(__global const float *a, __global const float *b, __global float *c) { int gid = get_global_id(0); c[gid] = a[gid] + b[gid]; } """).build() prg.sum(queue, a.shape, None, a_buf, b_buf, dest_buf) a_plus_b = numpy.empty_like(a) cl.enqueue_copy(queue, a_plus_b, dest_buf) print la.norm(a_plus_b - (a+b))



विशिष्ट लाइनें तुरंत दिखाई देती हैं: संदर्भ बनाना, निष्पादन कतारें, डिवाइस को बफ़र्स बनाना और कॉपी करना, साथ ही साथ कर्नेल को स्वयं संकलित करना और शुरू करना। आप प्रलेखन में OpenCL में संदर्भों और कतारों के बारे में अधिक पढ़ सकते हैं, और अपेक्षाकृत सरल कार्यक्रमों के लिए, आपको केवल एक कतार और एक संदर्भ की आवश्यकता होगी, जो उदाहरण में बहुत समान रेखाओं द्वारा बनाई जाएगी। सामान्य तौर पर, अक्सर ओपनसीएल कार्यक्रमों में गणना की संरचना कुछ इस तरह दिखाई देती है:



SHA1 हैशिंग


यह एक स्तर नीचे जाने और यह पता लगाने का समय है कि कर्नेल कोड कैसे काम करता है। किसके लिए, ओपनसीएल फ़ंक्शन को बाहर से चलाने के लिए, यह __k कर्नेल विशेषता द्वारा इंगित किया जाना चाहिए, एक शून्य मान प्रकार और कुछ निश्चित तर्क हैं, जो सीधे मान (int, float4, ...) या क्षेत्रों को इंगित कर सकते हैं। मेमोरी __global, __constant, __local। इसके अलावा, सुविधा के लिए, कर्नेल से कॉल किए जा सकने वाले अन्य कार्यों को प्रोग्राम में घोषित किया जा सकता है, और यह प्रदर्शन को प्रभावित नहीं करता है: सभी फ़ंक्शन स्वचालित रूप से प्रतिस्थापित किए जाते हैं (अर्थात, इनलाइन निर्देश के साथ)। इसका कारण यह है कि ओपनसीएल में पुनरावृत्ति का समर्थन नहीं किया जाता है।

इस तथ्य का उपयोग करते हुए कि ओपनसीएल एक संशोधित सी भाषा है, आप हैश फ़ंक्शन का एक तैयार-किए गए कार्यान्वयन को ले सकते हैं, उदाहरण के लिए, SHA1, और मामूली संशोधनों का उपयोग करें:

छिपा हुआ पाठ
 #define K1 0x5A827999 #define K2 0x6ED9EBA1 #define K3 0x8F1BBCDC #define K4 0xCA62C1D6 #define a0 0x67452301; #define b0 0xEFCDAB89; #define c0 0x98BADCFE; #define d0 0x10325476; #define e0 0xC3D2E1F0; #define f1(x,y,z) ( z ^ ( x & ( y ^ z ) ) ) /* Rounds 0-19 */ #define f2(x,y,z) ( x ^ y ^ z ) /* Rounds 20-39 */ #define f3(x,y,z) ( ( x & y ) | ( z & ( x | y ) ) ) /* Rounds 40-59 */ #define f4(x,y,z) ( x ^ y ^ z ) /* Rounds 60-79 */ #define ROTL(n,X) ( ( ( X ) << n ) | ( ( X ) >> ( 32 - n ) ) ) #define expand(W,i) ( W[ i & 15 ] = ROTL( 1, ( W[ i & 15 ] ^ W[ (i - 14) & 15 ] ^ \ W[ (i - 8) & 15 ] ^ W[ (i - 3) & 15 ] ) ) ) #define subRound(a, b, c, d, e, f, k, data) \ ( e += ROTL( 5, a ) + f( b, c, d ) + k + data, b = ROTL( 30, b ) ) #define REVERSE(value) value = ((value & 0xFF000000) >> 24) | ((value & 0x00FF0000) >> 8) | ((value & 0x0000FF00) << 8) | ((value & 0x000000FF) << 24) long sha1(uint *eData, const int length) { unsigned int A = a0; unsigned int B = b0; unsigned int C = c0; unsigned int D = d0; unsigned int E = e0; ((__local char *)eData)[length] = 0x80; for (int i = 0; i <= length / 4; i++) { REVERSE(eData[i]); } eData[14] = 0; eData[15] = length * 8; subRound( A, B, C, D, E, f1, K1, eData[ 0 ] ); subRound( E, A, B, C, D, f1, K1, eData[ 1 ] ); subRound( D, E, A, B, C, f1, K1, eData[ 2 ] ); subRound( C, D, E, A, B, f1, K1, eData[ 3 ] ); subRound( B, C, D, E, A, f1, K1, eData[ 4 ] ); subRound( A, B, C, D, E, f1, K1, eData[ 5 ] ); subRound( E, A, B, C, D, f1, K1, eData[ 6 ] ); subRound( D, E, A, B, C, f1, K1, eData[ 7 ] ); subRound( C, D, E, A, B, f1, K1, eData[ 8 ] ); subRound( B, C, D, E, A, f1, K1, eData[ 9 ] ); subRound( A, B, C, D, E, f1, K1, eData[ 10 ] ); subRound( E, A, B, C, D, f1, K1, eData[ 11 ] ); subRound( D, E, A, B, C, f1, K1, eData[ 12 ] ); subRound( C, D, E, A, B, f1, K1, eData[ 13 ] ); subRound( B, C, D, E, A, f1, K1, eData[ 14 ] ); subRound( A, B, C, D, E, f1, K1, eData[ 15 ] ); subRound( E, A, B, C, D, f1, K1, expand( eData, 16 ) ); subRound( D, E, A, B, C, f1, K1, expand( eData, 17 ) ); subRound( C, D, E, A, B, f1, K1, expand( eData, 18 ) ); subRound( B, C, D, E, A, f1, K1, expand( eData, 19 ) ); subRound( A, B, C, D, E, f2, K2, expand( eData, 20 ) ); subRound( E, A, B, C, D, f2, K2, expand( eData, 21 ) ); subRound( D, E, A, B, C, f2, K2, expand( eData, 22 ) ); subRound( C, D, E, A, B, f2, K2, expand( eData, 23 ) ); subRound( B, C, D, E, A, f2, K2, expand( eData, 24 ) ); subRound( A, B, C, D, E, f2, K2, expand( eData, 25 ) ); subRound( E, A, B, C, D, f2, K2, expand( eData, 26 ) ); subRound( D, E, A, B, C, f2, K2, expand( eData, 27 ) ); subRound( C, D, E, A, B, f2, K2, expand( eData, 28 ) ); subRound( B, C, D, E, A, f2, K2, expand( eData, 29 ) ); subRound( A, B, C, D, E, f2, K2, expand( eData, 30 ) ); subRound( E, A, B, C, D, f2, K2, expand( eData, 31 ) ); subRound( D, E, A, B, C, f2, K2, expand( eData, 32 ) ); subRound( C, D, E, A, B, f2, K2, expand( eData, 33 ) ); subRound( B, C, D, E, A, f2, K2, expand( eData, 34 ) ); subRound( A, B, C, D, E, f2, K2, expand( eData, 35 ) ); subRound( E, A, B, C, D, f2, K2, expand( eData, 36 ) ); subRound( D, E, A, B, C, f2, K2, expand( eData, 37 ) ); subRound( C, D, E, A, B, f2, K2, expand( eData, 38 ) ); subRound( B, C, D, E, A, f2, K2, expand( eData, 39 ) ); subRound( A, B, C, D, E, f3, K3, expand( eData, 40 ) ); subRound( E, A, B, C, D, f3, K3, expand( eData, 41 ) ); subRound( D, E, A, B, C, f3, K3, expand( eData, 42 ) ); subRound( C, D, E, A, B, f3, K3, expand( eData, 43 ) ); subRound( B, C, D, E, A, f3, K3, expand( eData, 44 ) ); subRound( A, B, C, D, E, f3, K3, expand( eData, 45 ) ); subRound( E, A, B, C, D, f3, K3, expand( eData, 46 ) ); subRound( D, E, A, B, C, f3, K3, expand( eData, 47 ) ); subRound( C, D, E, A, B, f3, K3, expand( eData, 48 ) ); subRound( B, C, D, E, A, f3, K3, expand( eData, 49 ) ); subRound( A, B, C, D, E, f3, K3, expand( eData, 50 ) ); subRound( E, A, B, C, D, f3, K3, expand( eData, 51 ) ); subRound( D, E, A, B, C, f3, K3, expand( eData, 52 ) ); subRound( C, D, E, A, B, f3, K3, expand( eData, 53 ) ); subRound( B, C, D, E, A, f3, K3, expand( eData, 54 ) ); subRound( A, B, C, D, E, f3, K3, expand( eData, 55 ) ); subRound( E, A, B, C, D, f3, K3, expand( eData, 56 ) ); subRound( D, E, A, B, C, f3, K3, expand( eData, 57 ) ); subRound( C, D, E, A, B, f3, K3, expand( eData, 58 ) ); subRound( B, C, D, E, A, f3, K3, expand( eData, 59 ) ); subRound( A, B, C, D, E, f4, K4, expand( eData, 60 ) ); subRound( E, A, B, C, D, f4, K4, expand( eData, 61 ) ); subRound( D, E, A, B, C, f4, K4, expand( eData, 62 ) ); subRound( C, D, E, A, B, f4, K4, expand( eData, 63 ) ); subRound( B, C, D, E, A, f4, K4, expand( eData, 64 ) ); subRound( A, B, C, D, E, f4, K4, expand( eData, 65 ) ); subRound( E, A, B, C, D, f4, K4, expand( eData, 66 ) ); subRound( D, E, A, B, C, f4, K4, expand( eData, 67 ) ); subRound( C, D, E, A, B, f4, K4, expand( eData, 68 ) ); subRound( B, C, D, E, A, f4, K4, expand( eData, 69 ) ); subRound( A, B, C, D, E, f4, K4, expand( eData, 70 ) ); subRound( E, A, B, C, D, f4, K4, expand( eData, 71 ) ); subRound( D, E, A, B, C, f4, K4, expand( eData, 72 ) ); subRound( C, D, E, A, B, f4, K4, expand( eData, 73 ) ); subRound( B, C, D, E, A, f4, K4, expand( eData, 74 ) ); subRound( A, B, C, D, E, f4, K4, expand( eData, 75 ) ); subRound( E, A, B, C, D, f4, K4, expand( eData, 76 ) ); subRound( D, E, A, B, C, f4, K4, expand( eData, 77 ) ); subRound( C, D, E, A, B, f4, K4, expand( eData, 78 ) ); subRound( B, C, D, E, A, f4, K4, expand( eData, 79 ) ); A += a0; B += b0; C += c0; D += d0; E += e0; return as_ulong((uint2)(D, E)); }



यहां कुछ स्पष्टीकरण की आवश्यकता है। क्रैकिंग के लिए "प्रायोगिक" पासवर्ड हैश के रूप में, मैंने लिंक्डइन हैश लीक किया, जो लगभग 6 मिलियन (अद्वितीय) हैं। सूची में उपलब्धता के लिए काफी त्वरित जांच के लिए कई विकल्प हैं, मैंने हैश टेबल का उपयोग किया (उस पर बाद में)। मेमोरी की खपत को कम करने के लिए और SHA1 के पूरे 20 बाइट्स को स्टोर करने के लिए विचार को गति न दें, लेकिन अंतिम 8 बाइट्स, अर्थात्। एक लंबी / लंबी कीमत। बेशक, इससे झूठे मेल की संभावना बढ़ जाती है, लेकिन यह बहुत छोटा है: मेरे पास मौजूद सभी पासवर्डों में से, मेरे पास केवल 6 ऐसे मामले थे, जो बिल्कुल भी महत्वपूर्ण नहीं है। इसलिए, फटा हुआ मान (अंतिम 8 बाइट्स) तुरंत उपरोक्त फ़ंक्शन से वापस आ गया है। अन्यथा, सब कुछ मानक है, SHA1 एल्गोरिथ्म 56 बाइट्स से कम तार के मामले के लिए विनिर्देश के अनुसार सीधे लागू किया जाता है।

खोज का संगठन


अगला, आपको स्वयं खोज को व्यवस्थित करने की आवश्यकता है। सबसे सरल विकल्प जानवर बल है, सभी वर्णों के एक ही सेट के लिए, और यह सीधे लागू किया जा सकता है, उदाहरण के लिए, एक समान स्थिति में:

छिपा हुआ पाठ
 __kernel void do_brute( __global const long *table, const ulong start_num, __global ulong *result, __global uint *res_ind) { char s[64]; uint *s_l = (__local uint *)s; int i, j; ulong _n, n; ulong start = start_num + get_global_id(0) * HASHES_PER_WORKITEM; for (j = 0; j < HASHES_PER_WORKITEM; j++) { n = _n = j + start; for (i = 15; i >= 0; i--) { s_l[i] = 0; } for (i = COMB_LEN - 1; i >= 0; i--) { s[i] = charset[n % CHARS_CNT]; n /= CHARS_CNT; } if (check_in_table(table, sha1(s_l, COMB_LEN))) { result[atomic_inc(res_ind)] = _n; } } }


यहां HASHES_PER_WORKITEM एक कार्य आइटम (स्ट्रीम) द्वारा एक रन में संसाधित हैश की संख्या है, COMB_LEN संयोजन की लंबाई है, चारसेट वर्णों की एक सरणी है, CHARS_CNT सरणी में वर्णों की संख्या है। जैसा कि आप देख सकते हैं, स्टार्टअप पर, पॉइंटर को हैश टेबल पर, पासवर्ड नंबर जिसमें से खोज शुरू होती है, और परिणाम को प्रदर्शित करने के लिए सरणी के लिए एक पॉइंटर और उसमें एक इंडेक्स भी इस कर्नेल को दिया जाता है।

ओपनसीएल में, एक समय में एक धागा लॉन्च नहीं किया जाता है, लेकिन उनमें से कुछ को वैश्विक कार्य आकार कहा जाता है, और सभी धागे एक ही तर्क प्राप्त करते हैं। उनमें से प्रत्येक के लिए प्रमुख स्थान के अपने हिस्से को क्रमबद्ध करने के लिए, स्ट्रिंग ulong start = start_num + get_global_id(0) * HASHES_PER_WORKITEM; एक विशिष्ट थ्रेड के लिए संख्या की गणना करता है (get_global_id (0) एक मानक फ़ंक्शन है जो थ्रेड इंडेक्स को 0 से वर्तमान वैश्विक कार्य आकार पर लौटाता है)।

अगला, प्रत्येक स्ट्रीम HASHES_PER_WORKITEM पासवर्ड की पुष्टि करता है, जिनमें से प्रत्येक को sha1 फ़ंक्शन द्वारा हैश किया जाता है और चेक_इन_टेबल फ़ंक्शन की उपस्थिति के लिए जाँच की जाती है, जिसे बाद में वर्णित किया गया है। इस रूप में और हैश टेबल के सबसे सरल कार्यान्वयन के साथ, मुझे प्रति सेकंड लगभग 20 मिलियन पासवर्ड का परिणाम मिला, जो स्पष्ट रूप से तुलना में पर्याप्त नहीं है, उदाहरण के लिए, ओक्लाहस्कैट, जो मेरे लैपटॉप पर 300 मिलियन से अधिक का उत्पादन करता है (यहां तक ​​कि सत्यापन हैश की एक बड़ी सूची से गुजरता है) यह उचित नहीं है)। आगे देखते हुए, मैं कहता हूं कि साधारण ब्रूट फोर्स के लिए मैं 160 मिलियन प्रति सेकंड की गति हासिल करने में कामयाब रहा, जो कि ओक्लासैट (एक हैश के साथ) की आधी से अधिक गति है।

हैश टेबल


तो, हैश के लिए जाँच कर रहा है। लागू किया गया पहला विकल्प ओपन एड्रेसिंग के साथ सबसे सरल हैश तालिका थी । इसे भर दिया गया था, इसलिए मामलों को जटिल नहीं करने के लिए, प्रोसेसर के साथ और वीडियो कार्ड के साथ नहीं; ओपेनसीएल में, केवल इसके लिए अनुरोध थे। यह मामला कुछ इस तरह था:

छिपा हुआ पाठ
 bool check_in_table( __global const long *table, const long value) { uint index = calc_index(value); uint step = calc_step(value); for (uint a = 1; ; a++) { index %= TABLE_SIZE; if (table[index] == 0) { return false; } if (table[index] == value) { return true; } index += a * step + 1; } }


मैंने विभिन्न तालिका आकार और जांच विधियों की कोशिश की, लेकिन गति में उल्लेखनीय सुधार नहीं हुआ। GPU हैश टेबल के बारे में सामग्री की तलाश में, मैं एक लेख "GPU पर एक कुशल हैश बिल्डिंग का निर्माण" वसीली वोल्कोव द्वारा प्रस्तुत किया गया है, जिसमें एक निश्चित कोयल हैशटेबल का उल्लेख है (मुझे नहीं पता कि रूसी में एक स्थापित अनुवाद है), जो मैंने पहले किया था नहीं सुना। संक्षेप में, इसका सार एक के बजाय कई हैश फ़ंक्शन के उपयोग में है और भरने का एक विशेष तरीका है, जिसके बाद तत्व k मेमोरी एक्सेस से अधिक नहीं के लिए पाया जाता है, जहां k हैश फ़ंक्शन की संख्या है। चूंकि कब्जे वाली मेमोरी की तुलना में गति मेरे लिए अधिक महत्वपूर्ण है, मैंने k = 2 का उपयोग किया। इसे भरना भी सीपीयू पर होता है।

निष्कर्ष


इसके अलावा, निश्चित रूप से, अनुकूलन ने कोड के दूसरे भाग को प्रभावित किया, अर्थात् पासवर्ड की पीढ़ी। उपरोक्त संस्करण में, कई गैर-इष्टतम स्थान तुरंत दिखाई देते हैं, उदाहरण के लिए, प्रत्येक अगला पासवर्ड खरोंच से उत्पन्न होता है, हालांकि आप पिछले एक को बदल सकते हैं। अनुकूलन के लिए अन्य स्थान हैं, विशेष रूप से, ओपनसीएल के लिए विशिष्ट: एक बहुत तेजी से स्थानीय एक के बजाय पात्रों की एक सरणी के लिए वैश्विक या निरंतर मेमोरी का उपयोग करना (विशिष्ट कार्यान्वयन के डेवलपर से सीधे मेमोरी क्षेत्रों के बारे में अधिक पढ़ना बेहतर है)। हालाँकि, यह कर्नेल कोड के विभिन्न अनुकूलन के बारे में एक अलग लेख लिखने के लायक है, लेकिन यहाँ मैं यह कहूँगा कि जब GPU के लिए प्रोग्रामिंग विभिन्न विकल्पों की कोशिश करने और उनकी गति को देखने के लायक है, क्योंकि यह हमेशा आंख से कहना संभव नहीं है कि यह तेजी से काम करेगा। कभी-कभी एक निश्चित निर्देश को हटाने से निष्पादन धीमा हो जाता है, और महत्वपूर्ण रूप से।

भविष्य में, मैंने अलग-अलग पदों के लिए अलग-अलग वर्णमाला के लिए समर्थन जोड़ा, साथ ही साथ शब्दों से अधिक अक्षर, और न केवल व्यक्तिगत वर्ण। सुविधा और लचीलेपन के लिए, कर्नेल कोड को माको टेम्पलेट इंजन द्वारा संसाधित किया जाता है। यह सब संग्रह में है (नीचे देखें)।

निष्कर्ष


इसलिए, आखिरकार मैंने क्या हासिल किया:


मूर्त परिणाम:


छापें:

पुनश्च:


विभिन्न टिप्पणियों और सिफारिशों का स्वागत है।

Source: https://habr.com/ru/post/In150289/

More articles:


All Articles