| 124 | 124 | data_type <- "typedef dust::no_data data_type;" |
|
| 125 | 125 | } else { |
|
| 126 | 126 | data_type <- c( |
|
| 127 | - | "struct ALIGN(16) data_type {", |
|
| 127 | + | "struct __align__(16) data_type {", |
|
| 128 | 128 | sprintf(" %s %s;", unname(dat$compare$data), names(dat$compare$data)), |
|
| 129 | 129 | "};") |
|
| 130 | 130 | } |
| 205 | 205 | "rng_state_type&" = dat$meta$dust$rng_state, |
|
| 206 | 206 | "real_type *" = dat$meta$result) |
|
| 207 | 207 | ||
| 208 | - | cpp_function("HOST void", "update", args, body) |
|
| 208 | + | cpp_function("void", "update", args, body) |
|
| 209 | 209 | } |
|
| 210 | 210 | ||
| 211 | 211 |
| 639 | 639 | dat$gpu <- generate_dust_gpu_storage(dat) |
|
| 640 | 640 | ||
| 641 | 641 | cpp_namespace("dust", |
|
| 642 | - | cpp_namespace("cuda", |
|
| 642 | + | cpp_namespace("gpu", |
|
| 643 | 643 | c(generate_dust_gpu_size(dat, rewrite), |
|
| 644 | 644 | generate_dust_gpu_copy(dat, rewrite), |
|
| 645 | 645 | generate_dust_gpu_update(dat), |
| 648 | 648 | ||
| 649 | 649 | ||
| 650 | 650 | generate_dust_gpu_update <- function(dat) { |
|
| 651 | - | name <- sprintf("update_device<%s>", dat$config$base) |
|
| 651 | + | name <- sprintf("update_gpu<%s>", dat$config$base) |
|
| 652 | 652 | ||
| 653 | 653 | args <- c( |
|
| 654 | 654 | "size_t" = dat$meta$time, |
|
| 655 | - | "const dust::cuda::interleaved<%s::real_type>" = dat$meta$state, |
|
| 656 | - | "dust::cuda::interleaved<int>" = dat$meta$dust$internal_int, |
|
| 657 | - | "dust::cuda::interleaved<%s::real_type>" = dat$meta$dust$internal_real, |
|
| 655 | + | "const dust::gpu::interleaved<%s::real_type>" = dat$meta$state, |
|
| 656 | + | "dust::gpu::interleaved<int>" = dat$meta$dust$internal_int, |
|
| 657 | + | "dust::gpu::interleaved<%s::real_type>" = dat$meta$dust$internal_real, |
|
| 658 | 658 | "const int *" = dat$meta$dust$shared_int, |
|
| 659 | 659 | "const %s::real_type *" = dat$meta$dust$shared_real, |
|
| 660 | 660 | "%s::rng_state_type&" = dat$meta$dust$rng_state, |
|
| 661 | - | "dust::cuda::interleaved<%s::real_type>" = dat$meta$result) |
|
| 661 | + | "dust::gpu::interleaved<%s::real_type>" = dat$meta$result) |
|
| 662 | 662 | names(args) <- sub("%s", dat$config$base, names(args), fixed = TRUE) |
|
| 663 | 663 | ||
| 664 | 664 | eqs <- generate_dust_equations(dat, NULL, dat$components$rhs$equations, |
| 668 | 668 | dust_flatten_eqs(eqs)) |
|
| 669 | 669 | ||
| 670 | 670 | c("template<>", |
|
| 671 | - | cpp_function("DEVICE void", name, args, body)) |
|
| 671 | + | cpp_function("__device__ void", name, args, body)) |
|
| 672 | 672 | } |
|
| 673 | 673 | ||
| 674 | 674 |
| 680 | 680 | code <- dat$compare$function_defn |
|
| 681 | 681 | ||
| 682 | 682 | base <- dat$config$base |
|
| 683 | - | return_type <- sprintf("DEVICE %s::real_type", base) |
|
| 684 | - | name <- sprintf("compare_device<%s>", base) |
|
| 683 | + | return_type <- sprintf("__device__ %s::real_type", base) |
|
| 684 | + | name <- sprintf("compare_gpu<%s>", base) |
|
| 685 | 685 | ||
| 686 | 686 | args <- c( |
|
| 687 | - | "const dust::cuda::interleaved<%s::real_type>" = "state", |
|
| 687 | + | "const dust::gpu::interleaved<%s::real_type>" = "state", |
|
| 688 | 688 | "const %s::data_type&" = "data", |
|
| 689 | - | "dust::cuda::interleaved<int>" = "internal_int", |
|
| 690 | - | "dust::cuda::interleaved<%s::real_type>" = "internal_real", |
|
| 689 | + | "dust::gpu::interleaved<int>" = "internal_int", |
|
| 690 | + | "dust::gpu::interleaved<%s::real_type>" = "internal_real", |
|
| 691 | 691 | "const int *" = "shared_int", |
|
| 692 | 692 | "const %s::real_type *" = "shared_real", |
|
| 693 | 693 | "%s::rng_state_type&" = "rng_state") |
| 705 | 705 | ||
| 706 | 706 | generate_dust_gpu_size <- function(dat, rewrite) { |
|
| 707 | 707 | dust_gpu_size <- function(x) { |
|
| 708 | - | name <- sprintf("dust::cuda::device_%s_size<%s>", |
|
| 708 | + | name <- sprintf("%s_size<%s>", |
|
| 709 | 709 | x$location, dat$config$base) |
|
| 710 | 710 | args <- set_names(dat$meta$dust$shared, |
|
| 711 | 711 | sprintf("dust::shared_ptr<%s>", dat$config$base)) |
| 719 | 719 | ||
| 720 | 720 | ||
| 721 | 721 | generate_dust_gpu_copy <- function(dat, rewrite) { |
|
| 722 | - | name <- sprintf("dust::cuda::device_shared_copy<%s>", dat$config$base) |
|
| 722 | + | name <- sprintf("dust::gpu::shared_copy<%s>", dat$config$base) |
|
| 723 | 723 | args <- c( |
|
| 724 | 724 | "dust::shared_ptr<%s>" = dat$meta$dust$shared, |
|
| 725 | 725 | "int *" = dat$meta$dust$shared_int, |
| 727 | 727 | names(args) <- sub("%s", dat$config$base, names(args), fixed = TRUE) |
|
| 728 | 728 | ||
| 729 | 729 | copy1 <- function(name, shared) { |
|
| 730 | - | sprintf("%s = dust::cuda::shared_copy(%s, %s);", shared, shared, |
|
| 730 | + | sprintf("%s = dust::gpu::shared_copy_data(%s, %s);", shared, shared, |
|
| 731 | 731 | vcapply(name, rewrite, USE.NAMES = FALSE)) |
|
| 732 | 732 | } |
|
| 733 | 733 |
| 827 | 827 | } else { |
|
| 828 | 828 | location <- dat$meta$state |
|
| 829 | 829 | } |
|
| 830 | - | type <- if (rank == 0) "real_type" else "dust::cuda::interleaved<real_type>" |
|
| 830 | + | type <- if (rank == 0) "real_type" else "dust::gpu::interleaved<real_type>" |
|
| 831 | 831 | c(x, list(type = type, rank = rank, location = location)) |
|
| 832 | 832 | } |
|
| 833 | 833 |
| 922 | 922 | location_dust <- sprintf("%s_%s", |
|
| 923 | 923 | location, if (type == "int") "int" else "real") |
|
| 924 | 924 | if (location == "internal") { |
|
| 925 | - | type_array <- sprintf("dust::cuda::interleaved<%s>", type_dust) |
|
| 925 | + | type_array <- sprintf("dust::gpu::interleaved<%s>", type_dust) |
|
| 926 | 926 | } else { |
|
| 927 | 927 | type_array <- sprintf("%s *", type_dust) |
|
| 928 | 928 | } |
| 115 | 115 | name = "odin_sum1", |
|
| 116 | 116 | declaration = c( |
|
| 117 | 117 | "template <typename real_type, typename container>", |
|
| 118 | - | paste("HOSTDEVICE real_type", |
|
| 118 | + | paste("__host__ __device__ real_type", |
|
| 119 | 119 | "odin_sum1(const container x, size_t from, size_t to);")), |
|
| 120 | 120 | definition = NULL) |
|
| 121 | 121 | } else { |
| 127 | 127 | ret <- lapply(odin:::generate_c_support_sum(rank), replace, tr) |
|
| 128 | 128 | for (v in c("declaration", "definition")) { |
|
| 129 | 129 | s <- ret[[v]] |
|
| 130 | - | s[[1L]] <- paste("HOSTDEVICE", s[[1L]]) |
|
| 130 | + | s[[1L]] <- paste("__host__ __device__", s[[1L]]) |
|
| 131 | 131 | ret[[v]] <- c(head, s) |
|
| 132 | 132 | } |
|
| 133 | 133 | } |
| Files | Coverage |
|---|---|
| R | 100.00% |
| Project Totals (7 files) | 100.00% |
richfitz
