FazBrowse GitHub Viewer | Trending |
URL:
| Home
Tools: [Download Repo ZIP]   [Original HTTPS Page]

updated sycl algorithm unittests · yhcpp/taskflow@f729e5f · GitHub

  1. yhcpp
  2. taskflow

Repository navigation

Commit f729e5f

Browse files
twhuang-utah
committed
updated sycl algorithm unittests
1 parent 418e19b commit f729e5f

8 files changed

Lines changed: 321 additions & 57 deletions

File tree

‎doxygen/releases/release-3.1.0.dox‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -47,13 +47,15 @@ To use %Taskflow v3.1.0, you need a compiler that supports C++17:
4747
@subsection release-3-1-0_cudaflow cudaFlow
4848

4949
@li optimized tf::cudaRoundRobinCapturing through an event-pruning heuristic
50+
@li optimized the default block size used in %cudaFlow algorithms
5051
@li added tf::cudaFlow::clear() to clean up a %cudaFlow
5152
@li added tf::cudaFlow::num_tasks() to query the task count in a %cudaFlow
5253
@li added tf::cudaTask::num_dependents() to query the dependent count in a %cudaTask
5354
@li added tf::cudaFlowCapturer::clear() to clean up a %cudaFlow capturer
5455
@li added tf::cudaFlowCapturer::num_tasks() to query the task count in a %cudaFlow capturer
5556
@li added %cudaFlow examples
5657
+ parallel reduction (examples/cuda/cuda_reduce.cu)
58+
+ parallel transform (examples/cuda/cuda_transform.cu)
5759

5860
@subsection release-3-1-0_syclflow syclFlow
5961

‎examples/cuda/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@ list(APPEND TF_CUDA_EXAMPLES
99
cuda_knn
1010
cuda_capturer
1111
cuda_reduce
12+
cuda_transform
1213
)
1314

1415
foreach(cuda_example IN LISTS TF_CUDA_EXAMPLES)

‎examples/cuda/cuda_transform.cu‎

Lines changed: 50 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,50 @@
1+
// This program demonstrates how to performs a parallel transform
2+
// using cudaFlow.
3+
4+
#include <taskflow/cudaflow.hpp>
5+
6+
int main(int argc, char* argv[]) {
7+
8+
if(argc != 2) {
9+
std::cerr << "usage: ./cuda_transform num_items\n";
10+
std::exit(EXIT_FAILURE);
11+
}
12+
13+
size_t N = std::atoi(argv[1]);
14+
15+
auto data = tf::cuda_malloc_shared<int>(N);
16+
auto src1 = tf::cuda_malloc_shared<int>(N);
17+
auto src2 = tf::cuda_malloc_shared<int>(N);
18+
auto src3 = tf::cuda_malloc_shared<int>(N);
19+
20+
// initialize the data
21+
for(size_t i=0; i<N; i++) {
22+
data[i] = 0;
23+
src1[i] = 1;
24+
src2[i] = 2;
25+
src3[i] = 3;
26+
}
27+
28+
// perform parallel transform
29+
tf::cudaFlow cudaflow;
30+
31+
// data[i] = src1[i] + src2[i] + src3[i]
32+
cudaflow.transform(
33+
data, data+N,
34+
[] __device__ (int a, int b, int c) { return a+b+c; },
35+
src1, src2, src3
36+
);
37+
38+
cudaflow.offload();
39+
40+
// inspect the result
41+
for(size_t i=0; i<N; i++) {
42+
if(data[i] != (src1[i] + src2[i] + src3[i])) {
43+
throw std::runtime_error("incorrect result");
44+
}
45+
}
46+
47+
std::cout << "correct result\n";
48+
49+
return 0;
50+
}

‎taskflow/cuda/cuda_capturer.hpp‎

Lines changed: 29 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -13,23 +13,6 @@
1313

