1 /*
2 * JtR format to crack "AS-REP" messages.
3 *
4 * This software is
5 * Copyright (c) 2017 Dhiru Kholia (dhiru at openwall.com),
6 * Copyright (c) 2018 magnum,
7 * and it is hereby released to the general public under the following terms:
8 *
9 * Redistribution and use in source and binary forms, with or without
10 * modification, are permitted.
11 *
12 * This file is based on krb5_asrep_fmt_plug.c and opencl_krb5pa-sha1_fmt_plug.c
13 * files.
14 */
15
16 #ifdef HAVE_OPENCL
17
18 #if FMT_EXTERNS_H
19 extern struct fmt_main fmt_opencl_krb5_asrep_aes;
20 #elif FMT_REGISTERS_H
21 john_register_one(&fmt_opencl_krb5_asrep_aes);
22 #else
23
24 #include "arch.h"
25 #include "misc.h"
26 #include "formats.h"
27 #include "options.h"
28 #include "common.h"
29 #include "config.h"
30 #include "aes.h"
31 #include "krb5_common.h"
32 #include "krb5_asrep_common.h"
33 #include "opencl_common.h"
34 #define MAX_OUTLEN 32
35 #include "opencl_pbkdf2_hmac_sha1.h"
36 #include "hmac_sha.h"
37
38 #define FORMAT_LABEL "krb5asrep-aes-opencl"
39 #define FORMAT_NAME "Kerberos 5 AS-REP etype 17/18"
40 #define ALGORITHM_NAME "PBKDF2-SHA1 OpenCL"
41 #define BENCHMARK_COMMENT ""
42 #define BENCHMARK_LENGTH 0x107
43 #define BINARY_SIZE 0
44 #define BINARY_ALIGN 1
45 #define SALT_SIZE sizeof(struct custom_salt *)
46 #define SALT_ALIGN sizeof(struct custom_salt *)
47
48 #define MIN_KEYS_PER_CRYPT 1
49 #define MAX_KEYS_PER_CRYPT 1
50
51 #define GETPOS(i, index) (((index) % ocl_v_width) * 4 + ((i) & ~3U) * ocl_v_width + (((i) & 3) ^ 3) + ((index) / ocl_v_width) * 64 * ocl_v_width)
52
53 static struct fmt_tests tests[] = {
54 // AS-REP-with-PA-unsupported-openwall.pcap
55 {"$krb5asrep$18$EXAMPLE.COMlulu$b49aa3de9314e2d8daafe323f2e84b9a4ddc361d99bf3bf3a99102f8bff5368bdefc9d7ae090532fdad2a508ac1271bfbd17363b3a1da23bf9db324a24c238634e3ab28d7f4eca009b4c3953c882f5a4206458a0b4238f3e538308d7339382f38412bbfe7b71e269274526edf7b802ea1ecdf7b8c17f9502b7a6750313329a68b8f8a2d039c8dfe74b9ead98684cfc86e5d0f77c18ba05718b01c33831db17191a0e77f9cef998bbb66a794915b03c94725aceabe9e2b5e25b665a37b5dd3a59a5552bd779dd5f0ae7295d232194eec1ca1ba0324bdc836ba623117e59fcfedab45a86d76d2c768341d327c035a1f5c756cfc06d76b6f7ea31c7a8e782eb48de0aab2fb373ffc2352c4192838323f8$a5245c7f39480a840da0e4c6", "openwall"},
56 // luser-18-12345678.pcap
57 {"$krb5asrep$18$EXAMPLE.COMluser$42e34732112be6cec1532177a6c93af5ec3b2fc7da106c004d6d89ddcb4131092aecbead3e9f30d07b593f4c7adc6478ab50b80fee07db3531471f5f1986c8882c45fef784258f9d43195108b83a74f6dcae1beed179c356c0da4e2d69f122efc579fd207d2b2b241a6c275997f2ec6fec95573a7518cb8b8528d932cc14186e4c5d46cef1eed4f2924ea316d80a62b0bcd98592a11eb69c04ef43b63aeae35e9f8bd8f842d0c9c33d768cd33c55914c2a1fb2f7c640b7270cf2274993c0ce4f413aac8e9d7a231c70dd0c6f8b9c16b47a90fae8d68982a66aa58e2eb8dde93d3504e87b5d4e33827c2aa501ed63544c0578032f395205c63b030cccc699aafb9132692c79a154d645fe83927b0eda$420973360c2e907b9053f1db", "12345678"},
58 // hero-17-abcd.pcap
59 {"$krb5asrep$17$EXAMPLE.COMhero$4e7c79214fd330b2e505a4c75e257e4686029136d54f92ce91bb69d5ffc064e64e925b3ae8bc1df431c74ccaf2075cb4a1a32151b0848964e147bf6f8e4a50caa7931faad50433991e016e312c70ad9007e38166f8df39eda3edd2445cce757e062d0919e663a67eb9fdb472b2a840cf521f18bd794947bcc0c0c6394cc5a60b860c963640867e623732206e7bf904d3b066a17b6f4ea3fd6d74f110ee80052e5297f7a19aaec22e22d582d183d43d6ca1792da187a3a182d1f479c5b4692841ccd701a63735d64584c4f8d199d67876dae5181f4eadfe75e454d0587d0953d7e16cb1b63265da6188b10c1746a2e83c41707bd03fcb2d460d1c6802826a0347b5ee7cdbe5384acad139b4395928bd$7ed0277ba9b853008cc62abe", "abcd"},
60 {NULL}
61 };
62
63 static cl_mem mem_in, mem_dk, mem_out, mem_salt, mem_state, mem_plaintext, mem_edata2;
64 static cl_mem pinned_in, pinned_out;
65 static cl_kernel pbkdf2_init, pbkdf2_loop, pbkdf2_final, asrep_final;
66 static struct fmt_main *self;
67
68 static struct custom_salt *cur_salt;
69
70 typedef struct {
71 pbkdf2_salt pbkdf2;
72 uint32_t etype;
73 uint32_t edata2len;
74 uint8_t edata1[16];
75 // edata2 is a separate __global buffer of variable size
76 } asrep_salt;
77
78 typedef struct {
79 unsigned int cracked;
80 } asrep_out;
81
82 static size_t key_buf_size;
83 static unsigned int *inbuffer;
84 static asrep_salt currentsalt;
85 static asrep_out *output;
86 static int edata_size = 4096;
87 static int new_keys;
88
89 #define ITERATIONS (4096 - 1)
90 #define HASH_LOOPS 105 // Must be made from factors 3, 3, 5, 7, 13
91 #define STEP 0
92 #define SEED 128
93
94 static const char * warn[] = {
95 "xfer: ", ", init: ", ", loop: ", ", final: ", ", asrep: ", ", res xfer: "
96 };
97
98 static int split_events[] = { 2, -1, -1 };
99
100 //This file contains auto-tuning routine(s). Has to be included after formats definitions.
101 #include "opencl_autotune.h"
102
103 /* ------- Helper functions ------- */
get_task_max_work_group_size()104 static size_t get_task_max_work_group_size()
105 {
106 size_t s;
107
108 s = autotune_get_task_max_work_group_size(FALSE, 0, pbkdf2_init);
109 s = MIN(s, autotune_get_task_max_work_group_size(FALSE, 0, pbkdf2_loop));
110 s = MIN(s, autotune_get_task_max_work_group_size(FALSE, 0, pbkdf2_final));
111 s = MIN(s, autotune_get_task_max_work_group_size(FALSE, 0, asrep_final));
112 return s;
113 }
114
create_clobj(size_t gws,struct fmt_main * self)115 static void create_clobj(size_t gws, struct fmt_main *self)
116 {
117 gws *= ocl_v_width;
118
119 key_buf_size = 64 * gws;
120
121 // Allocate memory
122 pinned_in = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, key_buf_size, NULL, &ret_code);
123 HANDLE_CLERROR(ret_code, "Error allocating pinned in");
124 mem_in = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, key_buf_size, NULL, &ret_code);
125 HANDLE_CLERROR(ret_code, "Error allocating mem in");
126 inbuffer = clEnqueueMapBuffer(queue[gpu_id], pinned_in, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, key_buf_size, 0, NULL, NULL, &ret_code);
127 HANDLE_CLERROR(ret_code, "Error mapping page-locked memory");
128
129 pinned_out = clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, sizeof(asrep_out) * gws, NULL, &ret_code);
130 HANDLE_CLERROR(ret_code, "Error allocating pinned out");
131 mem_out = clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, sizeof(asrep_out) * gws, NULL, &ret_code);
132 HANDLE_CLERROR(ret_code, "Error allocating mem out");
133 output = clEnqueueMapBuffer(queue[gpu_id], pinned_out, CL_TRUE, CL_MAP_READ, 0, sizeof(asrep_out) * gws, 0, NULL, NULL, &ret_code);
134 HANDLE_CLERROR(ret_code, "Error mapping page-locked memory");
135
136 mem_dk = clCreateBuffer(context[gpu_id], CL_MEM_READ_WRITE, sizeof(pbkdf2_out) * gws, NULL, &ret_code);
137 HANDLE_CLERROR(ret_code, "Error allocating mem dk");
138
139 mem_state = clCreateBuffer(context[gpu_id], CL_MEM_READ_WRITE, sizeof(pbkdf2_state) * gws, NULL, &ret_code);
140 HANDLE_CLERROR(ret_code, "Error allocating mem_state");
141
142 mem_salt = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(asrep_salt), ¤tsalt, &ret_code);
143 HANDLE_CLERROR(ret_code, "Error allocating mem setting");
144
145 mem_edata2 = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, edata_size, NULL, &ret_code);
146 HANDLE_CLERROR(ret_code, "Error allocating mem edata2");
147
148 mem_plaintext = clCreateBuffer(context[gpu_id], CL_MEM_READ_WRITE, edata_size * gws, NULL, &ret_code);
149 HANDLE_CLERROR(ret_code, "Error allocating mem plaintext");
150
151 HANDLE_CLERROR(clSetKernelArg(pbkdf2_init, 0, sizeof(mem_in), &mem_in), "Error while setting mem_in kernel argument");
152 HANDLE_CLERROR(clSetKernelArg(pbkdf2_init, 1, sizeof(mem_salt), &mem_salt), "Error while setting mem_salt kernel argument");
153 HANDLE_CLERROR(clSetKernelArg(pbkdf2_init, 2, sizeof(mem_state), &mem_state), "Error while setting mem_state kernel argument");
154
155 HANDLE_CLERROR(clSetKernelArg(pbkdf2_loop, 0, sizeof(mem_state), &mem_state), "Error while setting mem_state kernel argument");
156
157 HANDLE_CLERROR(clSetKernelArg(pbkdf2_final, 0, sizeof(mem_salt), &mem_salt), "Error while setting mem_salt kernel argument");
158 HANDLE_CLERROR(clSetKernelArg(pbkdf2_final, 1, sizeof(mem_dk), &mem_dk), "Error while setting mem_dk kernel argument");
159 HANDLE_CLERROR(clSetKernelArg(pbkdf2_final, 2, sizeof(mem_state), &mem_state), "Error while setting mem_state kernel argument");
160
161 HANDLE_CLERROR(clSetKernelArg(asrep_final, 0, sizeof(mem_salt), &mem_salt), "Error while setting mem_salt kernel argument");
162 HANDLE_CLERROR(clSetKernelArg(asrep_final, 1, sizeof(mem_dk), &mem_dk), "Error while setting mem_dk kernel argument");
163 HANDLE_CLERROR(clSetKernelArg(asrep_final, 2, sizeof(mem_edata2), &mem_edata2), "Error while setting mem_edata2 kernel argument");
164 HANDLE_CLERROR(clSetKernelArg(asrep_final, 3, sizeof(mem_plaintext), &mem_plaintext), "Error while setting mem_plaintext kernel argument");
165 HANDLE_CLERROR(clSetKernelArg(asrep_final, 4, sizeof(mem_out), &mem_out), "Error while setting mem_out kernel argument");
166 }
167
release_clobj(void)168 static void release_clobj(void)
169 {
170 if (mem_edata2) {
171 HANDLE_CLERROR(clEnqueueUnmapMemObject(queue[gpu_id], pinned_in, inbuffer, 0, NULL, NULL), "Error Unmapping mem in");
172 HANDLE_CLERROR(clEnqueueUnmapMemObject(queue[gpu_id], pinned_out, output, 0, NULL, NULL), "Error Unmapping mem out");
173 HANDLE_CLERROR(clFinish(queue[gpu_id]), "Error releasing memory mappings");
174
175 HANDLE_CLERROR(clReleaseMemObject(pinned_in), "Release pinned_in");
176 HANDLE_CLERROR(clReleaseMemObject(pinned_out), "Release pinned_out");
177 HANDLE_CLERROR(clReleaseMemObject(mem_in), "Release mem_in");
178 HANDLE_CLERROR(clReleaseMemObject(mem_out), "Release mem_out");
179 HANDLE_CLERROR(clReleaseMemObject(mem_dk), "Release mem_dk");
180 HANDLE_CLERROR(clReleaseMemObject(mem_salt), "Release mem_salt");
181 HANDLE_CLERROR(clReleaseMemObject(mem_state), "Release mem state");
182 HANDLE_CLERROR(clReleaseMemObject(mem_edata2), "Release mem_edata2");
183 HANDLE_CLERROR(clReleaseMemObject(mem_plaintext), "Release mem_plaintext");
184 mem_edata2 = NULL;
185 }
186 }
187
done(void)188 static void done(void)
189 {
190 if (autotuned) {
191 release_clobj();
192
193 HANDLE_CLERROR(clReleaseKernel(pbkdf2_init), "Release Kernel");
194 HANDLE_CLERROR(clReleaseKernel(pbkdf2_loop), "Release Kernel");
195 HANDLE_CLERROR(clReleaseKernel(pbkdf2_final), "Release Kernel");
196 HANDLE_CLERROR(clReleaseKernel(asrep_final), "Release Kernel");
197
198 HANDLE_CLERROR(clReleaseProgram(program[gpu_id]), "Release Program");
199
200 autotuned--;
201 }
202 }
203
init(struct fmt_main * _self)204 static void init(struct fmt_main *_self)
205 {
206 static char valgo[sizeof(ALGORITHM_NAME) + 12] = "";
207
208 self = _self;
209
210 opencl_prepare_dev(gpu_id);
211 /* VLIW5 does better with just 2x vectors due to GPR pressure */
212 if (!options.v_width && amd_vliw5(device_info[gpu_id]))
213 ocl_v_width = 2;
214 else
215 ocl_v_width = opencl_get_vector_width(gpu_id, sizeof(cl_int));
216
217 if (ocl_v_width > 1) {
218 /* Run vectorized kernel */
219 snprintf(valgo, sizeof(valgo),
220 ALGORITHM_NAME " %ux", ocl_v_width);
221 self->params.algorithm_name = valgo;
222 }
223 }
224
reset(struct db_main * db)225 static void reset(struct db_main *db)
226 {
227 if (!autotuned) {
228 char build_opts[128];
229
230 snprintf(build_opts, sizeof(build_opts),
231 "-DHASH_LOOPS=%u -DITERATIONS=%u -DMAX_OUTLEN=%u "
232 "-DPLAINTEXT_LENGTH=%u -DV_WIDTH=%u",
233 HASH_LOOPS, ITERATIONS, MAX_OUTLEN,
234 PLAINTEXT_LENGTH, ocl_v_width);
235 opencl_init("$JOHN/kernels/krb5_kernel.cl", gpu_id,
236 build_opts);
237
238 pbkdf2_init = clCreateKernel(program[gpu_id], "pbkdf2_init", &ret_code);
239 HANDLE_CLERROR(ret_code, "Error creating kernel");
240 crypt_kernel = pbkdf2_loop = clCreateKernel(program[gpu_id], "pbkdf2_loop", &ret_code);
241 HANDLE_CLERROR(ret_code, "Error creating kernel");
242 pbkdf2_final = clCreateKernel(program[gpu_id], "pbkdf2_final", &ret_code);
243 HANDLE_CLERROR(ret_code, "Error creating kernel");
244 asrep_final = clCreateKernel(program[gpu_id], "asrep_final", &ret_code);
245 HANDLE_CLERROR(ret_code, "Error creating kernel");
246
247 //Initialize openCL tuning (library) for this format.
248 opencl_init_auto_setup(SEED, 2 * HASH_LOOPS, split_events,
249 warn, 2, self, create_clobj,
250 release_clobj,
251 edata_size, 0, db);
252
253 //Auto tune execution from shared/included code.
254 autotune_run(self, 4 * ITERATIONS + 4, 0, 200);
255 }
256 }
257
valid(char * ciphertext,struct fmt_main * self)258 static int valid(char *ciphertext, struct fmt_main *self)
259 {
260 return krb5_asrep_valid(ciphertext, self, 0);
261 }
262
clear_keys(void)263 static void clear_keys(void) {
264 memset(inbuffer, 0, key_buf_size);
265 }
266
set_key(char * key,int index)267 static void set_key(char *key, int index)
268 {
269 int i;
270 int length = strlen(key);
271
272 for (i = 0; i < length; i++)
273 ((char*)inbuffer)[GETPOS(i, index)] = key[i];
274
275 new_keys = 1;
276 }
277
get_key(int index)278 static char *get_key(int index)
279 {
280 static char ret[PLAINTEXT_LENGTH + 1];
281 int i = 0;
282
283 while (i < PLAINTEXT_LENGTH &&
284 (ret[i] = ((char*)inbuffer)[GETPOS(i, index)]))
285 i++;
286 ret[i] = 0;
287
288 return ret;
289 }
290
set_salt(void * salt)291 static void set_salt(void *salt)
292 {
293 size_t buf_size;
294
295 cur_salt = *((struct custom_salt **)salt);
296 buf_size = (cur_salt->edata2len + 31) / 32 * 32;
297
298 if (buf_size > edata_size) {
299 edata_size = buf_size;
300 HANDLE_CLERROR(clReleaseMemObject(mem_plaintext), "Release mem_plaintext");
301 HANDLE_CLERROR(clReleaseMemObject(mem_edata2), "Release mem_edata2");
302 mem_plaintext = clCreateBuffer(context[gpu_id], CL_MEM_READ_WRITE, edata_size * global_work_size, NULL, &ret_code);
303 HANDLE_CLERROR(ret_code, "Error allocating mem plaintext");
304
305 mem_edata2 = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, edata_size, NULL, &ret_code);
306 HANDLE_CLERROR(ret_code, "Error allocating mem edata2");
307
308 HANDLE_CLERROR(clSetKernelArg(asrep_final, 2, sizeof(mem_edata2), &mem_edata2), "Error while setting mem_edata2 kernel argument");
309 HANDLE_CLERROR(clSetKernelArg(asrep_final, 3, sizeof(mem_plaintext), &mem_plaintext), "Error while setting mem_plaintext kernel argument");
310 }
311 currentsalt.pbkdf2.length = strlen((char*)cur_salt->salt);
312 currentsalt.pbkdf2.iterations = 4096;
313 currentsalt.pbkdf2.outlen = (cur_salt->etype == 17) ? 16 : 32;
314 currentsalt.etype = cur_salt->etype;
315 currentsalt.edata2len = cur_salt->edata2len;
316
317 memcpy(currentsalt.pbkdf2.salt, cur_salt->salt, currentsalt.pbkdf2.length);
318 memcpy(currentsalt.edata1, cur_salt->edata1, sizeof(currentsalt.edata1));
319 HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_salt, CL_FALSE, 0, sizeof(asrep_salt), ¤tsalt, 0, NULL, NULL), "Copy salt to gpu");
320 HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_edata2, CL_FALSE, 0, currentsalt.edata2len, cur_salt->edata2, 0, NULL, NULL), "Copy edata2 to gpu");
321 }
322
crypt_all(int * pcount,struct db_salt * salt)323 static int crypt_all(int *pcount, struct db_salt *salt)
324 {
325 const int count = *pcount;
326 int i, j;
327 size_t scalar_gws;
328 size_t *lws = local_work_size ? &local_work_size : NULL;
329 size_t gws = GET_NEXT_MULTIPLE(count, local_work_size);
330
331 scalar_gws = gws * ocl_v_width;
332
333 // Copy data to gpu
334 if (ocl_autotune_running || new_keys) {
335 BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0, key_buf_size, inbuffer, 0, NULL, multi_profilingEvent[0]), "Copy data to gpu");
336 new_keys = 0;
337 }
338
339 // Run kernel
340 BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], pbkdf2_init, 1, NULL, &gws, lws, 0, NULL, multi_profilingEvent[1]), "Run initial kernel");
341
342 for (j = 0; j < (ocl_autotune_running ? 1 : ((currentsalt.pbkdf2.outlen + 19) / 20)); j++) {
343 for (i = 0; i < (ocl_autotune_running ? 1 : ITERATIONS / HASH_LOOPS); i++) {
344 BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], pbkdf2_loop, 1, NULL, &gws, lws, 0, NULL, multi_profilingEvent[2]), "Run loop kernel");
345 BENCH_CLERROR(clFinish(queue[gpu_id]), "Error running loop kernel");
346 opencl_process_event();
347 }
348
349 BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], pbkdf2_final, 1, NULL, &gws, lws, 0, NULL, multi_profilingEvent[3]), "Run final pbkdf2 kernel");
350 }
351
352 BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], asrep_final, 1, NULL, &scalar_gws, lws, 0, NULL, multi_profilingEvent[4]), "Run final kernel (SHA1)");
353 BENCH_CLERROR(clFinish(queue[gpu_id]), "Failed running final kernel");
354
355 // Read the result back
356 BENCH_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_TRUE, 0, sizeof(asrep_out) * scalar_gws, output, 0, NULL, multi_profilingEvent[5]), "Copy result back");
357
358 return count;
359 }
360
cmp_all(void * binary,int count)361 static int cmp_all(void *binary, int count)
362 {
363 int index;
364
365 for (index = 0; index < count; index++)
366 if (output[index].cracked)
367 return 1;
368 return 0;
369 }
370
cmp_one(void * binary,int index)371 static int cmp_one(void *binary, int index)
372 {
373 return output[index].cracked;
374 }
375
cmp_exact(char * source,int index)376 static int cmp_exact(char *source, int index)
377 {
378 return 1;
379 }
380
381 struct fmt_main fmt_opencl_krb5_asrep_aes = {
382 {
383 FORMAT_LABEL,
384 FORMAT_NAME,
385 ALGORITHM_NAME,
386 BENCHMARK_COMMENT,
387 BENCHMARK_LENGTH,
388 0,
389 PLAINTEXT_LENGTH,
390 BINARY_SIZE,
391 BINARY_ALIGN,
392 SALT_SIZE,
393 SALT_ALIGN,
394 MIN_KEYS_PER_CRYPT,
395 MAX_KEYS_PER_CRYPT,
396 FMT_CASE | FMT_8_BIT | FMT_DYNA_SALT | FMT_HUGE_INPUT,
397 {NULL},
398 { FORMAT_TAG },
399 tests
400 }, {
401 init,
402 done,
403 reset,
404 fmt_default_prepare,
405 valid,
406 krb5_asrep_split,
407 fmt_default_binary,
408 krb5_asrep_get_salt,
409 {NULL},
410 fmt_default_source,
411 {
412 fmt_default_binary_hash
413 },
414 fmt_default_dyna_salt_hash,
415 NULL,
416 set_salt,
417 set_key,
418 get_key,
419 clear_keys,
420 crypt_all,
421 {
422 fmt_default_get_hash
423 },
424 cmp_all,
425 cmp_one,
426 cmp_exact
427 }
428 };
429
430 #endif /* plugin stanza */
431
432 #endif /* HAVE_OPENCL */
433