Graph Framework
Loading...
Searching...
No Matches
cuda_context.hpp
Go to the documentation of this file.
1//------------------------------------------------------------------------------
6//------------------------------------------------------------------------------
7
8#ifndef cuda_context_h
9#define cuda_context_h
10
11#include <array>
12#include <cstring>
13
14#include <cuda.h>
15#include <nvrtc.h>
16
17#include "random.hpp"
18#include "timing.hpp"
19#include "piecewise.hpp"
20
22#define MAX_REG 256
23
24namespace gpu {
25//------------------------------------------------------------------------------
30//------------------------------------------------------------------------------
31 static void check_nvrtc_error(nvrtcResult result,
32 const std::string &name) {
33#ifndef NDEBUG
34 assert(result == NVRTC_SUCCESS && nvrtcGetErrorString(result));
35#endif
36 }
37
38//------------------------------------------------------------------------------
43//------------------------------------------------------------------------------
44 static void check_error(CUresult result,
45 const std::string &name) {
46#ifndef NDEBUG
47 const char *error;
48 cuGetErrorString(result, &error);
49 if (result != CUDA_SUCCESS) {
50 std::cerr << name << " " << std::string(error) << std::endl;
51 }
52 assert(result == CUDA_SUCCESS && error);
53#endif
54 }
55
56//------------------------------------------------------------------------------
58//------------------------------------------------------------------------------
59 static CUresult cuda_init() {
60 const CUresult result = cuInit(0);
61 check_error(result, "cuInit");
62 return result;
63 }
65 static const CUresult result = cuda_init();
66
67//------------------------------------------------------------------------------
72//------------------------------------------------------------------------------
73 template<jit::float_scalar T, bool SAFE_MATH=false>
74 class cuda_context {
75 private:
77 CUdevice device;
79 CUcontext context;
81 CUmodule module;
83 std::unordered_map<graph::leaf_node<T, SAFE_MATH> *, CUdeviceptr> kernel_arguments;
84#ifdef USE_CUDA_TEXTURES
86 std::unordered_map<void *, CUtexObject> texture_arguments;
87#endif
89 CUdeviceptr result_buffer;
91 CUdeviceptr offset_buffer;
93 CUstream stream;
94
95//------------------------------------------------------------------------------
100//------------------------------------------------------------------------------
101 void check_error_async(CUresult result,
102 const std::string &name) {
103 check_error(result, name);
104#ifndef NDEBUG
105 std::string async_name = name + "_async";
106 check_error(cuStreamSynchronize(stream), async_name);
107#endif
108 }
109
110 public:
112 constexpr static size_t random_state_scale = 3000;
114 constexpr static size_t random_state_size = 1024*random_state_scale;
115
117 int remaining_const_memory;
118
119//------------------------------------------------------------------------------
123//------------------------------------------------------------------------------
124 static size_t max_concurrency() {
125 int count;
126 check_error(cuDeviceGetCount(&count), "cuDeviceGetCount");
127 return count;
128 }
129
130//------------------------------------------------------------------------------
132//------------------------------------------------------------------------------
133 static std::string device_type() {
134 return "Cuda GPU";
135 }
136
137//------------------------------------------------------------------------------
141//------------------------------------------------------------------------------
142 cuda_context(const size_t index) : result_buffer(0), module(0), offset_buffer(0) {
143 check_error(cuDeviceGet(&device, index), "cuDeviceGet");
144 check_error(cuDevicePrimaryCtxRetain(&context, device), "cuDevicePrimaryCtxRetain");
145 check_error(cuCtxSetCurrent(context), "cuCtxSetCurrent");
146 check_error(cuCtxSetCacheConfig(CU_FUNC_CACHE_PREFER_L1), "cuCtxSetCacheConfig");
147 check_error(cuStreamCreate(&stream, CU_STREAM_DEFAULT), "cuStreamCreate");
148 check_error(cuDeviceGetAttribute(&remaining_const_memory,
149 CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY,
150 device), "cuDeviceGetAttribute");
151 }
152
153//------------------------------------------------------------------------------
155//------------------------------------------------------------------------------
156 ~cuda_context() {
157 if (module) {
158 check_error(cuModuleUnload(module), "cuModuleUnload");
159 module = 0;
160 }
161
162 for (auto &[key, value] : kernel_arguments) {
163 check_error(cuMemFree(value), "cuMemFree");
164 }
165
166#ifdef USE_CUDA_TEXTURES
167 for (auto &[key, value] : texture_arguments) {
168 CUDA_RESOURCE_DESC resource;
169 check_error(cuTexObjectGetResourceDesc(&resource, value),
170 "cuTexObjectGetResourceDesc");
171
172 check_error(cuArrayDestroy(resource.res.array.hArray), "cuArrayDestroy");
173 check_error(cuTexObjectDestroy(value), "cuTexObjectDestroy");
174 }
175#endif
176
177 if (result_buffer) {
178 check_error(cuMemFree(result_buffer), "cuMemFree");
179 result_buffer = 0;
180 }
181 if (offset_buffer) {
182 check_error(cuMemFree(offset_buffer), "cuMemFree");
183 offset_buffer = 0;
184 }
185
186 check_error(cuStreamDestroy(stream), "cuStreamDestroy");
187 check_error(cuDevicePrimaryCtxRelease(device), "cuDevicePrimaryCtxRelease");
188 }
189
190//------------------------------------------------------------------------------
196//------------------------------------------------------------------------------
197 void compile(const std::string kernel_source,
198 std::vector<std::string> names,
199 const bool add_reduction=false) {
200 if (add_reduction) {
201 names.push_back("max_reduction");
202 }
203
204 nvrtcProgram kernel_program;
205 check_nvrtc_error(nvrtcCreateProgram(&kernel_program,
206 kernel_source.c_str(),
207 NULL, 0, NULL, NULL),
208 "nvrtcCreateProgram");
209
210 for (std::string &name : names) {
211 check_nvrtc_error(nvrtcAddNameExpression(kernel_program,
212 name.c_str()),
213 "nvrtcAddNameExpression");
214 }
215 if (add_reduction) {
216 check_nvrtc_error(nvrtcAddNameExpression(kernel_program,
217 "max_reduction"),
218 "nvrtcAddNameExpression");
219 }
220
221 std::ostringstream arch;
222 int compute_version;
223 check_error(cuDeviceGetAttribute(&compute_version,
224 CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,
225 device), "cuDeviceGetAttribute");
226 arch << "--gpu-architecture=compute_";
227 arch << compute_version;
228 if (jit::verbose) {
229 std::cout << "CUDA GPU info." << std::endl;
230 std::cout << " Major compute capability : " << compute_version << std::endl;
231
232 int value;
233 check_error(cuDeviceGetAttribute(&value,
234 CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK,
235 device), "cuDeviceGetAttribute");
236
237 std::cout << " Max shared memory : " << value << std::endl;
238
239 check_error(cuDeviceGetAttribute(&value,
240 CU_DEVICE_ATTRIBUTE_WARP_SIZE,
241 device), "cuDeviceGetAttribute");
242
243 std::cout << " Warp size : " << value << std::endl;
244 }
245
246 check_error(cuDeviceGetAttribute(&compute_version,
247 CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,
248 device), "cuDeviceGetAttribute");
249 arch << compute_version;
250 if (jit::verbose) {
251 std::cout << " Minor compute capability : " << compute_version << std::endl;
252 }
253
254 char device_name[100];
255 check_error(cuDeviceGetName(device_name, 100, device), "cuDeviceGetName");
256 if (jit::verbose) {
257 std::cout << " Device name : " << device_name << std::endl;
258 }
259
260 const std::string temp = arch.str();
261 std::array<const char *, 8> options({
262 temp.c_str(),
263 "--std=c++17",
264 "--relocatable-device-code=false",
265 "--include-path=" CUDA_INCLUDE,
266 "--include-path=" HEADER_DIR,
267 "--extra-device-vectorization",
268 "--device-as-default-execution-space",
269 "--use_fast_math"
270 });
271
272 if (nvrtcCompileProgram(kernel_program, options.size(), options.data())) {
273 size_t log_size;
274 check_nvrtc_error(nvrtcGetProgramLogSize(kernel_program, &log_size),
275 "nvrtcGetProgramLogSize");
276
277 char *log = static_cast<char *> (malloc(log_size));
278 check_nvrtc_error(nvrtcGetProgramLog(kernel_program, log),
279 "nvrtcGetProgramLog");
280 std::cout << log << std::endl;
281 free(log);
282 std::cout << kernel_source << std::endl;
283 }
284
285 check_error(cuDeviceGetAttribute(&compute_version,
286 CU_DEVICE_ATTRIBUTE_MANAGED_MEMORY,
287 device), "cuDeviceGetAttribute");
288 if (jit::verbose) {
289 std::cout << " Managed Memory : " << compute_version << std::endl;
290 }
291
292 size_t ptx_size;
293 check_nvrtc_error(nvrtcGetPTXSize(kernel_program, &ptx_size),
294 "nvrtcGetPTXSize");
295
296 char *ptx = static_cast<char *> (malloc(ptx_size));
297 check_nvrtc_error(nvrtcGetPTX(kernel_program, ptx), "nvrtcGetPTX");
298
299 check_nvrtc_error(nvrtcDestroyProgram(&kernel_program),
300 "nvrtcDestroyProgram");
301
302 std::array<CUjit_option, 3> module_options = {
303 CU_JIT_MAX_REGISTERS,
304 CU_JIT_LTO,
305 CU_JIT_POSITION_INDEPENDENT_CODE
306 };
307 std::array<void *, 3> module_values = {
308 reinterpret_cast<void *> (MAX_REG),
309 reinterpret_cast<void *> (1),
310 reinterpret_cast<void *> (0)
311 };
312
313 check_error(cuModuleLoadDataEx(&module, ptx, module_options.size(),
314 module_options.data(),
315 module_values.data()), "cuModuleLoadDataEx");
316
317 free(ptx);
318 }
319
320//------------------------------------------------------------------------------
332//------------------------------------------------------------------------------
333 std::function<void(void)> create_kernel_call(const std::string kernel_name,
334 graph::input_nodes<T, SAFE_MATH> inputs,
335 graph::output_nodes<T, SAFE_MATH> outputs,
336 graph::input_nodes<T, SAFE_MATH> atomics,
337 graph::shared_random_state<T, SAFE_MATH> state,
338 const size_t num_rays,
339 const jit::texture1d_list &tex1d_list,
340 const jit::texture2d_list &tex2d_list) {
341 CUfunction function;
342 check_error(cuModuleGetFunction(&function, module, kernel_name.c_str()), "cuModuleGetFunction");
343
344 std::vector<void *> buffers;
345 std::unordered_set<graph::leaf_node<T, SAFE_MATH> *> needed_buffers;
346
347 const size_t buffer_element_size = sizeof(T);
348 for (auto &input : inputs) {
349 if (!kernel_arguments.contains(input.get())) {
350 kernel_arguments.try_emplace(input.get());
351 check_error(cuMemAllocManaged(&kernel_arguments[input.get()],
352 input->size()*sizeof(T),
353 CU_MEM_ATTACH_GLOBAL),
354 "cuMemAllocManaged");
355 check_error(cuMemcpyHtoD(kernel_arguments[input.get()],
356 input->data(),
357 input->size()*sizeof(T)),
358 "cuMemcpyHtoD");
359 }
360 if (!needed_buffers.contains(input.get())) {
361 buffers.push_back(reinterpret_cast<void *> (&kernel_arguments[input.get()]));
362 needed_buffers.insert(input.get());
363 }
364 }
365 for (auto &output : outputs) {
366 if (!graph::atomic_accumulate_1D_cast(output).get()) {
367 if (!kernel_arguments.contains(output.get())) {
368 kernel_arguments.try_emplace(output.get());
369 check_error(cuMemAllocManaged(&kernel_arguments[output.get()],
370 num_rays*sizeof(T),
371 CU_MEM_ATTACH_GLOBAL),
372 "cuMemAllocManaged");
373 }
374 if (!needed_buffers.contains(output.get())) {
375 buffers.push_back(reinterpret_cast<void *> (&kernel_arguments[output.get()]));
376 needed_buffers.insert(output.get());
377 }
378 }
379 }
380 for (auto &atomic : atomics) {
381 if (!kernel_arguments.contains(atomic.get())) {
382 kernel_arguments.try_emplace(atomic.get());
383 check_error(cuMemAllocManaged(&kernel_arguments[atomic.get()],
384 atomic->size()*sizeof(T),
385 CU_MEM_ATTACH_GLOBAL),
386 "cuMemAllocManaged");
387 check_error(cuMemcpyHtoD(kernel_arguments[atomic.get()],
388 atomic->data(),
389 atomic->size()*sizeof(T)),
390 "cuMemcpyHtoD");
391 }
392 if (!needed_buffers.contains(atomic.get())) {
393 buffers.push_back(reinterpret_cast<void *> (&kernel_arguments[atomic.get()]));
394 needed_buffers.insert(atomic.get());
395 }
396 }
397
398 const size_t num_buffers = buffers.size();
399 if (state.get()) {
400 if (!kernel_arguments.contains(state.get())) {
401 kernel_arguments.try_emplace(state.get());
402 check_error(cuMemAllocManaged(&kernel_arguments[state.get()],
403 state->get_size_bytes(),
404 CU_MEM_ATTACH_GLOBAL),
405 "cuMemAllocManaged");
406 check_error(cuMemAlloc(&offset_buffer, sizeof(uint32_t)), "cuMemAlloc");
407 check_error(cuMemcpyHtoD(kernel_arguments[state.get()],
408 state->data(),
409 state->get_size_bytes()),
410 "cuMemcpyHtoD");
411 }
412 buffers.push_back(reinterpret_cast<void *> (&kernel_arguments[state.get()]));
413 buffers.push_back(reinterpret_cast<void *> (&offset_buffer));
414 }
415
416#ifdef USE_CUDA_TEXTURES
417 for (auto &[data, size] : tex1d_list) {
418 if (!texture_arguments.contains(data)) {
419 texture_arguments.try_emplace(data);
420 CUDA_RESOURCE_DESC resource_desc;
421 CUDA_TEXTURE_DESC texture_desc;
422 CUDA_ARRAY_DESCRIPTOR array_desc;
423
424 array_desc.Width = size;
425 array_desc.Height = 1;
426
427 memset(&resource_desc, 0, sizeof(CUDA_RESOURCE_DESC));
428 memset(&texture_desc, 0, sizeof(CUDA_TEXTURE_DESC));
429
430 resource_desc.resType = CU_RESOURCE_TYPE_ARRAY;
431 texture_desc.addressMode[0] = CU_TR_ADDRESS_MODE_BORDER;
432 texture_desc.addressMode[1] = CU_TR_ADDRESS_MODE_BORDER;
433 texture_desc.addressMode[2] = CU_TR_ADDRESS_MODE_BORDER;
434 if constexpr (jit::float_base<T>) {
435 array_desc.Format = CU_AD_FORMAT_FLOAT;
436 if constexpr (jit::complex_scalar<T>) {
437 array_desc.NumChannels = 2;
438 } else {
439 array_desc.NumChannels = 1;
440 }
441 } else {
442 array_desc.Format = CU_AD_FORMAT_UNSIGNED_INT32;
443 if constexpr (jit::complex_scalar<T>) {
444 array_desc.NumChannels = 4;
445 } else {
446 array_desc.NumChannels = 2;
447 }
448 }
449 check_error(cuArrayCreate(&resource_desc.res.array.hArray, &array_desc),
450 "cuArrayCreate");
451 check_error(cuMemcpyHtoA(resource_desc.res.array.hArray, 0, data,
452 size*sizeof(float)*array_desc.NumChannels),
453 "cuMemcpyHtoA");
454
455 check_error(cuTexObjectCreate(&texture_arguments[data],
456 &resource_desc, &texture_desc,
457 NULL),
458 "cuTexObjectCreate");
459 }
460 buffers.push_back(reinterpret_cast<void *> (&texture_arguments[data]));
461 }
462 for (auto &[data, size] : tex2d_list) {
463 if (!texture_arguments.contains(data)) {
464 texture_arguments.try_emplace(data);
465 CUDA_RESOURCE_DESC resource_desc;
466 CUDA_TEXTURE_DESC texture_desc;
467 CUDA_ARRAY_DESCRIPTOR array_desc;
468
469 array_desc.Width = size[0];
470 array_desc.Height = size[1];
471
472 memset(&resource_desc, 0, sizeof(CUDA_RESOURCE_DESC));
473 memset(&texture_desc, 0, sizeof(CUDA_TEXTURE_DESC));
474
475 resource_desc.resType = CU_RESOURCE_TYPE_ARRAY;
476 texture_desc.addressMode[0] = CU_TR_ADDRESS_MODE_BORDER;
477 texture_desc.addressMode[1] = CU_TR_ADDRESS_MODE_BORDER;
478 texture_desc.addressMode[2] = CU_TR_ADDRESS_MODE_BORDER;
479 const size_t total = size[0]*size[1];
480 if constexpr (jit::float_base<T>) {
481 array_desc.Format = CU_AD_FORMAT_FLOAT;
482 if constexpr (jit::complex_scalar<T>) {
483 array_desc.NumChannels = 2;
484 } else {
485 array_desc.NumChannels = 1;
486 }
487 } else {
488 array_desc.Format = CU_AD_FORMAT_UNSIGNED_INT32;
489 if constexpr (jit::complex_scalar<T>) {
490 array_desc.NumChannels = 4;
491 } else {
492 array_desc.NumChannels = 2;
493 }
494 }
495 check_error(cuArrayCreate(&resource_desc.res.array.hArray, &array_desc),
496 "cuArrayCreate");
497
498 CUDA_MEMCPY2D copy_desc;
499 memset(&copy_desc, 0, sizeof(copy_desc));
500
501 copy_desc.srcPitch = size[0]*sizeof(float)*array_desc.NumChannels;
502 copy_desc.srcMemoryType = CU_MEMORYTYPE_HOST;
503 copy_desc.srcHost = data;
504
505 copy_desc.dstMemoryType = CU_MEMORYTYPE_ARRAY;
506 copy_desc.dstArray = resource_desc.res.array.hArray;
507
508 copy_desc.WidthInBytes = copy_desc.srcPitch;
509 copy_desc.Height = size[0];
510
511 check_error(cuMemcpy2D(&copy_desc), "cuMemcpy2D");
512
513 check_error(cuTexObjectCreate(&texture_arguments[data],
514 &resource_desc, &texture_desc,
515 NULL),
516 "cuTexObjectCreate");
517 }
518 buffers.push_back(reinterpret_cast<void *> (&texture_arguments[data]));
519 }
520#endif
521
522 int value;
523 check_error(cuFuncGetAttribute(&value, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK,
524 function), "cuFuncGetAttribute");
525 int warp_size;
526 check_error(cuDeviceGetAttribute(&warp_size,
527 CU_DEVICE_ATTRIBUTE_WARP_SIZE,
528 device), "cuDeviceGetAttribute");
529
530 unsigned int total_parallel = state.get() ? state->size() : num_rays;
531 unsigned int threads_per_group = total_parallel < 1024 ? warp_size : value;
532 unsigned int thread_groups = total_parallel/threads_per_group + (total_parallel%threads_per_group ? 1 : 0);
533
534 int min_grid;
535 check_error(cuOccupancyMaxPotentialBlockSize(&min_grid, &value, function, 0, 0, 0),
536 "cuOccupancyMaxPotentialBlockSize");
537
538 if (jit::verbose) {
539 std::cout << " Kernel name : " << kernel_name << std::endl;
540 std::cout << " Threads per group : " << threads_per_group << std::endl;
541 std::cout << " Number of groups : " << thread_groups << std::endl;
542 std::cout << " Total problem size : " << threads_per_group*thread_groups << std::endl;
543 std::cout << " Total parallel : " << total_parallel << std::endl;
544 std::cout << " Min grid size : " << min_grid << std::endl;
545 std::cout << " Suggested Block size : " << value << std::endl;
546 }
547#ifdef PROFILE_KERNELS
548 timing::measure_diagnostic timer(kernel_name);
549#endif
550 if (state.get()) {
551 return [this, num_rays, function, thread_groups, threads_per_group, buffers
552#ifdef PROFILE_KERNELS
553 , timer
554#endif
555 ] () mutable {
556#ifdef PROFILE_KERNELS
557 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
558 reinterpret_cast<timing::measure_diagnostic *> (arg)->reset();
559 }, &timer), "cuLaunchHostFunc");
560#endif
561 for (uint32_t i = 0, ie = threads_per_group*thread_groups; i < num_rays; i += ie) {
562 check_error_async(cuStreamWriteValue32(stream, offset_buffer, i,
563 CU_STREAM_WRITE_VALUE_DEFAULT),
564 "cuStreamWriteValue32");
565 check_error_async(cuLaunchKernel(function,
566 thread_groups, 1, 1,
567 threads_per_group, 1, 1,
568 0, stream,
569 buffers.data(), NULL),
570 "cuLaunchKernel");
571 }
572#ifdef PROFILE_KERNELS
573 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
574 reinterpret_cast<timing::measure_diagnostic *> (arg)->print();
575 }, &timer), "cuLaunchHostFunc");
576#endif
577 };
578 } else {
579 return [this, function, thread_groups, threads_per_group, buffers
580#ifdef PROFILE_KERNELS
581 , timer
582#endif
583 ] () mutable {
584#ifdef PROFILE_KERNELS
585 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
586 reinterpret_cast<timing::measure_diagnostic *> (arg)->reset();
587 }, &timer), "cuLaunchHostFunc");
588#endif
589 check_error_async(cuLaunchKernel(function, thread_groups, 1, 1,
590 threads_per_group, 1, 1, 0, stream,
591 buffers.data(), NULL),
592 "cuLaunchKernel");
593#ifdef PROFILE_KERNELS
594 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
595 reinterpret_cast<timing::measure_diagnostic *> (arg)->print();
596 }, &timer), "cuLaunchHostFunc");
597#endif
598 };
599 }
600 }
601
602//------------------------------------------------------------------------------
608//------------------------------------------------------------------------------
609 std::function<T(void)> create_max_call(graph::shared_leaf<T, SAFE_MATH> &argument,
610 std::function<void(void)> run) {
611 check_error(cuMemAllocManaged(&result_buffer, sizeof(T),
612 CU_MEM_ATTACH_GLOBAL),
613 "cuMemAllocManaged");
614
615 std::vector<void *> buffers;
616
617 buffers.push_back(reinterpret_cast<void *> (&kernel_arguments[argument.get()]));
618 buffers.push_back(reinterpret_cast<void *> (&result_buffer));
619
620 CUfunction function;
621 check_error(cuModuleGetFunction(&function, module, "max_reduction"),
622 "cuModuleGetFunction");
623
624 int value;
625 int min_grid;
626 check_error(cuOccupancyMaxPotentialBlockSize(&min_grid, &value, function, 0, 0, 0),
627 "cuOccupancyMaxPotentialBlockSize");
628
629 if (jit::verbose) {
630 std::cout << " Kernel name : max_reduction" << std::endl;
631 std::cout << " Min grid size : " << min_grid << std::endl;
632 std::cout << " Suggested Block size : " << value << std::endl;
633 }
634
635 return [this, function, run, buffers] () mutable {
636 run();
637 check_error_async(cuLaunchKernel(function, 1, 1, 1,
638 1024, 1, 1, 0, stream,
639 buffers.data(), NULL),
640 "cuLaunchKernel");
641 wait();
642
643 return reinterpret_cast<T *> (result_buffer)[0];
644 };
645 }
646
647//------------------------------------------------------------------------------
652//------------------------------------------------------------------------------
653 std::function<void(void)> create_zero_call(graph::input_nodes<T, SAFE_MATH> &inputs) {
654 std::vector<CUdeviceptr> buffers;
655 for (auto &input : inputs) {
656 if (!kernel_arguments.contains(input.get())) {
657 kernel_arguments.try_emplace(input.get());
658 check_error(cuMemAllocManaged(&kernel_arguments[input.get()],
659 input->size()*sizeof(T),
660 CU_MEM_ATTACH_GLOBAL),
661 "cuMemAllocManaged");
662 check_error(cuMemcpyHtoD(kernel_arguments[input.get()],
663 input->data(),
664 input->size()*sizeof(T)),
665 "cuMemcpyHtoD");
666 }
667 buffers.push_back(kernel_arguments[input.get()]);
668 }
669
670 std::vector<size_t> sizes;
671 for (CUdeviceptr &buffer : buffers) {
672 size_t size;
673 check_error(cuMemGetAddressRange(NULL, &size, buffer),
674 "cuMemGetAddressRange");
675 sizes.push_back(size);
676 }
677#ifdef PROFILE_KERNELS
678 timing::measure_diagnostic timer("zero buffer");
679#endif
680 return [this, buffers, sizes
681#ifdef PROFILE_KERNELS
682 , timer
683#endif
684 ] () mutable {
685#ifdef PROFILE_KERNELS
686 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
687 reinterpret_cast<timing::measure_diagnostic *> (arg)->reset();
688 }, &timer), "cuLaunchHostFunc");
689#endif
690 for (size_t i = 0, ie = buffers.size(); i < ie; i++) {
691 check_error_async(cuMemsetD8Async(buffers[i], 0, sizes[i],
692 stream),
693 "cuMemsetD8Async");
694 }
695#ifdef PROFILE_KERNELS
696 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
697 reinterpret_cast<timing::measure_diagnostic *> (arg)->print();
698 }, &timer), "cuLaunchHostFunc");
699#endif
700 };
701 }
702
703//------------------------------------------------------------------------------
708//------------------------------------------------------------------------------
709 std::function<void(void)> create_copy_call(graph::copy_nodes<T, SAFE_MATH> &setters) {
710 std::vector<CUdeviceptr> sources;
711 std::vector<CUdeviceptr> destinations;
712
713 for (auto &[out, in] : setters) {
714 if (!kernel_arguments.contains(in.get())) {
715 kernel_arguments.try_emplace(in.get());
716 check_error(cuMemAllocManaged(&kernel_arguments[in.get()],
717 in->size()*sizeof(T),
718 CU_MEM_ATTACH_GLOBAL),
719 "cuMemAllocManaged");
720 check_error(cuMemcpyHtoD(kernel_arguments[in.get()],
721 in->data(),
722 in->size()*sizeof(T)),
723 "cuMemcpyHtoD");
724 }
725 destinations.push_back(kernel_arguments[in.get()]);
726
727 if (!kernel_arguments.contains(out.get())) {
728 kernel_arguments.try_emplace(out.get());
729 check_error(cuMemAllocManaged(&kernel_arguments[out.get()],
730 out->size()*sizeof(T),
731 CU_MEM_ATTACH_GLOBAL),
732 "cuMemAllocManaged");
733 check_error(cuMemcpyHtoD(kernel_arguments[out.get()],
734 out->data(),
735 out->size()*sizeof(T)),
736 "cuMemcpyHtoD");
737 }
738 sources.push_back(kernel_arguments[out.get()]);
739 }
740
741 std::vector<size_t> sizes;
742 for (CUdeviceptr &buffer : sources) {
743 size_t size;
744 check_error(cuMemGetAddressRange(NULL, &size, buffer),
745 "cuMemGetAddressRange");
746 sizes.push_back(size);
747 }
748#ifdef PROFILE_KERNELS
749 timing::measure_diagnostic timer("copy buffer");
750#endif
751 return [this, sources, destinations, sizes
752#ifdef PROFILE_KERNELS
753 , timer
754#endif
755 ] () mutable {
756#ifdef PROFILE_KERNELS
757 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
758 reinterpret_cast<timing::measure_diagnostic *> (arg)->reset();
759 }, &timer), "cuLaunchHostFunc");
760#endif
761 for (size_t i = 0, ie = sources.size(); i < ie; i++) {
762 check_error_async(cuMemcpyDtoDAsync(destinations[i],
763 sources[i],
764 sizes[i], stream),
765 "cuMemcpyDtoDAsync");
766 }
767#ifdef PROFILE_KERNELS
768 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
769 reinterpret_cast<timing::measure_diagnostic *> (arg)->print();
770 }, &timer), "cuLaunchHostFunc");
771#endif
772 };
773 }
774
775//------------------------------------------------------------------------------
777//------------------------------------------------------------------------------
778 void wait() {
779 check_error_async(cuStreamSynchronize(stream), "cuStreamSynchronize");
780 check_error(cuCtxSynchronize(), "cuCtxSynchronize");
781 }
782
783//------------------------------------------------------------------------------
788//------------------------------------------------------------------------------
789 std::function<void(void)> run_function(std::function<void(void)> callback) {
790#ifdef PROFILE_KERNELS
791 timing::measure_diagnostic timer("callback");
792#endif
793 return [this, callback
794#ifdef PROFILE_KERNELS
795 , timer
796#endif
797 ]() mutable {
798#ifdef PROFILE_KERNELS
799 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
800 reinterpret_cast<timing::measure_diagnostic *> (arg)->reset();
801 }, &timer), "cuLaunchHostFunc");
802#endif
803 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
804 reinterpret_cast<std::function<void(void)> *> (arg)->operator()();
805 }, &callback), "cuLaunchHostFunc");
806#ifdef PROFILE_KERNELS
807 check_error_async(cuLaunchHostFunc(stream, [](void *arg) {
808 reinterpret_cast<timing::measure_diagnostic *> (arg)->print();
809 }, &timer), "cuLaunchHostFunc");
810#endif
811 };
812 }
813
814//------------------------------------------------------------------------------
819//------------------------------------------------------------------------------
820 void print_results(const size_t index,
821 const graph::output_nodes<T, SAFE_MATH> &nodes) {
822 wait();
823 for (auto &out : nodes) {
824 const T temp = reinterpret_cast<T *> (kernel_arguments[out.get()])[index];
825 if constexpr (jit::complex_scalar<T>) {
826 std::cout << std::real(temp) << " " << std::imag(temp) << " ";
827 } else {
828 std::cout << temp << " ";
829 }
830 }
831 std::cout << std::endl;
832 }
833
834//------------------------------------------------------------------------------
840//------------------------------------------------------------------------------
841 T check_value(const size_t index,
842 const graph::shared_leaf<T, SAFE_MATH> &node) {
843 wait();
844 return reinterpret_cast<T *> (kernel_arguments[node.get()])[index];
845 }
846
847//------------------------------------------------------------------------------
852//------------------------------------------------------------------------------
853 void copy_to_device(graph::shared_leaf<T, SAFE_MATH> node,
854 T *source) {
855 size_t size;
856 check_error(cuMemGetAddressRange(NULL, &size, kernel_arguments[node.get()]), "cuMemGetAddressRange");
857 check_error_async(cuMemcpyHtoDAsync(kernel_arguments[node.get()], source, size, stream), "cuMemcpyHtoDAsync");
858 }
859
860//------------------------------------------------------------------------------
865//------------------------------------------------------------------------------
866 void copy_to_host(graph::shared_leaf<T, SAFE_MATH> node,
867 T *destination) {
868 size_t size;
869 check_error(cuMemGetAddressRange(NULL, &size, kernel_arguments[node.get()]), "cuMemGetAddressRange");
870 check_error_async(cuMemcpyDtoHAsync(destination, kernel_arguments[node.get()], size, stream), "cuMemcpyDtoHAsync");
871 }
872
873//------------------------------------------------------------------------------
877//------------------------------------------------------------------------------
878 void create_header(std::ostringstream &source_buffer) {
879 source_buffer << "typedef unsigned int uint32_t;" << std::endl
880 << "typedef unsigned short uint16_t;" << std::endl
881 << "typedef short int16_t;" << std::endl
882 << "typedef unsigned char uint8_t;" << std::endl
883 << "template<typename T, size_t S>" << std::endl
884 << "class array {" << std::endl
885 << "private:" << std::endl
886 << " T _buffer[S];" << std::endl
887 << "public:" << std::endl
888 << " T operator[] (const size_t index) const {" << std::endl
889 << " return _buffer[index];" << std::endl
890 << " }" << std::endl
891 << " T &operator[] (const size_t index) {" << std::endl
892 << " return _buffer[index];" << std::endl
893 << " }" << std::endl
894 << "};" << std::endl;
895 if constexpr (jit::complex_scalar<T>) {
896 source_buffer << "#define CUDA_DEVICE_CODE" << std::endl
897 << "#define M_PI " << M_PI << std::endl
898 << "#include <cuda/std/complex>" << std::endl
899 << "#include <special_functions.hpp>" << std::endl;
900#ifdef USE_CUDA_TEXTURES
901 if constexpr (jit::float_base<T>) {
902 source_buffer << "static __inline__ __device__ complex<float> to_cmp_float(float2 p) {"
903 << std::endl
904 << " return ";
905 jit::add_type<T> (source_buffer);
906 source_buffer << " (p.x, p.y);" << std::endl
907 << "}" << std::endl;
908 } else {
909 source_buffer << "static __inline__ __device__ complex<double> to_cmp_double(uint4 p) {"
910 << std::endl
911 << " return ";
912 jit::add_type<T> (source_buffer);
913 source_buffer << " (__hiloint2double(p.y, p.x), __hiloint2double(p.w, p.z));"
914 << std::endl
915 << "}" << std::endl;
916 }
917 } else if constexpr (jit::double_base<T>) {
918 source_buffer << "static __inline__ __device__ double to_double(uint2 p) {"
919 << std::endl
920 << " return __hiloint2double(p.y, p.x);"
921 << std::endl
922 << "}" << std::endl;
923#endif
924 }
925 }
926
927//------------------------------------------------------------------------------
945//------------------------------------------------------------------------------
946 void create_kernel_prefix(std::ostringstream &source_buffer,
947 const std::string name,
948 graph::input_nodes<T, SAFE_MATH> &inputs,
949 graph::output_nodes<T, SAFE_MATH> &outputs,
950 graph::input_nodes<T, SAFE_MATH> atomics,
951 graph::shared_random_state<T, SAFE_MATH> state,
952 const size_t size,
953 const std::vector<bool> &is_constant,
954 jit::register_map &registers,
955 const jit::register_usage &usage,
956 jit::texture1d_list &textures1d,
957 jit::texture2d_list &textures2d,
958 jit::argument_set &thread_shared,
959 jit::register_map &thread_mem,
960 const size_t iterations=1) {
961 source_buffer << std::endl;
962 source_buffer << "extern \"C\" __global__ void "
963 << name << "(" << std::endl;
964
965 int used_thread_mem = 0;
966 int max_shared_mem;
967 check_error(cuDeviceGetAttribute(&max_shared_mem,
968 CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK,
969 device), "cuDeviceGetAttribute");
970
971 int warp_size;
972 check_error(cuDeviceGetAttribute(&warp_size,
973 CU_DEVICE_ATTRIBUTE_WARP_SIZE,
974 device), "cuDeviceGetAttribute");
975
976 jit::argument_set used_args;
977 if (inputs.size()) {
978#ifdef USE_INPUT_CACHE
979 if (!is_constant[0] && iterations > 1) {
980 const size_t needed_mem = inputs[0]->size() > 1024 ?
981 1024*sizeof(T) :
982 warp_size*sizeof(T);
983 if (used_thread_mem + needed_mem < max_shared_mem) {
984 used_thread_mem += needed_mem;
985 thread_shared.insert(inputs[0].get());
986 }
987 } else if (is_constant[0] &&
988 inputs[0]->size() < size &&
989 inputs[0]->size() < 1024) {
990 const size_t needed_mem = inputs[0]->size()*sizeof(T);
991 if (used_thread_mem + needed_mem < max_shared_mem) {
992 used_thread_mem += needed_mem;
993 thread_shared.insert(inputs[0].get());
994 thread_mem[inputs[0].get()] = jit::to_string('t', inputs[0].get());
995 }
996 }
997#endif
998 source_buffer << " ";
999 if (is_constant[0]) {
1000 source_buffer << "const ";
1001 }
1002 jit::add_type<T> (source_buffer);
1003 source_buffer << " * __restrict__ "
1004 << jit::to_string('v', inputs[0].get());
1005 used_args.insert(inputs[0].get());
1006 }
1007 for (size_t i = 1, ie = inputs.size(); i < ie; i++) {
1008#ifdef USE_INPUT_CACHE
1009 if (!is_constant[i] && iterations > 1) {
1010 const size_t needed_mem = inputs[i]->size() > 1024 ?
1011 1024*sizeof(T) :
1012 warp_size*sizeof(T);
1013 if (used_thread_mem + needed_mem < max_shared_mem) {
1014 used_thread_mem += needed_mem;
1015 thread_shared.insert(inputs[i].get());
1016 }
1017 } else if (is_constant[i] &&
1018 inputs[i]->size() < size &&
1019 inputs[i]->size() < 1024) {
1020 const size_t needed_mem = inputs[i]->size()*sizeof(T);
1021 if (used_thread_mem + needed_mem < max_shared_mem) {
1022 used_thread_mem += needed_mem;
1023 thread_shared.insert(inputs[i].get());
1024 thread_mem[inputs[i].get()] = jit::to_string('t', inputs[i].get());
1025 }
1026 }
1027#endif
1028 if (!used_args.contains(inputs[i].get())) {
1029 inputs[i]->endline(source_buffer, usage, ',');
1030 source_buffer << " ";
1031 if (is_constant[i]) {
1032 source_buffer << "const ";
1033 }
1034 jit::add_type<T> (source_buffer);
1035 source_buffer << " * __restrict__ "
1036 << jit::to_string('v', inputs[i].get());
1037 used_args.insert(inputs[i].get());
1038 }
1039 }
1040 for (size_t i = 0, ie = outputs.size(); i < ie; i++) {
1041 if (!used_args.contains(outputs[i].get()) &&
1042 !graph::atomic_accumulate_1D_cast(outputs[i]).get()) {
1043 if (i == 0) {
1044 if (inputs.size()) {
1045 inputs[inputs.size() - 1]->endline(source_buffer,
1046 usage, ',');
1047 }
1048 } else {
1049 source_buffer << "," << std::endl;
1050 }
1051
1052 source_buffer << " ";
1053 jit::add_type<T> (source_buffer);
1054 source_buffer << " * __restrict__ "
1055 << jit::to_string('o', outputs[i].get());
1056 used_args.insert(outputs[i].get());
1057 }
1058 }
1059 for (size_t i = 0, ie = atomics.size(); i < ie; i++) {
1060 if (!used_args.contains(atomics[i].get())) {
1061 if (i == 0) {
1062 if (outputs.size()) {
1063 outputs[outputs.size() - 1]->endline(source_buffer,
1064 usage, ',');
1065 } else if (inputs.size()) {
1066 inputs[inputs.size() - 1]->endline(source_buffer,
1067 usage, ',');
1068 }
1069 } else {
1070 source_buffer << "," << std::endl;
1071 }
1072
1073 source_buffer << " ";
1074 jit::add_type<T> (source_buffer);
1075 source_buffer << " * __restrict__ "
1076 << jit::to_string('v', atomics[i].get());
1077 used_args.insert(atomics[i].get());
1078 }
1079 }
1080 if (state.get()) {
1081 source_buffer << "," << std::endl
1082 << " mt_state * __restrict__ "
1083 << jit::to_string('s', state.get())
1084 << "," << std::endl
1085 << " const uint32_t * __restrict__ offset"
1086 << std::endl;
1087 }
1088#ifdef USE_CUDA_TEXTURES
1089 for (auto &[key, value] : textures1d) {
1090 source_buffer << "," << std::endl;
1091 source_buffer << " cudaTextureObject_t "
1092 << jit::to_string('a', key);
1093 }
1094 for (auto &[key, value] : textures2d) {
1095 source_buffer << "," << std::endl;
1096 source_buffer << " cudaTextureObject_t "
1097 << jit::to_string('a', key);
1098 }
1099#endif
1100 source_buffer << ") {" << std::endl;
1101 if (thread_shared.size()) {
1102 source_buffer << " const int t_index = threadIdx.x;"
1103 << std::endl;
1104 }
1105 source_buffer << " const int index = blockIdx.x*blockDim.x + threadIdx.x;"
1106 << std::endl;
1107
1108 source_buffer << " if (";
1109 if (state.get()) {
1110 source_buffer << "offset[0] + ";
1111 }
1112 source_buffer << "index < " << size << ") {" << std::endl;
1113
1114 for (size_t i = 0, ie = inputs.size(); i < ie; i++) {
1115 if (is_constant[i]) {
1116#ifdef USE_INPUT_CACHE
1117 if (usage.at(inputs[i].get()) && inputs[i]->size() == size) {
1118 registers[inputs[i].get()] = jit::to_string('r', inputs[i].get());
1119 source_buffer << " const ";
1120 jit::add_type<T> (source_buffer);
1121 source_buffer << " " << registers[inputs[i].get()] << " = "
1122 << jit::to_string('v', inputs[i].get())
1123 << "[index]";
1124 inputs[i]->endline(source_buffer, usage);
1125 }
1126#else
1127 registers[inputs[i].get()] = jit::to_string('v', inputs[i].get())
1128 + "["
1129 + (state.get() ? "offset[0] + " : "")
1130 + "index]";
1131#endif
1132 }
1133 }
1134
1135 if (thread_shared.size()) {
1136 for (size_t i = 0, ie = inputs.size(); i < ie; i++) {
1137 if (thread_shared.contains(inputs[i].get()) && is_constant[i]) {
1138 source_buffer << " __shared__ ";
1139 jit::add_type<T> (source_buffer);
1140 source_buffer << " " << jit::to_string('t', inputs[i].get())
1141 << "[" << inputs[i]->size() << "]";
1142 inputs[i]->endline(source_buffer, usage);
1143 }
1144 }
1145 for (size_t i = 0, ie = inputs.size(); i < ie; i++) {
1146 if (thread_shared.contains(inputs[i].get()) && is_constant[i]) {
1147 source_buffer << " for(int j = t_index; j < "
1148 << inputs[i]->size()
1149 << "; j += blockDim.x) {" << std::endl
1150 << " "
1151 << jit::to_string('t', inputs[i].get())
1152 << "[j] = "
1153 << jit::to_string('v', inputs[i].get())
1154 << "[j]";
1155 inputs[i]->endline(source_buffer, usage);
1156 source_buffer << " }" << std::endl;
1157 }
1158 }
1159 for (size_t i = 0, ie = inputs.size(); i < ie; i++) {
1160 if (thread_shared.contains(inputs[i].get()) && is_constant[i]) {
1161 source_buffer << " __syncthreads();"
1162 << std::endl;
1163 break;
1164 }
1165 }
1166 for (size_t i = 0, ie = inputs.size(); i < ie; i++) {
1167 if (thread_shared.contains(inputs[i].get()) && is_constant[i]) {
1168 thread_shared.erase(inputs[i].get());
1169 }
1170 }
1171 for (size_t i = 0, ie = inputs.size(); i < ie; i++) {
1172 if (thread_shared.contains(inputs[i].get()) && !is_constant[i]) {
1173 source_buffer << " __shared__ ";
1174 jit::add_type<T> (source_buffer);
1175 source_buffer << " " << jit::to_string('t', inputs[i].get())
1176 << "["
1177 << (inputs[i]->size() > 1024 ? 1024 : warp_size)
1178 << "]";
1179 inputs[i]->endline(source_buffer, usage);
1180 source_buffer << " "
1181 << jit::to_string('t', inputs[i].get())
1182 << "[t_index] = "
1183 << jit::to_string('v', inputs[i].get())
1184 << "[index]";
1185 inputs[i]->endline(source_buffer, usage);
1186 }
1187 }
1188 }
1189
1190 if (iterations > 1) {
1191 source_buffer << " for (size_t j = 0; j < " << iterations << "; j++) {" << std::endl;
1192 }
1193
1194 for (size_t i = 0, ie = inputs.size(); i < ie; i++) {
1195 if (!is_constant[i]) {
1196#ifdef USE_INPUT_CACHE
1197 if (usage.at(inputs[i].get())) {
1198 registers[inputs[i].get()] = jit::to_string('r', inputs[i].get());
1199 source_buffer << " const ";
1200 jit::add_type<T> (source_buffer);
1201 source_buffer << " " << registers[inputs[i].get()] << " = ";
1202 if (thread_shared.contains(inputs[i].get())) {
1203 source_buffer << jit::to_string('t', inputs[i].get())
1204 << "[t_index]";
1205 } else {
1206 source_buffer << jit::to_string('v', inputs[i].get())
1207 << "[";
1208 if (state.get()) {
1209 source_buffer << "offset[0] + ";
1210 }
1211 source_buffer << "index]";
1212 }
1213 inputs[i]->endline(source_buffer, usage);
1214 }
1215#else
1216 if (thread_shared.contains(inputs[i].get())) {
1217 registers[inputs[i].get()] = jit::to_string('t', inputs[i].get()) + "[t_index]";
1218 } else {
1219 registers[inputs[i].get()] = jit::to_string('v', inputs[i].get()) + "["
1220 + (state.get() ? "offset[0] + " : "")
1221 + "index]";
1222 }
1223#endif
1224 }
1225 }
1226 if (state.get()) {
1227#ifdef USE_INPUT_CACHE
1228 registers[state.get()] = jit::to_string('r', state.get());
1229 source_buffer << " mt_state &" << registers[state.get()] << " = "
1230 << jit::to_string('s', state.get())
1231 << "[index]";
1232 state->endline(source_buffer, usage);
1233#else
1234 registers[state.get()] = jit::to_string('s', state.get()) + "[threadIdx.x]";
1235#endif
1236 }
1237 }
1238
1239//------------------------------------------------------------------------------
1251//------------------------------------------------------------------------------
1252 void create_kernel_postfix(std::ostringstream &source_buffer,
1253 graph::output_nodes<T, SAFE_MATH> &outputs,
1254 graph::map_nodes<T, SAFE_MATH> &setters,
1255 graph::shared_random_state<T, SAFE_MATH> state,
1256 jit::register_map &registers,
1257 const jit::register_usage &usage,
1258 const jit::argument_set &thread_shared,
1259 jit::register_map &thread_mem,
1260 const size_t iterations=1) {
1261 std::unordered_set<void *> out_registers;
1262 for (auto &[out, in] : setters) {
1263 if (!out->is_match(in)) {
1264 auto a = out->compile(source_buffer, registers,
1265 thread_mem, usage);
1266 source_buffer << " ";
1267 if (thread_shared.contains(in.get())) {
1268 source_buffer << jit::to_string('t', in.get())
1269 << "[t_index] = ";
1270 } else {
1271 source_buffer << jit::to_string('v', in.get())
1272 << "[";
1273 if (state.get()) {
1274 source_buffer << "offset[0] + ";
1275 }
1276 source_buffer << "index] = ";
1277 }
1278 if constexpr (SAFE_MATH) {
1279 if constexpr (jit::complex_scalar<T>) {
1280 jit::add_type<T> (source_buffer);
1281 source_buffer << " (";
1282 if (!graph::random_cast(a).get()) {
1283 source_buffer << "isnan(real("
1284 << registers[a.get()]
1285 << ")) ? 0.0 : real("
1286 << registers[a.get()]
1287 << "), isnan(imag("
1288 << registers[a.get()]
1289 << ")) ? 0.0 : imag("
1290 << registers[a.get()]
1291 << ")";
1292 } else {
1293 source_buffer << registers[a.get()];
1294 }
1295 source_buffer << ");" << std::endl;
1296 } else {
1297 if (!graph::random_cast(a).get()) {
1298 source_buffer << "isnan(" << registers[a.get()]
1299 << ") ? 0.0 : ";
1300 }
1301 source_buffer << registers[a.get()]
1302 << ";" << std::endl;
1303 }
1304 } else {
1305 source_buffer << registers[a.get()] << ";" << std::endl;
1306 }
1307 out_registers.insert(out.get());
1308 }
1309 }
1310
1311 for (auto &out : outputs) {
1312 if (!graph::variable_cast(out).get() &&
1313 !graph::atomic_accumulate_1D_cast(out).get() &&
1314 !out_registers.contains(out.get())) {
1315 auto a = out->compile(source_buffer, registers,
1316 thread_mem, usage);
1317 source_buffer << " "
1318 << jit::to_string('o', out.get())
1319 << "[";
1320 if (state.get()) {
1321 source_buffer << "offset[0] + ";
1322 }
1323 source_buffer << "index] = ";
1324 if constexpr (SAFE_MATH) {
1325 if constexpr (jit::complex_scalar<T>) {
1326 jit::add_type<T> (source_buffer);
1327 source_buffer << " (";
1328 if (!graph::random_cast(a).get()) {
1329 source_buffer << "isnan(real("
1330 << registers[a.get()]
1331 << ")) ? 0.0 : real("
1332 << registers[a.get()]
1333 << "), isnan(imag("
1334 << registers[a.get()]
1335 << ")) ? 0.0 : imag("
1336 << registers[a.get()]
1337 << ")";
1338 } else {
1339 source_buffer << registers[a.get()];
1340 }
1341 source_buffer << ");" << std::endl;
1342 } else {
1343 if (!graph::random_cast(a).get()) {
1344 source_buffer << "isnan(" << registers[a.get()]
1345 << ") ? 0.0 : ";
1346 }
1347 source_buffer << registers[a.get()]
1348 << ";" << std::endl;
1349 }
1350 } else {
1351 source_buffer << registers[a.get()] << ";" << std::endl;
1352 }
1353 out_registers.insert(out.get());
1354 }
1355 }
1356
1357 if (iterations > 1) {
1358 source_buffer << " }" << std::endl;
1359 }
1360 for (auto &[out, in] : setters) {
1361 if (thread_shared.contains(in.get())) {
1362 source_buffer << " "
1363 << jit::to_string('v', in.get())
1364 << "[";
1365 if (state.get()) {
1366 source_buffer << "offset[0] + ";
1367 }
1368 source_buffer << "index] = "
1369 << jit::to_string('t', in.get())
1370 << "[t_index];" << std::endl;
1371 }
1372 }
1373 source_buffer << " }" << std::endl << "}" << std::endl;
1374 }
1375
1376//------------------------------------------------------------------------------
1381//------------------------------------------------------------------------------
1382 void create_reduction(std::ostringstream &source_buffer,
1383 const size_t size) {
1384 source_buffer << std::endl;
1385 source_buffer << "extern \"C\" __global__ void max_reduction(" << std::endl;
1386 source_buffer << " const ";
1387 jit::add_type<T> (source_buffer);
1388 source_buffer << " * __restrict__ input," << std::endl;
1389 source_buffer << " ";
1390 jit::add_type<T> (source_buffer);
1391 source_buffer << " * __restrict__ result) {" << std::endl;
1392 source_buffer << " const unsigned int i = threadIdx.x;" << std::endl;
1393 source_buffer << " const unsigned int j = threadIdx.x/32;" << std::endl;
1394 source_buffer << " const unsigned int k = threadIdx.x%32;" << std::endl;
1395 source_buffer << " if (i < " << size << ") {" << std::endl;
1396 source_buffer << " " << jit::type_to_string<T> () << " sub_max = ";
1397 if constexpr (jit::complex_scalar<T>) {
1398 source_buffer << "abs(input[i]);" << std::endl;
1399 } else {
1400 source_buffer << "input[i];" << std::endl;
1401 }
1402 source_buffer << " for (size_t index = i + 1024; index < " << size <<"; index += 1024) {" << std::endl;
1403 if constexpr (jit::complex_scalar<T>) {
1404 source_buffer << " sub_max = max(sub_max, abs(input[index]));" << std::endl;
1405 } else {
1406 source_buffer << " sub_max = max(sub_max, input[index]);" << std::endl;
1407 }
1408 source_buffer << " }" << std::endl;
1409 source_buffer << " __shared__ " << jit::type_to_string<T> () << " thread_max[32];" << std::endl;
1410 source_buffer << " for (int index = 16; index > 0; index /= 2) {" << std::endl;
1411 source_buffer << " sub_max = max(sub_max, __shfl_down_sync(__activemask(), sub_max, index));" << std::endl;
1412 source_buffer << " }" << std::endl;
1413 source_buffer << " thread_max[j] = sub_max;" << std::endl;
1414 source_buffer << " __syncthreads();" << std::endl;
1415 source_buffer << " if (j == 0) {" << std::endl;
1416 source_buffer << " for (int index = 16; index > 0; index /= 2) {" << std::endl;
1417 source_buffer << " thread_max[k] = max(thread_max[k], __shfl_down_sync(__activemask(), thread_max[k], index));" << std::endl;
1418 source_buffer << " }" << std::endl;
1419 source_buffer << " *result = thread_max[0];" << std::endl;
1420 source_buffer << " }" << std::endl;
1421 source_buffer << " }" << std::endl;
1422 source_buffer << "}" << std::endl << std::endl;
1423 }
1424
1425//------------------------------------------------------------------------------
1429//------------------------------------------------------------------------------
1430 T *get_buffer(graph::shared_leaf<T, SAFE_MATH> &node) {
1431 return reinterpret_cast<T *> (kernel_arguments[node.get()]);
1432 }
1433 };
1434}
1435
1436#endif /* cuda_context_h */
gpu::cuda_context
Class representing a cuda gpu context.
Definition cuda_context.hpp:74
gpu::cuda_context::get_buffer
T * get_buffer(graph::shared_leaf< T, SAFE_MATH > &node)
Get the buffer for a node.
Definition cuda_context.hpp:1430
gpu::cuda_context::~cuda_context
~cuda_context()
Cuda context destructor.
Definition cuda_context.hpp:156
gpu::cuda_context::remaining_const_memory
int remaining_const_memory
Remaining constant memory in bytes.
Definition cuda_context.hpp:117
gpu::cuda_context::create_zero_call
std::function< void(void)> create_zero_call(graph::input_nodes< T, SAFE_MATH > &inputs)
Create kernel call that will be memset a buffer to zero.
Definition cuda_context.hpp:653
gpu::cuda_context::copy_to_host
void copy_to_host(graph::shared_leaf< T, SAFE_MATH > node, T *destination)
Copy buffer contents to host.
Definition cuda_context.hpp:866
gpu::cuda_context::create_copy_call
std::function< void(void)> create_copy_call(graph::copy_nodes< T, SAFE_MATH > &setters)
Create kernel call that will to copy one buffer to another.
Definition cuda_context.hpp:709
gpu::cuda_context::create_reduction
void create_reduction(std::ostringstream &source_buffer, const size_t size)
Create reduction.
Definition cuda_context.hpp:1382
gpu::cuda_context::run_function
std::function< void(void)> run_function(std::function< void(void)> callback)
Run a callback function in the queue.
Definition cuda_context.hpp:789
gpu::cuda_context::create_max_call
std::function< T(void)> create_max_call(graph::shared_leaf< T, SAFE_MATH > &argument, std::function< void(void)> run)
Create a max compute kernel calling function.
Definition cuda_context.hpp:609
gpu::cuda_context::create_kernel_postfix
void create_kernel_postfix(std::ostringstream &source_buffer, graph::output_nodes< T, SAFE_MATH > &outputs, graph::map_nodes< T, SAFE_MATH > &setters, graph::shared_random_state< T, SAFE_MATH > state, jit::register_map &registers, const jit::register_usage &usage, const jit::argument_set &thread_shared, jit::register_map &thread_mem, const size_t iterations=1)
Create kernel postfix.
Definition cuda_context.hpp:1252
gpu::cuda_context::create_kernel_call
std::function< void(void)> create_kernel_call(const std::string kernel_name, graph::input_nodes< T, SAFE_MATH > inputs, graph::output_nodes< T, SAFE_MATH > outputs, graph::input_nodes< T, SAFE_MATH > atomics, graph::shared_random_state< T, SAFE_MATH > state, const size_t num_rays, const jit::texture1d_list &tex1d_list, const jit::texture2d_list &tex2d_list)
Create a kernel calling function.
Definition cuda_context.hpp:333
gpu::cuda_context::cuda_context
cuda_context(const size_t index)
Cuda context constructor.
Definition cuda_context.hpp:142
gpu::cuda_context::max_concurrency
static size_t max_concurrency()
Get the maximum number of concurrent instances.
Definition cuda_context.hpp:124
gpu::cuda_context::copy_to_device
void copy_to_device(graph::shared_leaf< T, SAFE_MATH > node, T *source)
Copy buffer contents to the device.
Definition cuda_context.hpp:853
gpu::cuda_context::device_type
static std::string device_type()
Device description.
Definition cuda_context.hpp:133
gpu::cuda_context::create_kernel_prefix
void create_kernel_prefix(std::ostringstream &source_buffer, const std::string name, graph::input_nodes< T, SAFE_MATH > &inputs, graph::output_nodes< T, SAFE_MATH > &outputs, graph::input_nodes< T, SAFE_MATH > atomics, graph::shared_random_state< T, SAFE_MATH > state, const size_t size, const std::vector< bool > &is_constant, jit::register_map &registers, const jit::register_usage &usage, jit::texture1d_list &textures1d, jit::texture2d_list &textures2d, jit::argument_set &thread_shared, jit::register_map &thread_mem, const size_t iterations=1)
Create kernel prefix.
Definition cuda_context.hpp:946
gpu::cuda_context::print_results
void print_results(const size_t index, const graph::output_nodes< T, SAFE_MATH > &nodes)
Print out the results.
Definition cuda_context.hpp:820
gpu::cuda_context::random_state_size
static constexpr size_t random_state_size
Size of random state needed.
Definition cuda_context.hpp:114
gpu::cuda_context::random_state_scale
static constexpr size_t random_state_scale
Random state size multiplyer.
Definition cuda_context.hpp:112
gpu::cuda_context::check_value
T check_value(const size_t index, const graph::shared_leaf< T, SAFE_MATH > &node)
Check the value.
Definition cuda_context.hpp:841
gpu::cuda_context::wait
void wait()
Hold the current thread until the stream has completed.
Definition cuda_context.hpp:778
gpu::cuda_context::create_header
void create_header(std::ostringstream &source_buffer)
Create the source header.
Definition cuda_context.hpp:878
gpu::cuda_context::compile
void compile(const std::string kernel_source, std::vector< std::string > names, const bool add_reduction=false)
Compile the kernels.
Definition cuda_context.hpp:197
timing::measure_diagnostic
A timing object.
Definition timing.hpp:18
jit::complex_scalar
Complex scalar concept.
Definition register.hpp:25
jit::double_base
Double base concept.
Definition register.hpp:43
jit::float_base
float base concept.
Definition register.hpp:38
MAX_REG
#define MAX_REG
Maximum number of registers to use.
Definition cuda_context.hpp:22
assert
subroutine assert(test, message)
Assert check.
Definition f_binding_test.f90:38
gpu
Name space for GPU backends.
Definition cpu_context.hpp:52
graph::atomic_accumulate_1D_cast
shared_atomic_accumulate_1D< T, SAFE_MATH > atomic_accumulate_1D_cast(shared_leaf< T, SAFE_MATH > x)
Cast to a atomic accumulate 1D node.
Definition piecewise.hpp:2426
graph::shared_random_state
std::shared_ptr< random_state_node< T, SAFE_MATH > > shared_random_state
Convenience type alias for shared sqrt nodes.
Definition random.hpp:269
graph::input_nodes
std::vector< shared_variable< T, SAFE_MATH > > input_nodes
Convenience type alias for a vector of inputs.
Definition node.hpp:2149
graph::variable_cast
shared_variable< T, SAFE_MATH > variable_cast(shared_leaf< T, SAFE_MATH > x)
Cast to a variable node.
Definition node.hpp:2169
graph::copy_nodes
std::vector< std::pair< shared_variable< T, SAFE_MATH >, shared_variable< T, SAFE_MATH > > > copy_nodes
Convenience type alias for copying buffers.
Definition node.hpp:2157
graph::shared_leaf
std::shared_ptr< leaf_node< T, SAFE_MATH > > shared_leaf
Convenience type alias for shared leaf nodes.
Definition node.hpp:691
graph::map_nodes
std::vector< std::pair< shared_leaf< T, SAFE_MATH >, shared_variable< T, SAFE_MATH > > > map_nodes
Convenience type alias for mapping end codes back to inputs.
Definition node.hpp:2153
graph::random_cast
shared_random< T, SAFE_MATH > random_cast(shared_leaf< T, SAFE_MATH > x)
Cast to a random node.
Definition random.hpp:539
graph::output_nodes
std::vector< shared_leaf< T, SAFE_MATH > > output_nodes
Convenience type alias for a vector of output nodes.
Definition node.hpp:706
jit::register_usage
std::unordered_map< void *, size_t > register_usage
Type alias for counting register usage.
Definition register.hpp:261
jit::register_map
std::unordered_map< void *, std::string > register_map
Type alias for mapping node pointers to register names.
Definition register.hpp:259
jit::texture2d_list
std::unordered_map< void *, std::array< size_t, 2 > > texture2d_list
Type alias for indexing 2D textures.
Definition register.hpp:267
jit::texture1d_list
std::unordered_map< void *, size_t > texture1d_list
Type alias for indexing 1D textures.
Definition register.hpp:265
jit::to_string
std::string to_string(const char prefix, const NODE *pointer)
Convert a graph::leaf_node pointer to a string.
Definition register.hpp:247
jit::argument_set
std::unordered_set< void * > argument_set
Type for tacking thread shared memory.
Definition register.hpp:271
output
Name space for output files.
Definition output.hpp:16
piecewise.hpp
Piecewise nodes.
random.hpp
Random constants and distributions.
timing.hpp
Routines to time the ray execution.