Skip to content

Commit a7d7bbe

Browse files
LU-JOHNLU-JOHN
authored
[SYCLomatic] Migrate thrust::gather_if (#752)
* Migrate thrust::gather_if Signed-off-by: Lu, John <[email protected]> * Record unsupported thrust overloads cleanly Signed-off-by: Lu, John <[email protected]> * Add test_api_level testcase Signed-off-by: Lu, John <[email protected]> --------- Signed-off-by: Lu, John <[email protected]>
1 parent a09591c commit a7d7bbe

File tree

7 files changed

+271
-9
lines changed

7 files changed

+271
-9
lines changed

‎clang/lib/DPCT/APINamesThrust.inc‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -54,6 +54,11 @@ thrustFactory("thrust::adjacent_difference",
5454
{4,PolicyState::NoPolicy ,3,"oneapi::dpl::adjacent_difference", HelperFeatureEnum::no_feature_helper },
5555
{3,PolicyState::NoPolicy ,3,"oneapi::dpl::adjacent_difference", HelperFeatureEnum::no_feature_helper }}),
5656

57+
// thrust::gather_if
58+
thrustFactory("thrust::gather_if",
59+
{{7,PolicyState::HasPolicy,5,"dpct::gather_if", HelperFeatureEnum::DplExtrasAlgorithm_gather_if},
60+
{6,PolicyState::NoPolicy ,5,"dpct::gather_if", HelperFeatureEnum::DplExtrasAlgorithm_gather_if}}),
61+
5762
// thrust::inclusive_scan
5863
CONDITIONAL_FACTORY_ENTRY(
5964
CheckArgCount(5),

‎clang/lib/DPCT/APINames_thrust.inc‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -73,7 +73,7 @@ ENTRY(thrust::for_each, thrust::for_each, true, NO_FLAG, P4, "Successful")
7373
ENTRY(thrust::for_each_n, thrust::for_each_n, true, NO_FLAG, P4, "comment")
7474
ENTRY(thrust::free, thrust::free, false, NO_FLAG, P4, "comment")
7575
ENTRY(thrust::gather, thrust::gather, true, NO_FLAG, P4, "Successful")
76-
ENTRY(thrust::gather_if, thrust::gather_if, false, NO_FLAG, P4, "comment")
76+
ENTRY(thrust::gather_if, thrust::gather_if, true, NO_FLAG, P4, "Successful")
7777
ENTRY(thrust::generate, thrust::generate, true, NO_FLAG, P4, "Successful")
7878
ENTRY(thrust::generate_n, thrust::generate_n, true, NO_FLAG, P4, "Successful")
7979
ENTRY(thrust::get, thrust::get, true, NO_FLAG, P4, "Successful")

‎clang/test/dpct/test_api_level/DplExtrasAlgorithm/api_test22.cu‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
// UNSUPPORTED: cuda-8.0
2+
// UNSUPPORTED: v8.0
3+
// RUN: dpct --format-range=none --usm-level=none --use-custom-helper=api -out-root %T/DplExtrasAlgorithm/api_test22_out %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only -std=c++17 -fsized-deallocation
4+
// RUN: grep "IsCalled" %T/DplExtrasAlgorithm/api_test22_out/MainSourceFiles.yaml | wc -l > %T/DplExtrasAlgorithm/api_test22_out/count.txt
5+
// RUN: FileCheck --input-file %T/DplExtrasAlgorithm/api_test22_out/count.txt --match-full-lines %s
6+
// RUN: rm -rf %T/DplExtrasAlgorithm/api_test22_out
7+
8+
// CHECK: 41
9+
// TEST_FEATURE: DplExtrasAlgorithm_gather_if
10+
11+
#include <thrust/gather.h>
12+
#include <thrust/device_vector.h>
13+
14+
int main() {
15+
int init[10] = {0, 1, 2, 3};
16+
thrust::device_vector<int> MD(init, init + 4);
17+
thrust::device_vector<int> SD(init, init + 4);
18+
thrust::device_vector<int> ID(init, init + 4);
19+
thrust::device_vector<int> RD(4);
20+
21+
thrust::gather_if(MD.begin(), MD.end(), SD.begin(), ID.begin(), RD.begin());
22+
return 0;
23+
}

‎clang/test/dpct/thrust_gather.cu‎

Lines changed: 124 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,124 @@
1+
// UNSUPPORTED: cuda-8.0
2+
// UNSUPPORTED: v8.0
3+
// RUN: dpct -out-root %T/thrust_gather %s --cuda-include-path="%cuda-path/include" --usm-level=none
4+
// RUN: FileCheck --input-file %T/thrust_gather/thrust_gather.dp.cpp --match-full-lines %s
5+
6+
#include <thrust/gather.h>
7+
#include <thrust/device_vector.h>
8+
#include <thrust/host_vector.h>
9+
#include <thrust/execution_policy.h>
10+
struct is_less_than_zero
11+
{
12+
__host__ __device__
13+
bool operator()(int x) const
14+
{
15+
return x < 0;
16+
}
17+
};
18+
19+
int main(void) {
20+
21+
thrust::device_vector<int> AD(4);
22+
thrust::device_vector<int> BD(4);
23+
thrust::device_vector<int> SD(4);
24+
thrust::device_vector<int> RD(4);
25+
thrust::host_vector<int> AH(4);
26+
thrust::host_vector<int> BH(4);
27+
thrust::host_vector<int> SH(4);
28+
thrust::host_vector<int> RH(4);
29+
30+
is_less_than_zero pred;
31+
32+
int *h_ptr;
33+
int *d_ptr;
34+
35+
h_ptr = (int*)std::malloc(20 * sizeof(int));
36+
cudaMalloc(&d_ptr, 20 * sizeof(int));
37+
38+
/*******************************************************************************************
39+
1. Test gather_if
40+
2. Test four VERSIONs (with/without exec argument with/without predicate)
41+
3. Test each VERSION with (device_vector/host_vector/malloc-ed memory/cudaMalloc-ed memory)
42+
*******************************************************************************************/
43+
44+
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
45+
// CHECK:dpct::gather_if(oneapi::dpl::execution::seq, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin(), pred);
46+
// CHECK-NEXT:dpct::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin(), pred);
47+
// CHECK-NEXT:if (dpct::is_device_ptr(h_ptr + 4)) {
48+
// CHECK-NEXT: dpct::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), dpct::device_pointer<int>(h_ptr), dpct::device_pointer<int>(h_ptr + 4), dpct::device_pointer<>(SH.begin()), dpct::device_pointer<>(BH.begin()), dpct::device_pointer<>(RH.begin()), pred);
49+
// CHECK-NEXT:} else {
50+
// CHECK-NEXT: dpct::gather_if(oneapi::dpl::execution::seq, h_ptr, h_ptr + 4, SH.begin(), BH.begin(), RH.begin(), pred);
51+
// CHECK-NEXT:};
52+
// VERSION first last stencil input result pred
53+
thrust::gather_if( AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin(), pred);
54+
thrust::gather_if( AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin(), pred);
55+
thrust::gather_if( h_ptr, h_ptr+4, SH.begin(), BH.begin(), RH.begin(), pred);
56+
// Overload not supported with thrust
57+
// thrust::gather_if( d_ptr, d_ptr+4, SD.begin(), BD.begin(), RD.begin(), pred);
58+
59+
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
60+
// CHECK:dpct::gather_if(oneapi::dpl::execution::seq, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin(), pred);
61+
// CHECK-NEXT:dpct::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin(), pred);
62+
// CHECK-NEXT:if (dpct::is_device_ptr(h_ptr)) {
63+
// CHECK-NEXT: dpct::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), dpct::device_pointer<int>(h_ptr), dpct::device_pointer<int>(h_ptr + 4), dpct::device_pointer<>(SH.begin()), dpct::device_pointer<>(BH.begin()), dpct::device_pointer<>(RH.begin()), pred);
64+
// CHECK-NEXT:} else {
65+
// CHECK-NEXT: dpct::gather_if(oneapi::dpl::execution::seq, h_ptr, h_ptr + 4, SH.begin(), BH.begin(), RH.begin(), pred);
66+
// CHECK-NEXT:};
67+
// CHECK-NEXT:if (dpct::is_device_ptr(d_ptr)) {
68+
// CHECK-NEXT: dpct::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), dpct::device_pointer<int>(d_ptr), dpct::device_pointer<int>(d_ptr + 4), dpct::device_pointer<>(SD.begin()), dpct::device_pointer<>(BD.begin()), dpct::device_pointer<>(RD.begin()), pred);
69+
// CHECK-NEXT:} else {
70+
// CHECK-NEXT: dpct::gather_if(oneapi::dpl::execution::seq, d_ptr, d_ptr + 4, SD.begin(), BD.begin(), RD.begin(), pred);
71+
// CHECK-NEXT:};
72+
// VERSION exec first last stencil input result pred
73+
thrust::gather_if(thrust::host, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin(), pred);
74+
thrust::gather_if(thrust::device, AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin(), pred);
75+
thrust::gather_if(thrust::host, h_ptr, h_ptr+4, SH.begin(), BH.begin(), RH.begin(), pred);
76+
thrust::gather_if(thrust::device, d_ptr, d_ptr+4, SD.begin(), BD.begin(), RD.begin(), pred);
77+
78+
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
79+
80+
// CHECK: /*
81+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
82+
// CHECK-NEXT: */
83+
// CHECK-NEXT: thrust::gather_if(AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin());
84+
// CHECK-NEXT: /*
85+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
86+
// CHECK-NEXT: */
87+
// CHECK-NEXT: thrust::gather_if(AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin());
88+
// CHECK-NEXT: /*
89+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
90+
// CHECK-NEXT: */
91+
// CHECK-NEXT: thrust::gather_if(h_ptr, h_ptr + 4, SH.begin(), BH.begin(), RH.begin());
92+
// VERSION first last stencil input result
93+
thrust::gather_if( AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin());
94+
thrust::gather_if( AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin());
95+
thrust::gather_if( h_ptr, h_ptr+4, SH.begin(), BH.begin(), RH.begin());
96+
// Overload not supported with thrust
97+
// thrust::gather_if( d_ptr, d_ptr+4, SD.begin(), BD.begin(), RD.begin());
98+
99+
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
100+
101+
// CHECK: /*
102+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
103+
// CHECK-NEXT: */
104+
// CHECK-NEXT: thrust::gather_if(oneapi::dpl::execution::seq, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin());
105+
// CHECK-NEXT: /*
106+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
107+
// CHECK-NEXT: */
108+
// CHECK-NEXT: thrust::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin());
109+
// CHECK-NEXT: /*
110+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
111+
// CHECK-NEXT: */
112+
// CHECK-NEXT: thrust::gather_if(oneapi::dpl::execution::seq, h_ptr, h_ptr + 4, SH.begin(), BH.begin(), RH.begin());
113+
// CHECK-NEXT: /*
114+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
115+
// CHECK-NEXT: */
116+
// CHECK-NEXT: thrust::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), d_ptr, d_ptr + 4, SD.begin(), BD.begin(), RD.begin());
117+
// VERSION exec first last stencil input result
118+
thrust::gather_if(thrust::host, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin());
119+
thrust::gather_if(thrust::device, AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin());
120+
thrust::gather_if(thrust::host, h_ptr, h_ptr+4, SH.begin(), BH.begin(), RH.begin());
121+
thrust::gather_if(thrust::device, d_ptr, d_ptr+4, SD.begin(), BD.begin(), RD.begin());
122+
123+
return 0;
124+
}

‎clang/test/dpct/thrust_gather_usm.cu‎

Lines changed: 114 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,114 @@
1+
// UNSUPPORTED: cuda-8.0
2+
// UNSUPPORTED: v8.0
3+
// RUN: dpct -out-root %T/thrust_gather_usm %s --cuda-include-path="%cuda-path/include" --usm-level=restricted
4+
// RUN: FileCheck --input-file %T/thrust_gather_usm/thrust_gather_usm.dp.cpp --match-full-lines %s
5+
6+
#include <thrust/gather.h>
7+
#include <thrust/device_vector.h>
8+
#include <thrust/host_vector.h>
9+
#include <thrust/execution_policy.h>
10+
struct is_less_than_zero
11+
{
12+
__host__ __device__
13+
bool operator()(int x) const
14+
{
15+
return x < 0;
16+
}
17+
};
18+
19+
int main(void) {
20+
21+
thrust::device_vector<int> AD(4);
22+
thrust::device_vector<int> BD(4);
23+
thrust::device_vector<int> SD(4);
24+
thrust::device_vector<int> RD(4);
25+
thrust::host_vector<int> AH(4);
26+
thrust::host_vector<int> BH(4);
27+
thrust::host_vector<int> SH(4);
28+
thrust::host_vector<int> RH(4);
29+
30+
is_less_than_zero pred;
31+
32+
int *h_ptr;
33+
int *d_ptr;
34+
35+
h_ptr = (int*)std::malloc(20 * sizeof(int));
36+
cudaMalloc(&d_ptr, 20 * sizeof(int));
37+
38+
/*******************************************************************************************
39+
1. Test gather_if
40+
2. Test four VERSIONs (with/without exec argument) AND (with/without predicate argument)
41+
3. Test each VERSION with (device_vector/host_vector/malloc-ed memory/cudaMalloc-ed memory)
42+
*******************************************************************************************/
43+
44+
/*********** gather_if ***********************************************************************************************************************************************/
45+
46+
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
47+
// CHECK:dpct::gather_if(oneapi::dpl::execution::seq, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin(), pred);
48+
// CHECK-NEXT:dpct::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin(), pred);
49+
// CHECK-NEXT:dpct::gather_if(oneapi::dpl::execution::seq, h_ptr, h_ptr + 4, SH.begin(), BH.begin(), RH.begin(), pred);
50+
// VERSION first last stencil input result pred
51+
thrust::gather_if( AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin(), pred);
52+
thrust::gather_if( AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin(), pred);
53+
thrust::gather_if( h_ptr, h_ptr+4, SH.begin(), BH.begin(), RH.begin(), pred);
54+
// Overload not supported with thrust
55+
// thrust::gather_if( d_ptr, d_ptr+4, SD.begin(), BD.begin(), RD.begin(), pred);
56+
57+
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
58+
// CHECK:dpct::gather_if(oneapi::dpl::execution::seq, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin(), pred);
59+
// CHECK-NEXT:dpct::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin(), pred);
60+
// CHECK-NEXT:dpct::gather_if(oneapi::dpl::execution::seq, h_ptr, h_ptr + 4, SH.begin(), BH.begin(), RH.begin(), pred);
61+
// CHECK-NEXT:dpct::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), d_ptr, d_ptr + 4, SD.begin(), BD.begin(), RD.begin(), pred);
62+
// VERSION exec first last stencil input result pred
63+
thrust::gather_if(thrust::host, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin(), pred);
64+
thrust::gather_if(thrust::device, AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin(), pred);
65+
thrust::gather_if(thrust::host, h_ptr, h_ptr+4, SH.begin(), BH.begin(), RH.begin(), pred);
66+
thrust::gather_if(thrust::device, d_ptr, d_ptr+4, SD.begin(), BD.begin(), RD.begin(), pred);
67+
68+
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
69+
70+
// CHECK: /*
71+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
72+
// CHECK-NEXT: */
73+
// CHECK-NEXT: thrust::gather_if(AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin());
74+
// CHECK-NEXT: /*
75+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
76+
// CHECK-NEXT: */
77+
// CHECK-NEXT: thrust::gather_if(AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin());
78+
// CHECK-NEXT: /*
79+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
80+
// CHECK-NEXT: */
81+
// CHECK-NEXT: thrust::gather_if(h_ptr, h_ptr + 4, SH.begin(), BH.begin(), RH.begin());
82+
// VERSION first last stencil input result
83+
thrust::gather_if( AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin());
84+
thrust::gather_if( AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin());
85+
thrust::gather_if( h_ptr, h_ptr+4, SH.begin(), BH.begin(), RH.begin());
86+
// Overload not supported with thrust
87+
// thrust::gather_if( d_ptr, d_ptr+4, SD.begin(), BD.begin(), RD.begin());
88+
89+
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
90+
91+
// CHECK: /*
92+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
93+
// CHECK-NEXT: */
94+
// CHECK-NEXT: thrust::gather_if(oneapi::dpl::execution::seq, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin());
95+
// CHECK-NEXT: /*
96+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
97+
// CHECK-NEXT: */
98+
// CHECK-NEXT: thrust::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin());
99+
// CHECK-NEXT: /*
100+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
101+
// CHECK-NEXT: */
102+
// CHECK-NEXT: thrust::gather_if(oneapi::dpl::execution::seq, h_ptr, h_ptr + 4, SH.begin(), BH.begin(), RH.begin());
103+
// CHECK-NEXT: /*
104+
// CHECK-NEXT: DPCT1107:{{[0-9]+}}: Migration for this overload of thrust::gather_if is not supported.
105+
// CHECK-NEXT: */
106+
// CHECK-NEXT: thrust::gather_if(oneapi::dpl::execution::make_device_policy(q_ct1), d_ptr, d_ptr + 4, SD.begin(), BD.begin(), RD.begin());
107+
// VERSION exec first last stencil input result
108+
thrust::gather_if(thrust::host, AH.begin(), AH.end(), SH.begin(), BH.begin(), RH.begin());
109+
thrust::gather_if(thrust::device, AD.begin(), AD.end(), SD.begin(), BD.begin(), RD.begin());
110+
thrust::gather_if(thrust::host, h_ptr, h_ptr+4, SH.begin(), BH.begin(), RH.begin());
111+
thrust::gather_if(thrust::device, d_ptr, d_ptr+4, SD.begin(), BD.begin(), RD.begin());
112+
113+
return 0;
114+
}

‎clang/test/dpct/thrust_replace.cu‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -151,10 +151,8 @@ int main(void) {
151151
thrust::replace_copy_if( AH.begin(), AH.end(), SH.begin(), BH.begin(), pred, 0);
152152
thrust::replace_copy_if( AD.begin(), AD.end(), SD.begin(), BD.begin(), pred, 0);
153153
thrust::replace_copy_if( h_ptr, h_ptr+4, SH.begin(), BH.begin(), pred, 0);
154-
#ifdef ADD_BUG
155-
// This fails with nvcc
156-
thrust::replace_copy_if( d_ptr, d_ptr+4, SD.begin(), BD.begin(), pred, 0);
157-
#endif
154+
// Overload not supported with thrust
155+
// thrust::replace_copy_if( d_ptr, d_ptr+4, SD.begin(), BD.begin(), pred, 0);
158156

159157
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
160158
// CHECK:oneapi::dpl::replace_copy_if(oneapi::dpl::execution::seq, AH.begin(), AH.end(), BH.begin(), pred, 0);

‎clang/test/dpct/thrust_replace_usm.cu‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -107,10 +107,8 @@ int main(void) {
107107
thrust::replace_copy_if( AH.begin(), AH.end(), SH.begin(), BH.begin(), pred, 0);
108108
thrust::replace_copy_if( AD.begin(), AD.end(), SD.begin(), BD.begin(), pred, 0);
109109
thrust::replace_copy_if( h_ptr, h_ptr+4, SH.begin(), BH.begin(), pred, 0);
110-
#ifdef ADD_BUG
111-
// This fails with nvcc
112-
thrust::replace_copy_if( d_ptr, d_ptr+4, SD.begin(), BD.begin(), pred, 0);
113-
#endif
110+
// Overload not supported with thrust
111+
// thrust::replace_copy_if( d_ptr, d_ptr+4, SD.begin(), BD.begin(), pred, 0);
114112

115113
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
116114
// CHECK:oneapi::dpl::replace_copy_if(oneapi::dpl::execution::seq, AH.begin(), AH.end(), BH.begin(), pred, 0);

0 commit comments

Comments
 (0)