1414
namespace tf {
1515

16-
/**
17-
@brief queries the maximum threads allowed per block
18-
*/
19-
constexpr size_t cuda_default_max_threads_per_block() {
20-
return 512;
21-
}
22-
23-
/**
24-
@brief queries the default number of threads per block to run N elements
25-
*/
26-
constexpr size_t cuda_default_block_size(size_t N) {
27-
if(N <= 32) return 32;
28-
else {
29-
return std::min(cuda_default_max_threads_per_block(), next_pow2(N));
30-
}
31-
}
32-
3316
// ----------------------------------------------------------------------------
3417
// class definition: cudaFlowCapturerBase
3518
// ----------------------------------------------------------------------------
@@ -511,6 +494,8 @@ class cudaFlowCapturer {
511494
void offload();
512495

513496
private:
497+
498+
const size_t _MAX_BLOCK_SIZE;
514499

515500
handle_t _handle;
516501

@@ -527,18 +512,26 @@ class cudaFlowCapturer {
527512
cudaGraph_t _capture();
528513

529514
void _destroy_executable();
515+
516+
size_t _default_block_size(size_t) const;
530517
};
531518

532519
// constructs a cudaFlow capturer from a taskflow
533520
inline cudaFlowCapturer::cudaFlowCapturer(cudaGraph& g) :
534-
_handle{std::in_place_type_t<Internal>{}},
535-
_graph {g} {
521+
_MAX_BLOCK_SIZE {
522+
cuda_get_device_max_threads_per_block(cuda_get_device())
523+
},
524+
_handle {std::in_place_type_t<Internal>{}},
525+
_graph {g} {
536526
}
537527

538528
// constructs a standalone cudaFlow capturer
539529
inline cudaFlowCapturer::cudaFlowCapturer() :
540-
_handle{std::in_place_type_t<External>{}},
541-
_graph {std::get<External>(_handle).graph} {
530+
_MAX_BLOCK_SIZE {
531+
cuda_get_device_max_threads_per_block(cuda_get_device())
532+
},
533+
_handle {std::in_place_type_t<External>{}},
534+
_graph {std::get<External>(_handle).graph} {
542535
}
543536

544537
inline cudaFlowCapturer::~cudaFlowCapturer() {
@@ -547,6 +540,11 @@ inline cudaFlowCapturer::~cudaFlowCapturer() {
547540
}
548541
}
549542

543+
// Function: _default_block_size
544+
inline size_t cudaFlowCapturer::_default_block_size(size_t N) const {
545+
return N <= 32u ? 32u : std::min(_MAX_BLOCK_SIZE, next_pow2(N));
546+
}
547+
550548
// Function: empty
551549
inline bool cudaFlowCapturer::empty() const {
552550
return _graph.empty();
@@ -642,10 +640,10 @@ cudaTask cudaFlowCapturer::single_task(C&& callable) {
642640
// Function: for_each
643641
template <typename I, typename C>
644642
cudaTask cudaFlowCapturer::for_each(I first, I last, C&& c) {
645-
return on([first, last, c=std::forward<C>(c)](cudaStream_t stream) mutable {
643+
return on([&, first, last, c=std::forward<C>(c)](cudaStream_t stream) mutable {
646644
// TODO: special case for N == 0?
647645
size_t N = std::distance(first, last);
648-
size_t B = cuda_default_block_size(N);
646+
size_t B = _default_block_size(N);
649647
cuda_for_each<I, C><<<(N+B-1)/B, B, 0, stream>>>(first, N, c);
650648
});
651649
}
@@ -658,22 +656,22 @@ cudaTask cudaFlowCapturer::for_each_index(I beg, I end, I inc, C&& c) {
658656
TF_THROW("invalid range [", beg, ", ", end, ") with inc size ", inc);
659657
}
660658

661-
return on([beg, end, inc, c=std::forward<C>(c)] (cudaStream_t stream) mutable {
659+
return on([&, beg, end, inc, c=std::forward<C>(c)] (cudaStream_t stream) mutable {
662660
// TODO: special case when N is 0?
663661
size_t N = distance(beg, end, inc);
664-
size_t B = cuda_default_block_size(N);
662+
size_t B = _default_block_size(N);
665663
cuda_for_each_index<I, C><<<(N+B-1)/B, B, 0, stream>>>(beg, inc, N, c);
666664
});
667665
}
668666

669667
// Function: transform
670668
template <typename I, typename C, typename... S>
671669
cudaTask cudaFlowCapturer::transform(I first, I last, C&& c, S... srcs) {
672-
return on([first, last, c=std::forward<C>(c), srcs...]
670+
return on([&, first, last, c=std::forward<C>(c), srcs...]
673671
(cudaStream_t stream) mutable {
674672
// TODO: special case when N is 0?
675673
size_t N = std::distance(first, last);
676-
size_t B = cuda_default_block_size(N);
674+
size_t B = _default_block_size(N);
677675
cuda_transform<I, C, S...><<<(N+B-1)/B, B, 0, stream>>>(first, N, c, srcs...);
678676
});
679677
}
@@ -682,13 +680,13 @@ cudaTask cudaFlowCapturer::transform(I first, I last, C&& c, S... srcs) {
682680
template <typename I, typename T, typename C>
683681
cudaTask cudaFlowCapturer::reduce(I first, I last, T* result, C&& c) {
684682

685-
return on([first, last, result, c=std::forward<C>(c)]
683+
return on([&, first, last, result, c=std::forward<C>(c)]
686684
(cudaStream_t stream) mutable {
687685
//using value_t = std::decay_t<decltype(*std::declval<I>())>;
688686

689687
// TODO: special case N == 0?
690688
size_t N = std::distance(first, last);
691-
size_t B = cuda_default_block_size(N);
689+
size_t B = _default_block_size(N);
692690

693691
cuda_reduce<I, T, C, false><<<1, B, B*sizeof(T), stream>>>(
694692
first, N, result, c
@@ -702,13 +700,13 @@ cudaTask cudaFlowCapturer::uninitialized_reduce(
702700
I first, I last, T* result, C&& c
703701
) {
704702

705-
return on([first, last, result, c=std::forward<C>(c)]
703+
return on([&, first, last, result, c=std::forward<C>(c)]
706704
(cudaStream_t stream) mutable {
707705
//using value_t = std::decay_t<decltype(*std::declval<I>())>;
708706

709707
// TODO: special case N == 0?
710708
size_t N = std::distance(first, last);
711-
size_t B = cuda_default_block_size(N);
709+
size_t B = _default_block_size(N);
712710

713711
cuda_reduce<I, T, C, true><<<1, B, B*sizeof(T), stream>>>(
714712
first, N, result, c

‎taskflow/cuda/cuda_flow.hpp‎

Lines changed: 20 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -567,17 +567,24 @@ class cudaFlow {
567567

568568
private:
569569

570+
const size_t _MAX_BLOCK_SIZE;
571+
570572
handle_t _handle;
571573

572574
cudaGraph& _graph;
573575

574576
cudaGraphExec_t _executable {nullptr};
575577

576578
cudaFlow(cudaGraph&);
579+
580+
size_t _default_block_size(size_t N) const;
577581
};
578582

579583
// Construct a standalone cudaFlow
580584
inline cudaFlow::cudaFlow() :
585+
_MAX_BLOCK_SIZE {
586+
cuda_get_device_max_threads_per_block(cuda_get_device())
587+
},
581588
_handle {std::in_place_type_t<External>{}},
582589
_graph {std::get<External>(_handle).graph} {
583590

@@ -589,6 +596,9 @@ inline cudaFlow::cudaFlow() :
589596

590597
// Construct the cudaFlow from executor (internal graph)
591598
inline cudaFlow::cudaFlow(cudaGraph& g) :
599+
_MAX_BLOCK_SIZE {
600+
cuda_get_device_max_threads_per_block(cuda_get_device())
601+
},
592602
_handle {std::in_place_type_t<Internal>{}},
593603
_graph {g} {
594604

@@ -609,6 +619,11 @@ inline cudaFlow::~cudaFlow() {
609619
_graph._native_handle = nullptr;
610620
}
611621

622+
// Function: _default_block_size
623+
inline size_t cudaFlow::_default_block_size(size_t N) const {
624+
return N <= 32u ? 32u : std::min(_MAX_BLOCK_SIZE, next_pow2(N));
625+
}
626+
612627
// Procedure: clear
613628
inline void cudaFlow::clear() {
614629

@@ -1017,7 +1032,7 @@ template <typename I, typename C>
10171032
cudaTask cudaFlow::for_each(I first, I last, C&& c) {
10181033

10191034
size_t N = std::distance(first, last);
1020-
size_t B = cuda_default_block_size(N);
1035+
size_t B = _default_block_size(N);
10211036

10221037
// TODO: special case when N is 0?
10231038

@@ -1037,7 +1052,7 @@ cudaTask cudaFlow::for_each_index(I beg, I end, I inc, C&& c) {
10371052
// TODO: special case when N is 0?
10381053

10391054
size_t N = distance(beg, end, inc);
1040-
size_t B = cuda_default_block_size(N);
1055+
size_t B = _default_block_size(N);
10411056

10421057
return kernel(
10431058
(N+B-1) / B, B, 0, cuda_for_each_index<I, C>, beg, inc, N, std::forward<C>(c)
@@ -1051,7 +1066,7 @@ cudaTask cudaFlow::transform(I first, I last, C&& c, S... srcs) {
10511066
// TODO: special case when N is 0?
10521067

10531068
size_t N = std::distance(first, last);
1054-
size_t B = cuda_default_block_size(N);
1069+
size_t B = _default_block_size(N);
10551070

10561071
return kernel(
10571072
(N+B-1) / B, B, 0, cuda_transform<I, C, S...>,
@@ -1067,7 +1082,7 @@ cudaTask cudaFlow::reduce(I first, I last, T* result, C&& op) {
10671082

10681083
// TODO: special case N == 0?
10691084
size_t N = std::distance(first, last);
1070-
size_t B = cuda_default_block_size(N);
1085+
size_t B = _default_block_size(N);
10711086

10721087
return kernel(
10731088
1, B, B*sizeof(T), cuda_reduce<I, T, C, false>,
@@ -1083,7 +1098,7 @@ cudaTask cudaFlow::uninitialized_reduce(I first, I last, T* result, C&& op) {
10831098

10841099
// TODO: special case N == 0?
10851100
size_t N = std::distance(first, last);
1086-
size_t B = cuda_default_block_size(N);
1101+
size_t B = _default_block_size(N);
10871102

10881103
return kernel(
10891104
1, B, B*sizeof(T), cuda_reduce<I, T, C, true>,

‎taskflow/sycl/sycl_flow.hpp‎

Lines changed: 1 addition & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -378,12 +378,7 @@ inline syclFlow::syclFlow(Executor& e, syclGraph& g, sycl::queue& queue) :
378378

379379
// Function: _default_group_size
380380
inline size_t syclFlow::_default_group_size(size_t N) const {
381-
if(N <= 32) {
382-
return 32;
383-
}
384-
else {
385-
return std::min(_MAX_WORK_GROUP_SIZE, next_pow2(N));
386-
}
381+
return N <= 32u ? 32u : std::min(_MAX_WORK_GROUP_SIZE, next_pow2(N));
387382
}
388383

389384
// Function: empty

‎unittests/cuda/cuda_algorithm.cu‎

Lines changed: 0 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -465,15 +465,6 @@ TEST_CASE("capture_for_each_index.double" * doctest::timeout(300)) {
465465
// transform
466466
// ----------------------------------------------------------------------------
467467

468-
struct TransformFunc {
469-
__device__ int operator () (int& v1, float& v2, double& v3) {
470-
v1 = 1;
471-
v2 = 3.0f;
472-
v3 = 5.0;
473-
return 17;
474-
}
475-
};
476-
477468
template <typename F>
478469
void transform() {
479470

0 commit comments

Comments
 (0)
Back | FazBrowse Home | New Git URL