gpu_mandelbrot_between.cu 5.4 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129
  1. /* StarPU --- Runtime system for heterogeneous multicore architectures.
  2. *
  3. * Copyright (C) 2019 Mael Keryell
  4. *
  5. * StarPU is free software; you can redistribute it and/or modify
  6. * it under the terms of the GNU Lesser General Public License as published by
  7. * the Free Software Foundation; either version 2.1 of the License, or (at
  8. * your option) any later version.
  9. *
  10. * StarPU is distributed in the hope that it will be useful, but
  11. * WITHOUT ANY WARRANTY; without even the implied warranty of
  12. * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
  13. *
  14. * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  15. */
  16. #include <stdio.h>
  17. #include <stdint.h>
  18. #include <starpu.h>
  19. #define THREADS_PER_BLOCK 64
  20. static inline long long jlstarpu_max(long long a, long long b)
  21. {
  22. return (a > b) ? a : b;
  23. }
  24. static inline long long jlstarpu_interval_size(long long start, long long step, long long stop)
  25. {
  26. if (stop >= start){
  27. return jlstarpu_max(0, (stop - start + 1) / step);
  28. } else {
  29. return jlstarpu_max(0, (stop - start - 1) / step);
  30. }
  31. }
  32. __device__ static inline long long jlstarpu_max__device(long long a, long long b)
  33. {
  34. return (a > b) ? a : b;
  35. }
  36. __device__ static inline long long jlstarpu_interval_size__device(long long start, long long step, long long stop)
  37. {
  38. if (stop >= start){
  39. return jlstarpu_max__device(0, (stop - start + 1) / step);
  40. } else {
  41. return jlstarpu_max__device(0, (stop - start - 1) / step);
  42. }
  43. }
  44. __global__ void mandelbrot(int64_t kernel_ids__start_1, int64_t kernel_ids__step_1, int64_t kernel_ids__dim_1, int64_t kernel_ids__start_2,
  45. int64_t kernel_ids__step_2, int64_t kernel_ids__dim_2, double* ptr_hF6lCYyJ, int64_t local_width,
  46. int64_t* ptr_qoUGBRtY, int64_t local_height, double conv_limit, int64_t* ptr_A5zD9sJZ,
  47. uint32_t ld_A5zD9sJZ)
  48. {
  49. int64_t THREAD_ID = (int64_t) ((((blockIdx).x) * ((blockDim).x)) + ((threadIdx).x));
  50. if ((THREAD_ID) >= (((1) * (kernel_ids__dim_2)) * (kernel_ids__dim_1)))
  51. {
  52. return ;
  53. };
  54. int64_t kernel_ids__index_1 = (int64_t) (((THREAD_ID) / ((1) * (kernel_ids__dim_2))) % (kernel_ids__dim_1));
  55. int64_t kernel_ids__index_2 = (int64_t) (((THREAD_ID) / (1)) % (kernel_ids__dim_2));
  56. int64_t x = (int64_t) ((kernel_ids__start_1) + ((kernel_ids__index_1) * (kernel_ids__step_1)));
  57. int64_t y = (int64_t) ((kernel_ids__start_2) + ((kernel_ids__index_2) * (kernel_ids__step_2)));
  58. double max_iterations = (double) (ptr_hF6lCYyJ[(5) - (1)]);
  59. double zoom = (double) ((ptr_hF6lCYyJ[(3) - (1)]) * (0.25296875));
  60. int64_t X = (int64_t) ((x) + ((local_width) * ((ptr_qoUGBRtY[(2) - (1)]) - (1))));
  61. int64_t Y = (int64_t) ((y) + ((local_height) * ((ptr_qoUGBRtY[(1) - (1)]) - (1))));
  62. double cr = (double) ((ptr_hF6lCYyJ[(1) - (1)]) + (((X) - ((ptr_hF6lCYyJ[(3) - (1)]) / (2))) / (zoom)));
  63. double zr = (double) (cr);
  64. double ci = (double) ((ptr_hF6lCYyJ[(2) - (1)]) + (((Y) - ((ptr_hF6lCYyJ[(4) - (1)]) / (2))) / (zoom)));
  65. double zi = (double) (ci);
  66. int64_t n = (int64_t) (0);
  67. int64_t b1 = (int64_t) (((n) < (max_iterations)) + ((((zr) * (zr)) + ((zi) * (zi))) < ((conv_limit) * (conv_limit))));
  68. while ((b1) >= (2))
  69. {
  70. double tmp = (double) ((((zr) * (zr)) - ((zi) * (zi))) + (cr));
  71. zi = ((2) * (zr) * (zi)) + (ci);
  72. zr = tmp;
  73. n = (n) + (1);
  74. b1 = ((n) <= (max_iterations)) + ((((zr) * (zr)) + ((zi) * (zi))) <= ((conv_limit) * (conv_limit)));
  75. }
  76. ;
  77. if ((n) < (max_iterations))
  78. {
  79. ptr_A5zD9sJZ[((y) + (((x) - (1)) * (ld_A5zD9sJZ))) - (1)] = ((255) * (n)) / (max_iterations);
  80. } else
  81. {
  82. ptr_A5zD9sJZ[((y) + (((x) - (1)) * (ld_A5zD9sJZ))) - (1)] = 0;
  83. }
  84. ;
  85. }
  86. extern "C" void CUDA_mandelbrot(void** buffers_uwrYFDVe, void* cl_arg_uwrYFDVe)
  87. {
  88. uint32_t ld_A5zD9sJZ = (uint32_t) (STARPU_MATRIX_GET_LD(buffers_uwrYFDVe[(1) - (1)]));
  89. int64_t* ptr_A5zD9sJZ = (int64_t*) (STARPU_MATRIX_GET_PTR(buffers_uwrYFDVe[(1) - (1)]));
  90. double* ptr_hF6lCYyJ = (double*) (STARPU_VECTOR_GET_PTR(buffers_uwrYFDVe[(2) - (1)]));
  91. int64_t* ptr_qoUGBRtY = (int64_t*) (STARPU_VECTOR_GET_PTR(buffers_uwrYFDVe[(3) - (1)]));
  92. int64_t local_width = (int64_t) (STARPU_MATRIX_GET_NY(buffers_uwrYFDVe[(1) - (1)]));
  93. int64_t local_height = (int64_t) (STARPU_MATRIX_GET_NX(buffers_uwrYFDVe[(1) - (1)]));
  94. double conv_limit = (double) (2.0);
  95. int64_t kernel_ids__start_1 = (int64_t) (1);
  96. int64_t kernel_ids__step_1 = (int64_t) (1);
  97. int64_t kernel_ids__dim_1 = (int64_t) (jlstarpu_interval_size(kernel_ids__start_1, kernel_ids__step_1, local_width));
  98. int64_t kernel_ids__start_2 = (int64_t) (1);
  99. int64_t kernel_ids__step_2 = (int64_t) (1);
  100. int64_t kernel_ids__dim_2 = (int64_t) (jlstarpu_interval_size(kernel_ids__start_2, kernel_ids__step_2, local_height));
  101. int64_t nthreads = (int64_t) (((1) * (kernel_ids__dim_1)) * (kernel_ids__dim_2));
  102. int64_t nblocks = (int64_t) ((((nthreads) + (THREADS_PER_BLOCK)) - (1)) / (THREADS_PER_BLOCK));
  103. mandelbrot
  104. <<< nblocks, THREADS_PER_BLOCK, 0, starpu_cuda_get_local_stream()
  105. >>> (kernel_ids__start_1, kernel_ids__step_1, kernel_ids__dim_1, kernel_ids__start_2,
  106. kernel_ids__step_2, kernel_ids__dim_2, ptr_hF6lCYyJ, local_width,
  107. ptr_qoUGBRtY, local_height, conv_limit, ptr_A5zD9sJZ,
  108. ld_A5zD9sJZ);
  109. ;
  110. cudaStreamSynchronize(starpu_cuda_get_local_stream());
  111. }