Loading [MathJax]/extensions/tex2jax.js
Parallel numerical verification of the σ_odd problem  October 6, 2018
•All Classes Namespaces Files Functions Variables Typedefs Macros
opencl.cpp
Go to the documentation of this file.
1 /* -*- coding: latin-1 -*- */
2 /** \file opencl/opencl/opencl.cpp (February 6, 2018)
3  *
4  * GPLv3 --- Copyright (C) 2017, 2018 Olivier Pirson
5  * http://www.opimedia.be/
6  */
7 
8 // \cond
9 #include <cassert>
10 
11 #include <algorithm>
12 // \endcond
13 
14 #include "../../common/helper/helper.hpp"
15 #include "../../sequential/sequential/sequential.hpp"
16 
17 #include "opencl.hpp"
18 
19 
20 
21 // #define BUILD_LOG
22 
23 
24 
25 namespace opencl {
26  /* ***********
27  * Functions *
28  *************/
29 
30  std::set<nat_type>
31  opencl_check_gentle_varsigma_odd(nat_type first_n, nat_type last_n,
32  unsigned int opencl_nb_number,
33  bool print_bad, bool print_time) {
34  assert(sigmaodd::is_odd(first_n));
35  assert(3 <= first_n);
36  assert(first_n <= last_n);
37  assert(last_n <= sigmaodd::MAX_POSSIBLE_N);
38 
39 #ifdef PRIME16
40  assert(n < 65536);
41 #endif
42 
43  assert(sizeof(nat_type) == 8);
44  assert(sizeof(nat_type) == sizeof(cl_ulong));
45 
46  assert(sizeof(prime_type) == 4);
47  assert(sizeof(prime_type) == sizeof(cl_uint));
48 
49  cl_ulong cumulative_time = 0;
50 
51  const std::string kernel_src(helper::file_to_string("opencl_src/sigmaodd.cl"));
52 
53 
54  // Init OpenCL
55  cl_int error;
56  const cl::Device device = opencl::get_first_device_gpu();
57  const VECTOR_CLASS<cl::Device> devices = {device};
58  const cl::Context context{device};
59 
60 
61  // Init OpenCL buffer
62  size_t primes_byte_size = sigmaodd::odd_primes_table_nb()*sizeof(prime_type);
63  size_t ns_byte_size = opencl_nb_number*sizeof(nat_type);
64  size_t results_byte_size = opencl_nb_number*sizeof(cl_uchar);
65 
66  cl::Buffer primes_buffer{context, CL_MEM_READ_ONLY, primes_byte_size};
67  cl::Buffer ns_buffer{context, CL_MEM_READ_ONLY, ns_byte_size};
68  cl::Buffer results_buffer{context, CL_MEM_WRITE_ONLY, results_byte_size};
69 
70 
71  // Command queue
72  const cl::CommandQueue queue{context, device, CL_QUEUE_PROFILING_ENABLE};
73 
74  error = queue.enqueueWriteBuffer(primes_buffer, CL_TRUE, 0, primes_byte_size,
75  sigmaodd::odd_primes_table_ptr());
76  print_error(error, "clEnqueueWriteBuffer");
77 
78 
79  // Program
80  cl::Program program{context, kernel_src};
81  error = program.build(devices, "-DONLY_MAIN_KERNEL -cl-single-precision-constant -cl-denorms-are-zero -cl-mad-enable -cl-no-signed-zeros -cl-unsafe-math-optimizations -cl-finite-math-only -cl-fast-relaxed-math"
82 #ifdef BUILD_LOG
83  " -cl-nv-verbose"
84 #endif
85  );
86  print_error(error, "clBuildProgram");
87 
88 #ifdef BUILD_LOG
89  {
90  std::string log;
91 
92  error = program.getBuildInfo(device, CL_PROGRAM_BUILD_LOG, &log);
93  print_error(error, "clGetProgramBuildInfo");
94  std::cout << log << std::endl;
95  }
96 #endif
97 
98  // cl::Kernel kernel{program, "check_ns"};
99  cl::Kernel kernel{program, "check_ns__optimized"};
100  // cl::Kernel kernel{program, "check_ns__simplified"};
101 
102 #ifdef BUILD_LOG
103  {
104  size_t work_group_size;
105 
106  error = kernel.getWorkGroupInfo(device, CL_KERNEL_WORK_GROUP_SIZE, &work_group_size);
107  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_WORK_GROUP_SIZE)");
108 
109  size_t compile_work_group_size[3];
110 
111  error = kernel.getWorkGroupInfo(device, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, &compile_work_group_size);
112  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_COMPILE_WORK_GROUP_SIZE)", &compile_work_group_size);
113 
114  cl_ulong local_mem_size;
115 
116  error = kernel.getWorkGroupInfo(device, CL_KERNEL_LOCAL_MEM_SIZE, &local_mem_size);
117  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_LOCAL_MEM_SIZE)");
118 
119  size_t preferred_work_group_size_multiple;
120 
121  error = kernel.getWorkGroupInfo(device, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, &preferred_work_group_size_multiple);
122  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE)");
123 
124  cl_ulong private_mem_size;
125 
126  error = kernel.getWorkGroupInfo(device, CL_KERNEL_PRIVATE_MEM_SIZE, &private_mem_size);
127  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_PRIVATE_MEM_SIZE)");
128  }
129 #endif
130 
131  error = kernel.setArg(0, primes_buffer);
132  print_error(error, "clSetKernelArg(0, primes_buffer)");
133 
134  error = kernel.setArg(1, ns_buffer);
135  print_error(error, "clSetKernelArg(1, ns_buffer)");
136 
137  error = kernel.setArg(3, results_buffer);
138  print_error(error, "clSetKernelArg(3, results_buffer)");
139 
140  std::vector<nat_type> ns(opencl_nb_number);
141  std::vector<unsigned char> results(opencl_nb_number);
142  cl::Event event;
143 
144  std::set<nat_type> bad_table;
145  unsigned int nb = 0;
146 
147  for (nat_type n = first_n; n <= last_n; n += 2) {
148  if (!sigmaodd::is_first_mersenne_prime_unitary_divide_or_square(n)) {
149  ns[nb] = n;
150  ++nb;
151  }
152 
153  if ((nb == opencl_nb_number) || (n == last_n)) {
154  error = kernel.setArg(2, nb);
155  print_error(error, "clSetKernelArg(2, nb)");
156 
157  const size_t nb_byte_size = nb*sizeof(nat_type);
158 
159  error = queue.enqueueWriteBuffer(ns_buffer, CL_TRUE, 0, nb_byte_size, ns.data());
160  print_error(error, "clEnqueueWriteBuffer");
161 
162  error = queue.enqueueNDRangeKernel(kernel,
163  cl::NullRange,
164  cl::NDRange(opencl_nb_number),
165  cl::NullRange,
166  nullptr,
167  &event);
168  print_error(error, "clEnqueueNDRangeKernel");
169 
170  cl_ulong start;
171  cl_ulong stop;
172 
173  error = event.wait();
174  print_error(error, "clWaitForEvents");
175 
176  error = event.getProfilingInfo(CL_PROFILING_COMMAND_START, &start);
177  print_error(error, "getProfilingInfo");
178  error = event.getProfilingInfo(CL_PROFILING_COMMAND_END, &stop);
179  print_error(error, "getProfilingInfo");
180 
181  cumulative_time += stop - start;
182 
183  error = queue.enqueueReadBuffer(results_buffer, CL_TRUE, 0, results_byte_size, results.data());
184  print_error(error, "clEnqueueReadBuffer");
185 
186  for (unsigned int i = 0; i < nb; ++i) {
187  if (!results[i]) {
188  const nat_type bad = ns[i];
189 
190  bad_table.insert(bad);
191  if (print_bad) {
192  std::cout << bad << std::endl;
193  }
194  }
195  }
196  if (print_bad) {
197  std::cout.flush();
198  }
199 
200  nb = 0;
201  }
202  }
203 
204  if (print_bad) {
205  std::cout << "# bad numbers: " << bad_table.size() << std::endl;
206  }
207 
208  if (print_time) {
209  std::cout << "Total OpenCL duration:\t"
210  << helper::duration_ms_to_string(static_cast<double>(cumulative_time)/(1000*1000))
211  << std::endl;
212  }
213 
214  return bad_table;
215  }
216 
217 
218  std::set<nat_type>
219  opencl_check_gentle_varsigma_odd__parallelize_factorization(nat_type first_n, nat_type last_n,
220  unsigned int opencl_nb,
221  bool print_bad, bool print_time) {
222  assert(sigmaodd::is_odd(first_n));
223  assert(3 <= first_n);
224  assert(first_n <= last_n);
225  assert(last_n <= sigmaodd::MAX_POSSIBLE_N);
226 
227  assert(opencl_nb >= 2);
228  assert(sigmaodd::divide_until_odd(opencl_nb) == 1); // power of 2
229 
230 #ifdef PRIME16
231  assert(n < 65536);
232 #endif
233 
234  assert(sizeof(nat_type) == 8);
235  assert(sizeof(nat_type) == sizeof(cl_ulong));
236 
237  assert(sizeof(prime_type) == 4);
238  assert(sizeof(prime_type) == sizeof(cl_uint));
239 
240  cl_ulong cumulative_time = 0;
241 
242  std::string kernel_src(helper::file_to_string("opencl_src/sigmaodd.cl"));
243 
244 
245  // Init OpenCL
246  cl_int error;
247  const cl::Device device = opencl::get_first_device_gpu();
248  const VECTOR_CLASS<cl::Device> devices = {device};
249  const cl::Context context{device};
250 
251 
252  // Init OpenCL buffer
253  size_t primes_byte_size = sigmaodd::odd_primes_table_nb()*sizeof(prime_type);
254  size_t factors_byte_size = opencl_nb*sizeof(nat_type);
255 
256  cl::Buffer primes_buffer{context, CL_MEM_READ_ONLY, primes_byte_size};
257  cl::Buffer factors_buffer{context, CL_MEM_READ_WRITE, factors_byte_size};
258 
259 
260  // Command queue
261  const cl::CommandQueue queue{context, device, CL_QUEUE_PROFILING_ENABLE};
262 
263  error = queue.enqueueWriteBuffer(primes_buffer, CL_TRUE, 0, primes_byte_size,
264  sigmaodd::odd_primes_table_ptr());
265  print_error(error, "clEnqueueWriteBuffer");
266 
267 
268  // Program
269  cl::Program program{context, kernel_src};
270  error = program.build(devices, "-DONLY_MAIN_KERNEL -cl-single-precision-constant -cl-denorms-are-zero -cl-mad-enable -cl-no-signed-zeros -cl-unsafe-math-optimizations -cl-finite-math-only -cl-fast-relaxed-math"
271 #ifdef BUILD_LOG
272  " -cl-nv-verbose"
273 #endif
274  );
275  print_error(error, "clBuildProgram");
276 
277 #ifdef BUILD_LOG
278  {
279  std::string log;
280 
281  error = program.getBuildInfo(device, CL_PROGRAM_BUILD_LOG, &log);
282  print_error(error, "clGetProgramBuildInfo");
283  std::cout << log << std::endl;
284  }
285 #endif
286 
287  cl::Kernel kernel{program, "compute_partial_sigma_odd"};
288 
289 #ifdef BUILD_LOG
290  {
291  size_t work_group_size;
292 
293  error = kernel.getWorkGroupInfo(device, CL_KERNEL_WORK_GROUP_SIZE, &work_group_size);
294  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_WORK_GROUP_SIZE)");
295 
296  size_t compile_work_group_size[3];
297 
298  error = kernel.getWorkGroupInfo(device, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, &compile_work_group_size);
299  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_COMPILE_WORK_GROUP_SIZE)", &compile_work_group_size);
300 
301  cl_ulong local_mem_size;
302 
303  error = kernel.getWorkGroupInfo(device, CL_KERNEL_LOCAL_MEM_SIZE, &local_mem_size);
304  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_LOCAL_MEM_SIZE)");
305 
306  size_t preferred_work_group_size_multiple;
307 
308  error = kernel.getWorkGroupInfo(device, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, &preferred_work_group_size_multiple);
309  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE)");
310 
311  cl_ulong private_mem_size;
312 
313  error = kernel.getWorkGroupInfo(device, CL_KERNEL_PRIVATE_MEM_SIZE, &private_mem_size);
314  print_error(error, "clGetKernelWorkGroupInfo(CL_KERNEL_PRIVATE_MEM_SIZE)");
315  }
316 #endif
317 
318  error = kernel.setArg(0, primes_buffer);
319  print_error(error, "clSetKernelArg(0, primes_buffer)");
320 
321  error = kernel.setArg(2, opencl_nb);
322  print_error(error, "clSetKernelArg(2, opencl_nb)");
323 
324  error = kernel.setArg(4, factors_buffer);
325  print_error(error, "clSetKernelArg(4, factors_buffer)");
326 
327  cl::Event event;
328 
329  std::set<nat_type> bad_table;
330 
331  for (nat_type n = first_n; n <= last_n; n += 2) {
332  if (!sigmaodd::is_first_mersenne_prime_unitary_divide_or_square(n)) {
333  unsigned int prime_offset = 0;
334  nat_type sqrt_n_divided = sigmaodd::floor_square_root(n);
335  nat_type n_divided = n;
336  nat_type varsigma_odd = 1;
337 
338  do {
339  error = kernel.setArg(1, prime_offset);
340 #ifndef NDEBUG
341  print_error(error, "clSetKernelArg(1, prime_offset)");
342 #endif
343 
344  assert(n_divided != 0);
345 
346  error = kernel.setArg(3, n_divided);
347 #ifndef NDEBUG
348  print_error(error, "clSetKernelArg(3, n_divided)");
349 #endif
350 
351  error = queue.enqueueNDRangeKernel(kernel,
352  cl::NullRange,
353  cl::NDRange(opencl_nb),
354  cl::NullRange,
355  nullptr,
356  &event);
357 #ifndef NDEBUG
358  print_error(error, "clEnqueueNDRangeKernel");
359 #endif
360 
361  cl_ulong start;
362  cl_ulong stop;
363 
364  error = event.wait();
365 #ifndef NDEBUG
366  print_error(error, "clWaitForEvents");
367 #endif
368 
369  error = event.getProfilingInfo(CL_PROFILING_COMMAND_START, &start);
370 #ifndef NDEBUG
371  print_error(error, "getProfilingInfo");
372 #endif
373  error = event.getProfilingInfo(CL_PROFILING_COMMAND_END, &stop);
374 #ifndef NDEBUG
375  print_error(error, "getProfilingInfo");
376 #endif
377 
378  cumulative_time += stop - start;
379 
380  nat_type results[2];
381 
382  error = queue.enqueueReadBuffer(factors_buffer, CL_TRUE, 0, sizeof(nat_type)*2, &results);
383 #ifndef NDEBUG
384  print_error(error, "clEnqueueReadBuffer");
385 #endif
386 
387  assert(results[0] != 0);
388  assert(results[0] <= n_divided);
389 
390  if (results[0] > 1) { // at least one prime factor found
391  assert(sigmaodd::is_divide(results[0], n_divided));
392 
393  n_divided /= results[0];
394  sqrt_n_divided = sigmaodd::floor_square_root(n_divided);
395 
396  assert(results[1] != 0);
397 
398  varsigma_odd *= sigmaodd::divide_until_odd(results[1]);
399 
400  if (varsigma_odd
401  * sequential::sequential_sigma_odd_upper_bound_with_sqrt(n_divided, bad_table, 0,
402  sqrt_n_divided) < n) {
403  n_divided = 1;
404  }
405  }
406  prime_offset = std::min(prime_offset + opencl_nb,
407  sigmaodd::odd_primes_table_nb() - opencl_nb);
408  } while ((n_divided > 1)
409  && (*(sigmaodd::odd_primes_table_ptr() + prime_offset) <= sqrt_n_divided));
410 
411  if (n_divided > 1) { // remain n_divided is prime
412  varsigma_odd *= sigmaodd::divide_until_odd(n_divided + 1);
413  }
414 
415  if (varsigma_odd > n) { // bad number
416  bad_table.insert(n);
417  if (print_bad) {
418  std::cout << n << std::endl;
419  std::cout.flush();
420  }
421  }
422  }
423  }
424 
425  if (print_bad) {
426  std::cout << "# bad numbers: " << bad_table.size() << std::endl;
427  }
428 
429  if (print_time) {
430  std::cout << "Total OpenCL duration:\t"
431  << helper::duration_ms_to_string(static_cast<double>(cumulative_time)/(1000*1000))
432  << std::endl;
433  }
434 
435  return bad_table;
436  }
437 
438 
439  std::vector<nat_type>
440  opencl_check_ns(const std::vector<nat_type> &ns) {
441  assert(sizeof(nat_type) == sizeof(cl_ulong));
442  assert(sizeof(prime_type) == sizeof(cl_uint));
443 
444  const std::string kernel_src(helper::file_to_string("opencl_src/sigmaodd.cl"));
445 
446  // Init OpenCL
447  cl_int error;
448  const cl::Device device = opencl::get_first_device_gpu();
449  const VECTOR_CLASS<cl::Device> devices = {device};
450  const cl::Context context{device};
451 
452 
453  // Init OpenCL buffer
454  size_t primes_byte_size = sigmaodd::odd_primes_table_nb()*sizeof(prime_type);
455  size_t ns_byte_size = ns.size()*sizeof(nat_type);
456 
457  cl::Buffer primes_buffer{context, CL_MEM_READ_ONLY, primes_byte_size};
458  cl::Buffer ns_buffer{context, CL_MEM_READ_ONLY, ns_byte_size};
459  cl::Buffer results_buffer{context, CL_MEM_WRITE_ONLY, ns_byte_size};
460 
461 
462  // Command queue
463  const cl::CommandQueue queue{context, device, CL_QUEUE_PROFILING_ENABLE};
464 
465  queue.enqueueWriteBuffer(primes_buffer, CL_TRUE, 0, primes_byte_size,
466  sigmaodd::odd_primes_table_ptr());
467 
468 
469  // Program
470  cl::Program program{context, kernel_src};
471 
472  error = program.build(devices);
473  print_error(error, "clBuildProgram");
474 
475  cl::Kernel kernel{program, "check_ns"};
476 
477  error = kernel.setArg(0, primes_buffer);
478  print_error(error, "clSetKernelArg(0, primes_buffer)");
479 
480  error = kernel.setArg(1, ns_buffer);
481  print_error(error, "clSetKernelArg(1, ns_buffer)");
482 
483  error = kernel.setArg(2, results_buffer);
484  print_error(error, "clSetKernelArg(2, results_buffer)");
485 
486  cl::Event event;
487 
488  queue.enqueueWriteBuffer(ns_buffer, CL_TRUE, 0, ns_byte_size, ns.data());
489  error = queue.enqueueNDRangeKernel(kernel,
490  cl::NullRange,
491  cl::NDRange(ns.size()),
492  cl::NullRange,
493  nullptr,
494  &event);
495  print_error(error, "clEnqueueNDRangeKernel");
496 
497  event.wait();
498 
499  std::vector<nat_type> results(ns.size());
500 
501  queue.enqueueReadBuffer(results_buffer, CL_TRUE, 0, ns_byte_size, results.data());
502 
503  std::vector<nat_type> bads;
504 
505  for (unsigned int i = 0; i < results.size(); ++i) {
506  if (!results[i]) {
507  bads.push_back(ns[i]);
508  }
509  }
510 
511  return bads;
512  }
513 
514 
515  std::vector<nat_type>
516  opencl_run_program_on_ns(const std::string &filename, const std::string &kernal_name,
517  const std::vector<nat_type> &ns) {
518  assert(sizeof(nat_type) == sizeof(cl_ulong));
519  assert(sizeof(prime_type) == sizeof(cl_uint));
520 
521  const std::string kernel_src(helper::file_to_string(filename));
522 
523  // Init OpenCL
524  const cl::Device device = opencl::get_first_device_gpu();
525  const VECTOR_CLASS<cl::Device> devices = {device};
526  const cl::Context context{device};
527 
528 
529  // Init OpenCL buffer
530  size_t primes_byte_size = sigmaodd::odd_primes_table_nb()*sizeof(prime_type);
531  size_t ns_byte_size = ns.size()*sizeof(nat_type);
532 
533  cl::Buffer primes_buffer{context, CL_MEM_READ_ONLY, primes_byte_size};
534  cl::Buffer ns_buffer{context, CL_MEM_READ_ONLY, ns_byte_size};
535  cl::Buffer results_buffer{context, CL_MEM_WRITE_ONLY, ns_byte_size};
536 
537 
538  // Command queue
539  const cl::CommandQueue queue{context, device};
540 
541  queue.enqueueWriteBuffer(primes_buffer, CL_TRUE, 0, primes_byte_size,
542  sigmaodd::odd_primes_table_ptr());
543  queue.enqueueWriteBuffer(ns_buffer, CL_TRUE, 0, ns_byte_size, ns.data());
544 
545 
546  // Program
547  cl::Program program{context, kernel_src};
548 
549  program.build(devices);
550 
551  cl::Kernel kernel{program, kernal_name.c_str()};
552 
553  kernel.setArg(0, primes_buffer);
554  kernel.setArg(1, ns_buffer);
555  kernel.setArg(2, results_buffer);
556 
557  queue.enqueueNDRangeKernel(kernel,
558  cl::NullRange,
559  cl::NDRange(ns.size()),
560  cl::NullRange,
561  nullptr,
562  nullptr);
563 
564  std::vector<nat_type> results(ns.size());
565 
566  queue.enqueueReadBuffer(results_buffer, CL_TRUE, 0, ns_byte_size, results.data());
567 
568  return results;
569  }
570 
571 } // namespace opencl
opencl::opencl_run_program_on_ns
std::vector< nat_type > opencl_run_program_on_ns(const std::string &filename, const std::string &kernal_name, const std::vector< nat_type > &ns)
Run the OpenCL program from filename on ns. Used for tests.
Definition: opencl.cpp:516
opencl::get_first_device_gpu
cl::Device get_first_device_gpu()
Return the first GPU device found. If not found then print an error message and exit.
Definition: helper.cpp:194
opencl::opencl_check_gentle_varsigma_odd
std::set< nat_type > opencl_check_gentle_varsigma_odd(nat_type first_n, nat_type last_n, unsigned int opencl_nb_number, bool print_bad, bool print_time)
Check in the order all odd gentle numbers between first_n and last_n, and if print_bad then print all...
Definition: opencl.cpp:31
sigmaodd::floor_square_root
nat_type floor_square_root(nat_type n)
Return the square root of n rounded to below.
Definition: helper.cpp:76
opencl::opencl_check_ns
std::vector< nat_type > opencl_check_ns(const std::vector< nat_type > &ns)
Check all numbers in ns and return vector of all bad numbers found.
Definition: opencl.cpp:440
sigmaodd::is_odd
constexpr bool is_odd(nat_type n)
Return true iff n is odd.
Definition: helper__inline.hpp:32
helper::duration_ms_to_string
std::string duration_ms_to_string(double duration_ms)
Return a string with the duration expressed in milliseconds, seconds, minutes and hours...
Definition: helper.cpp:48
sigmaodd::divide_until_odd
nat_type divide_until_odd(nat_type n)
Return n divided by 2 until the result is odd.
Definition: divisors__inline.hpp:132
opencl::nat_type
sigmaodd::nat_type nat_type
Definition: helper.hpp:28
sigmaodd::MAX_POSSIBLE_N
constexpr nat_type MAX_POSSIBLE_N
Lower bound of the bigger number such that it is possible to compute the result of the sigma function...
Definition: helper.hpp:54
sigmaodd::is_divide
constexpr bool is_divide(nat_type d, nat_type n)
Return true iff d divide n, i.e. if n is divisible by d.
Definition: divisors__inline.hpp:73
sigmaodd::odd_primes_table_ptr
const prime_type * odd_primes_table_ptr()
Return a pointer to the first number in the precalculated table.
Definition: primes__inline.hpp:66
opencl::print_error
void print_error(cl_int code, std::string message, bool only_if_error, bool exit_if_error)
Print an error message corresponding to the error code.
Definition: helper.cpp:352
sequential::sequential_sigma_odd_upper_bound_with_sqrt
constexpr nat_type sequential_sigma_odd_upper_bound_with_sqrt(nat_type n, const std::set< nat_type > &bad_table, nat_type bad_first_n, nat_type sqrt_n)
Return an upper bound of varsigma_odd(n).
Definition: sequential__inline.hpp:76
helper::file_to_string
std::string file_to_string(std::string filename)
Read the file and return its content to a string. If failed then print a error message and exit...
Definition: helper.cpp:75
sigmaodd::is_first_mersenne_prime_unitary_divide_or_square
constexpr bool is_first_mersenne_prime_unitary_divide_or_square(nat_type n)
Return true iff is_first_mersenne_prime_unitary_divide(n) or is_square(n).
Definition: divisors__inline.hpp:92
sigmaodd::odd_primes_table_nb
constexpr unsigned int odd_primes_table_nb()
Return the number of odd prime numbers in the precalculated table.
Definition: primes__inline.hpp:20
opencl::opencl_check_gentle_varsigma_odd__parallelize_factorization
std::set< nat_type > opencl_check_gentle_varsigma_odd__parallelize_factorization(nat_type first_n, nat_type last_n, unsigned int opencl_nb, bool print_bad, bool print_time)
Like opencl_check_gentle_varsigma_odd() but instead parallelize on group of opencl_nb numbers...
Definition: opencl.cpp:219
opencl::prime_type
sigmaodd::prime_type prime_type
Definition: helper.hpp:29
opencl.hpp
Implementation of the OpenCL parallel algorithms presented in the report.
varsigma_odd
nat_type varsigma_odd(__global const prime_type *primes, nat_type n)
Return varsigma_odd(n), i.e. the sum of all odd divisors of n, divided by 2 until to be odd...
Definition: sigmaodd.cl:497
Generated by   doxygen 1.8.13