master_kernel.h 11 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272
  1. const char *kernel_source_master = " \
  2. int is_intersecting(float a, float b, float c, float d, float p, float q, float r, float s) { \
  3. float det, gamma, lambda; \
  4. \
  5. det = (c - a) * (s - q) - (r - p) * (d - b); \
  6. \
  7. if (det == 0) { \
  8. return 1; \
  9. } \
  10. else { \
  11. lambda = ((s - q) * (r - a) + (p - r) * (s - b)) / det; \
  12. gamma = ((b - d) * (r - a) + (c - a) * (s - b)) / det; \
  13. return (0 < lambda && lambda < 1) && (0 < gamma && gamma < 1); \
  14. } \
  15. } \
  16. \
  17. void norm(float dest[2], float a[2]){ \
  18. float magnitude = sqrt(a[0] * a[0] + a[1] * a[1]); \
  19. \
  20. dest[0] = a[0] / magnitude; \
  21. dest[1] = a[1] / magnitude; \
  22. } \
  23. \
  24. void sub(float dest[2], float a[2], float b[2]){ \
  25. dest[0] = a[0] - b[0]; \
  26. dest[1] = a[1] - b[1]; \
  27. } \
  28. \
  29. void add(float dest[2], float a[2], float b[2]){ \
  30. dest[0] = a[0] + b[0]; \
  31. dest[1] = a[1] + b[1]; \
  32. } \
  33. \
  34. __kernel void raster_kernel(__global int* objects, __global float* verticies, __global float* normals, __global float* uvs, __global int* triangles, __global uint* frame_buffer, __global uchar* textures, __global int* texture_info, int object_count, int width, int height, int show_buffer){ \
  35. int id_x = get_global_id(0) - width / 2; \
  36. int id_y = get_global_id(1) - height / 2; \
  37. \
  38. float z_pixel = 0; \
  39. uint frame_pixel = 0x22222222; \
  40. \
  41. float highest_z = 800;\
  42. float lowest_z = 0;\
  43. \
  44. int has_normals = 1;\
  45. int has_uvs = 1;\
  46. \
  47. float w0;\
  48. float w1;\
  49. float w2;\
  50. \
  51. \
  52. for (int object = 0; object < object_count; object++){ \
  53. int base = object * 15;\
  54. \
  55. int object_x = objects[base + 0]; \
  56. int object_y = objects[base + 1]; \
  57. int object_z = objects[base + 2]; \
  58. int object_r_x = objects[base + 3]; \
  59. int object_r_y = objects[base + 4]; \
  60. int object_r_z = objects[base + 5]; \
  61. int object_s_x = objects[base + 6]; \
  62. int object_s_y = objects[base + 7]; \
  63. int object_s_z = objects[base + 8]; \
  64. \
  65. int triangle_count = objects[base + 9];\
  66. int triangle_index = objects[base + 10];\
  67. int vertex_index = objects[base + 11];\
  68. int normal_index = objects[base + 12];\
  69. int uv_index = objects[base + 13];\
  70. int texture_index = objects[base + 14];\
  71. \
  72. for (int triangle = 0; triangle < triangle_count; triangle++){\
  73. int triangle_base = (triangle + triangle_index) * 9; \
  74. \
  75. int i0 = (vertex_index + triangles[triangle_base + 0]) * 3;\
  76. int i1 = (vertex_index + triangles[triangle_base + 1]) * 3;\
  77. int i2 = (vertex_index + triangles[triangle_base + 2]) * 3;\
  78. \
  79. int i3 = (normal_index + triangles[triangle_base + 3]) * 3;\
  80. int i4 = (normal_index + triangles[triangle_base + 4]) * 3;\
  81. int i5 = (normal_index + triangles[triangle_base + 5]) * 3;\
  82. \
  83. int i6 = (uv_index + triangles[triangle_base + 6]) * 3;\
  84. int i7 = (uv_index + triangles[triangle_base + 7]) * 3;\
  85. int i8 = (uv_index + triangles[triangle_base + 8]) * 3;\
  86. \
  87. float z0 = verticies[i0 + 2] * object_s_z + object_z;\
  88. float x0 = (verticies[i0 + 0] * object_s_x + object_x);\
  89. float y0 = (verticies[i0 + 1] * object_s_y + object_y);\
  90. \
  91. float z1 = verticies[i1 + 2] * object_s_z + object_z;\
  92. float x1 = (verticies[i1 + 0] * object_s_x + object_x);\
  93. float y1 = (verticies[i1 + 1] * object_s_y + object_y);\
  94. \
  95. float z2 = verticies[i2 + 2] * object_s_z + object_z;\
  96. float x2 = (verticies[i2 + 0] * object_s_x + object_x);\
  97. float y2 = (verticies[i2 + 1] * object_s_y + object_y);\
  98. \
  99. \
  100. if (i3 < 0 || i4 < 0 || i5 < 0){\
  101. has_normals = 0;\
  102. }\
  103. if (i6 < 0 || i7 < 0 || i8 < 0){\
  104. has_uvs = 0;\
  105. }\
  106. \
  107. if (isinf(x0) || isinf(y0) || isinf(z0) || isinf(x1) || isinf(y1) || isinf(z1) || isinf(x2) || isinf(y2) || isinf(z2)){\
  108. continue;\
  109. }\
  110. \
  111. float smallest_x = x0; \
  112. float largest_x = x0; \
  113. float smallest_y = y0; \
  114. float largest_y = y0; \
  115. \
  116. if (x0 > largest_x) largest_x = x0;\
  117. if (x1 > largest_x) largest_x = x1;\
  118. if (x2 > largest_x) largest_x = x2;\
  119. \
  120. if (x0 < smallest_x) smallest_x = x0;\
  121. if (x1 < smallest_x) smallest_x = x1;\
  122. if (x2 < smallest_x) smallest_x = x2;\
  123. \
  124. if (y0 > largest_y) largest_y = y0;\
  125. if (y1 > largest_y) largest_y = y1;\
  126. if (y2 > largest_y) largest_y = y2;\
  127. \
  128. if (y0 < smallest_y) smallest_y = y0;\
  129. if (y1 < smallest_y) smallest_y = y1;\
  130. if (y2 < smallest_y) smallest_y = y2;\
  131. \
  132. smallest_x = fmin(smallest_x, 0); \
  133. largest_x = fmax(largest_x, width); \
  134. smallest_y = fmin(smallest_y, 0); \
  135. largest_y = fmax(largest_y, height); \
  136. \
  137. if (id_x >= smallest_x && id_x <= largest_x && id_y >= smallest_y && id_y <= largest_y){ \
  138. int intersections = 0; \
  139. \
  140. intersections += is_intersecting(id_x, id_y, 10000, 100000, x0, y0, x1, y1); \
  141. intersections += is_intersecting(id_x, id_y, 10000, 100000, x1, y1, x2, y2); \
  142. intersections += is_intersecting(id_x, id_y, 10000, 100000, x2, y2, x0, y0); \
  143. \
  144. if (intersections % 2 == 0){ \
  145. continue; \
  146. } \
  147. \
  148. float denominator = (y1 - y2) * (x0 - x2) + (x2 - x1) * (y0 - y2); \
  149. \
  150. if (denominator <= 0) { \
  151. continue; \
  152. } \
  153. w0 = ((y1 - y2) * (id_x - x2) + (x2 - x1) * (id_y - y2)) / denominator; \
  154. w1 = ((y2 - y0) * (id_x - x0) + (x0 - x2) * (id_y - y0)) / denominator; \
  155. w2 = 1.0 - w0 - w1; \
  156. \
  157. float z = w0 * z0 + w1 * z1 + w2 * z2; \
  158. \
  159. if (z > z_pixel){ \
  160. z_pixel = z; \
  161. \
  162. float n_x0 = normals[i3 + 0];\
  163. float n_y0 = normals[i3 + 1];\
  164. float n_z0 = normals[i3 + 2];\
  165. \
  166. float n_x1 = normals[i4 + 0];\
  167. float n_y1 = normals[i4 + 1];\
  168. float n_z1 = normals[i4 + 2];\
  169. \
  170. float n_x2 = normals[i5 + 0];\
  171. float n_y2 = normals[i5 + 1];\
  172. float n_z2 = normals[i5 + 2];\
  173. \
  174. float u_x0 = uvs[i6 + 0];\
  175. float u_y0 = uvs[i6 + 1];\
  176. float u_z0 = uvs[i6 + 2];\
  177. \
  178. float u_x1 = uvs[i7 + 0];\
  179. float u_y1 = uvs[i7 + 1];\
  180. float u_z1 = uvs[i7 + 2];\
  181. \
  182. float u_x2 = uvs[i8 + 0];\
  183. float u_y2 = uvs[i8 + 1];\
  184. float u_z2 = uvs[i8 + 2];\
  185. \
  186. switch (show_buffer){\
  187. case 0:{\
  188. float ux = w0 * u_x0 + w1 * u_x1 + w2 * u_x2;\
  189. float uy = w0 * u_y0 + w1 * u_y1 + w2 * u_y2;\
  190. \
  191. int texture_width = texture_info[texture_index * 2];\
  192. int texture_height = texture_info[texture_index * 2 + 1];\
  193. \
  194. int ix = clamp((int)ux * texture_width, 0, texture_width - 1);\
  195. int iy = clamp((int)uy * texture_height, 0, texture_height - 1);\
  196. \
  197. int pixel_offset = (iy * texture_width + ix) * 4;\
  198. \
  199. uchar r = textures[0 + 0];\
  200. uchar g = textures[0 + 1];\
  201. uchar b = textures[0 + 2];\
  202. uchar a = textures[0 + 3];\
  203. \
  204. frame_pixel = (a << 24) | (r << 16) | (g << 8) | b;\
  205. \
  206. break;}\
  207. case 1:{\
  208. float z = clamp(z_pixel, 0.0f, highest_z);\
  209. \
  210. float norm_z = z / highest_z;\
  211. \
  212. uchar intensity = (uchar)(norm_z * 255.0f);\
  213. \
  214. frame_pixel = 0xFF000000 | (intensity << 16) | (intensity << 8) | intensity;\
  215. \
  216. break;}\
  217. case 2:{\
  218. float nx = w0 * n_x0 + w1 * n_x1 + w2 * n_x2;\
  219. float ny = w0 * n_y0 + w1 * n_y1 + w2 * n_y2;\
  220. float nz = w0 * n_z0 + w1 * n_z1 + w2 * n_z2;\
  221. nx = clamp((nx * 0.5f + 0.5f) * 255.0f, 0.0f, 255.0f);\
  222. ny = clamp((ny * 0.5f + 0.5f) * 255.0f, 0.0f, 255.0f);\
  223. nz = clamp((nz * 0.5f + 0.5f) * 255.0f, 0.0f, 255.0f);\
  224. \
  225. uchar r = (uchar)nx;\
  226. uchar g = (uchar)ny;\
  227. uchar b = (uchar)nz;\
  228. \
  229. if (!has_normals){\
  230. r = 20;\
  231. g = 20;\
  232. b = 20;\
  233. }\
  234. \
  235. frame_pixel = 0xFF000000 | (r << 16) | (g << 8) | b;\
  236. \
  237. break;}\
  238. case 3:{\
  239. float ux = w0 * u_x0 + w1 * u_x1 + w2 * u_x2;\
  240. float uy = w0 * u_y0 + w1 * u_y1 + w2 * u_y2;\
  241. \
  242. uchar r = (uchar)clamp(ux * 255.0f, 0.0f, 255.0f);\
  243. uchar g = (uchar)clamp(uy * 255.0f, 0.0f, 255.0f);\
  244. uchar b = 0;\
  245. \
  246. if (!has_uvs){\
  247. r = 20;\
  248. g = 20;\
  249. b = 20;\
  250. }\
  251. \
  252. frame_pixel = 0xFF000000 | (r << 16) | (g << 8) | b;\
  253. \
  254. break;}\
  255. default:{\
  256. frame_pixel = 0xFF00FFFF;\
  257. \
  258. break;}\
  259. }\
  260. } \
  261. }\
  262. }\
  263. }\
  264. \
  265. int pixel_coord = (height * 0.5 - id_y) * width + id_x + width * 0.5;\
  266. \
  267. if (pixel_coord >= width * height || pixel_coord < 0){\
  268. return;\
  269. }\
  270. frame_buffer[pixel_coord] = frame_pixel; \
  271. \
  272. }\n";