Skip to content

Commit 0a66244

Browse files
LU-JOHNLU-JOHN
authored
[SYCLomatic] Refactor more thrust migrations (#767)
Signed-off-by: Lu, John <[email protected]>
1 parent 289da99 commit 0a66244

File tree

1 file changed

+11
-233
lines changed

1 file changed

+11
-233
lines changed

‎clang/lib/DPCT/APINamesThrust.inc‎

Lines changed: 11 additions & 233 deletions
Original file line numberDiff line numberDiff line change
@@ -63,243 +63,21 @@ thrustFactory("thrust::gather_if",
6363
{6,PolicyState::NoPolicy ,5,"dpct::gather_if", HelperFeatureEnum::DplExtrasAlgorithm_gather_if}}),
6464

6565
// thrust::inclusive_scan
66-
CONDITIONAL_FACTORY_ENTRY(
67-
CheckArgCount(5),
68-
CONDITIONAL_FACTORY_ENTRY(
69-
makeCheckAnd(CheckIsPtr(1), makeCheckNot(checkIsUSM())),
70-
//Handling case: thrust::inclusive_scan(policy, ptr, ptr, ptr, op);
71-
IFELSE_FACTORY_ENTRY(
72-
"thrust::inclusive_scan",
73-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::Memory_is_device_ptr,
74-
CALL_FACTORY_ENTRY("thrust::inclusive_scan", CALL(MapNames::getDpctNamespace() + "is_device_ptr", ARG(1)))),
75-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasMemory_device_pointer_forward_decl,
76-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
77-
CALL("oneapi::dpl::inclusive_scan",
78-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
79-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(1)), ARG(1)),
80-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(2)), ARG(2)),
81-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(3)), ARG(3)),
82-
ARG(4)))),
83-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
84-
CALL("oneapi::dpl::inclusive_scan",
85-
ARG("oneapi::dpl::execution::seq"),
86-
ARG(1), ARG(2), ARG(3), ARG(4)))),
87-
//Handling case: thrust::inclusive_scan(policy, device.begin(), device.end(), result, op);
88-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
89-
CALL("oneapi::dpl::inclusive_scan",
90-
makeMappedThrustPolicyEnum(0),
91-
ARG(1), ARG(2), ARG(3), ARG(4)))
92-
),
93-
CONDITIONAL_FACTORY_ENTRY(
94-
CheckArgCount(3),
95-
CONDITIONAL_FACTORY_ENTRY(
96-
makeCheckAnd(CheckIsPtr(1), makeCheckNot(checkIsUSM())),
97-
//Handling case: thrust::inclusive_scan(ptr, ptr, ptr);
98-
IFELSE_FACTORY_ENTRY(
99-
"thrust::inclusive_scan",
100-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::Memory_is_device_ptr,
101-
CALL_FACTORY_ENTRY("thrust::inclusive_scan", CALL(MapNames::getDpctNamespace() + "is_device_ptr", ARG(1)))),
102-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasMemory_device_pointer_forward_decl,
103-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
104-
CALL("oneapi::dpl::inclusive_scan",
105-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
106-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(0)), ARG(0)),
107-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(1)), ARG(1)),
108-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(2)), ARG(2))))),
109-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
110-
CALL("oneapi::dpl::inclusive_scan",
111-
ARG("oneapi::dpl::execution::seq"),
112-
ARG(0), ARG(1), ARG(2)))),
113-
//Handling case: thrust::inclusive_scan(host.begin(), host.end(), result);
114-
CONDITIONAL_FACTORY_ENTRY(
115-
CheckThrustArgType(1, "thrust::device_ptr"),
116-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
117-
CALL("oneapi::dpl::inclusive_scan",
118-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
119-
ARG(0), ARG(1), ARG(2))),
120-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
121-
CALL("oneapi::dpl::inclusive_scan",
122-
ARG("oneapi::dpl::execution::seq"),
123-
ARG(0), ARG(1), ARG(2))))
124-
),
125-
CONDITIONAL_FACTORY_ENTRY(
126-
makeCheckAnd(CheckIsPtr(1), makeCheckNot(checkIsUSM())),
127-
CONDITIONAL_FACTORY_ENTRY(
128-
IsPolicyArgType(0),
129-
//Handling case: thrust::inclusive_scan(policy, ptr, ptr, ptr);
130-
IFELSE_FACTORY_ENTRY(
131-
"thrust::inclusive_scan",
132-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::Memory_is_device_ptr,
133-
CALL_FACTORY_ENTRY("thrust::inclusive_scan", CALL(MapNames::getDpctNamespace() + "is_device_ptr", ARG(1)))),
134-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasMemory_device_pointer_forward_decl,
135-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
136-
CALL("oneapi::dpl::inclusive_scan",
137-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
138-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(1)), ARG(1)),
139-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(2)), ARG(2)),
140-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(3)), ARG(3))))),
141-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
142-
CALL("oneapi::dpl::inclusive_scan",
143-
ARG("oneapi::dpl::execution::seq"),
144-
ARG(1), ARG(2), ARG(3)))),
145-
//Handling case: thrust::inclusive_scan(ptr, ptr, ptr, op);
146-
IFELSE_FACTORY_ENTRY(
147-
"thrust::inclusive_scan",
148-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::Memory_is_device_ptr,
149-
CALL_FACTORY_ENTRY("thrust::inclusive_scan", CALL(MapNames::getDpctNamespace() + "is_device_ptr", ARG(1)))),
150-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasMemory_device_pointer_forward_decl,
151-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
152-
CALL("oneapi::dpl::inclusive_scan",
153-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
154-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(0)), ARG(0)),
155-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(1)), ARG(1)),
156-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(2)), ARG(2)),
157-
ARG(3)))),
158-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
159-
CALL("oneapi::dpl::inclusive_scan",
160-
ARG("oneapi::dpl::execution::seq"),
161-
ARG(0), ARG(1), ARG(2), ARG(3))))),
162-
CONDITIONAL_FACTORY_ENTRY(
163-
IsPolicyArgType(0),
164-
//Handling case: thrust::inclusive_scan(policy, device.begin(), device.end(), result);
165-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
166-
CALL("oneapi::dpl::inclusive_scan",
167-
makeMappedThrustPolicyEnum(0),
168-
ARG(1), ARG(2), ARG(3))),
169-
//Handling case: thrust::inclusive_scan(host.begin(), host.end(), result, op);
170-
CONDITIONAL_FACTORY_ENTRY(
171-
CheckThrustArgType(1, "thrust::device_ptr"),
172-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
173-
CALL("oneapi::dpl::inclusive_scan",
174-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
175-
ARG(0), ARG(1), ARG(2), ARG(3))),
176-
CALL_FACTORY_ENTRY("thrust::inclusive_scan",
177-
CALL("oneapi::dpl::inclusive_scan",
178-
ARG("oneapi::dpl::execution::seq"),
179-
ARG(0), ARG(1), ARG(2), ARG(3)))))
180-
)
181-
)
182-
)
66+
thrustFactory("thrust::inclusive_scan",
67+
{{5,PolicyState::HasPolicy,3,"oneapi::dpl::inclusive_scan", HelperFeatureEnum::no_feature_helper },
68+
{4,PolicyState::HasPolicy,3,"oneapi::dpl::inclusive_scan", HelperFeatureEnum::no_feature_helper },
69+
{4,PolicyState::NoPolicy ,3,"oneapi::dpl::inclusive_scan", HelperFeatureEnum::no_feature_helper },
70+
{3,PolicyState::NoPolicy ,3,"oneapi::dpl::inclusive_scan", HelperFeatureEnum::no_feature_helper }}),
18371

18472
// thrust::gather
185-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasAlgorithm_gather,
186-
CONDITIONAL_FACTORY_ENTRY(
187-
CheckArgCount(5),
188-
CONDITIONAL_FACTORY_ENTRY(
189-
makeCheckAnd(CheckIsPtr(1), makeCheckNot(checkIsUSM())),
190-
// Handling case: thrust::gather(policy, ptr, ptr, ptr, ptr);
191-
IFELSE_FACTORY_ENTRY(
192-
"thrust::gather",
193-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::Memory_is_device_ptr,
194-
CALL_FACTORY_ENTRY("thrust::gather", CALL(MapNames::getDpctNamespace() + "is_device_ptr", ARG(1)))),
195-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasMemory_device_pointer_forward_decl,
196-
CALL_FACTORY_ENTRY("thrust::gather",
197-
CALL(MapNames::getDpctNamespace() + "gather",
198-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
199-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(1)), ARG(1)),
200-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(2)), ARG(2)),
201-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(3)), ARG(3)),
202-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(4)), ARG(4))))),
203-
CALL_FACTORY_ENTRY("thrust::gather",
204-
CALL(MapNames::getDpctNamespace() + "gather",
205-
ARG("oneapi::dpl::execution::seq"),
206-
ARG(1), ARG(2), ARG(3), ARG(4)))),
207-
// Handling case: thrust::gather(thrust::device, d_map.begin(), d_map.end(), d_values.begin(),d_output.begin());
208-
CALL_FACTORY_ENTRY("thrust::gather",
209-
CALL(MapNames::getDpctNamespace() + "gather",
210-
makeMappedThrustPolicyEnum(0),
211-
ARG(1), ARG(2), ARG(3), ARG(4)))
212-
),
213-
CONDITIONAL_FACTORY_ENTRY(
214-
makeCheckAnd(CheckIsPtr(1), makeCheckNot(checkIsUSM())),
215-
// Handling case: thrust::gather(ptr, ptr, ptr, ptr);
216-
IFELSE_FACTORY_ENTRY(
217-
"thrust::gather",
218-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::Memory_is_device_ptr,
219-
CALL_FACTORY_ENTRY("thrust::gather", CALL(MapNames::getDpctNamespace() + "is_device_ptr", ARG(1)))),
220-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasMemory_device_pointer_forward_decl,
221-
CALL_FACTORY_ENTRY("thrust::gather",
222-
CALL(MapNames::getDpctNamespace() + "gather",
223-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
224-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(0)), ARG(0)),
225-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(1)), ARG(1)),
226-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(2)), ARG(2)),
227-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(3)), ARG(3))))),
228-
CALL_FACTORY_ENTRY("thrust::gather",
229-
CALL(MapNames::getDpctNamespace() + "gather",
230-
ARG("oneapi::dpl::execution::seq"),
231-
ARG(0), ARG(1), ARG(2), ARG(3)))),
232-
// Handling case: thrust::gather(h_map.begin(), h_map.end(), h_values.begin(),h_output.begin());
233-
CONDITIONAL_FACTORY_ENTRY(
234-
CheckThrustArgType(1, "thrust::device_ptr"),
235-
CALL_FACTORY_ENTRY("thrust::gather", CALL(MapNames::getDpctNamespace() + "gather",
236-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
237-
ARG(0), ARG(1), ARG(2), ARG(3))),
238-
CALL_FACTORY_ENTRY("thrust::gather", CALL(MapNames::getDpctNamespace() + "gather", ARG("oneapi::dpl::execution::seq"),
239-
ARG(0), ARG(1), ARG(2), ARG(3))))
240-
)
241-
)
242-
)
73+
thrustFactory("thrust::gather",
74+
{{5,PolicyState::HasPolicy,4,MapNames::getDpctNamespace()+"gather", HelperFeatureEnum::DplExtrasAlgorithm_gather },
75+
{4,PolicyState::NoPolicy ,4,MapNames::getDpctNamespace()+"gather", HelperFeatureEnum::DplExtrasAlgorithm_gather }}),
24376

24477
// thrust::scatter
245-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasAlgorithm_scatter,
246-
CONDITIONAL_FACTORY_ENTRY(
247-
CheckArgCount(5),
248-
CONDITIONAL_FACTORY_ENTRY(
249-
makeCheckAnd(CheckIsPtr(1), makeCheckNot(checkIsUSM())),
250-
// Handling case: thrust::scatter(policy, ptr, ptr, ptr, ptr);
251-
IFELSE_FACTORY_ENTRY(
252-
"thrust::scatter",
253-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::Memory_is_device_ptr,
254-
CALL_FACTORY_ENTRY("thrust::scatter", CALL(MapNames::getDpctNamespace() + "is_device_ptr", ARG(1)))),
255-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasMemory_device_pointer_forward_decl,
256-
CALL_FACTORY_ENTRY("thrust::scatter",
257-
CALL(MapNames::getDpctNamespace() + "scatter",
258-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
259-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(1)), ARG(1)),
260-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(2)), ARG(2)),
261-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(3)), ARG(3)),
262-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(4)), ARG(4))))),
263-
CALL_FACTORY_ENTRY("thrust::scatter",
264-
CALL(MapNames::getDpctNamespace() + "scatter",
265-
ARG("oneapi::dpl::execution::seq"),
266-
ARG(1), ARG(2), ARG(3), ARG(4)))),
267-
// Handling case: thrust::scatter(policy, d_values.begin(), d_values.end(), d_map.begin(), d_output.begin());
268-
CALL_FACTORY_ENTRY("thrust::scatter",
269-
CALL(MapNames::getDpctNamespace() + "scatter",
270-
makeMappedThrustPolicyEnum(0),
271-
ARG(1), ARG(2), ARG(3), ARG(4)))
272-
),
273-
CONDITIONAL_FACTORY_ENTRY(
274-
makeCheckAnd(CheckIsPtr(1), makeCheckNot(checkIsUSM())),
275-
// Handling case: thrust::scatter(ptr, ptr, ptr, ptr);
276-
IFELSE_FACTORY_ENTRY(
277-
"thrust::scatter",
278-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::Memory_is_device_ptr,
279-
CALL_FACTORY_ENTRY("thrust::scatter", CALL(MapNames::getDpctNamespace() + "is_device_ptr", ARG(1)))),
280-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasMemory_device_pointer_forward_decl,
281-
CALL_FACTORY_ENTRY("thrust::scatter",
282-
CALL(MapNames::getDpctNamespace() + "scatter",
283-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
284-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(0)), ARG(0)),
285-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(1)), ARG(1)),
286-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(2)), ARG(2)),
287-
CALL(TEMPLATED_CALLEE_WITH_ARGS(MapNames::getDpctNamespace() + "device_pointer", getDerefedType(3)), ARG(3))))),
288-
CALL_FACTORY_ENTRY("thrust::scatter",
289-
CALL(MapNames::getDpctNamespace() + "scatter",
290-
ARG("oneapi::dpl::execution::seq"),
291-
ARG(0), ARG(1), ARG(2), ARG(3)))),
292-
// Handling case: thrust::scatter(d_values.begin(), d_values.end(), d_map.begin(), d_output.begin());
293-
CONDITIONAL_FACTORY_ENTRY(
294-
CheckThrustArgType(1, "thrust::device_ptr"),
295-
CALL_FACTORY_ENTRY("thrust::scatter", CALL(MapNames::getDpctNamespace() + "scatter",
296-
CALL("oneapi::dpl::execution::make_device_policy", QUEUESTR),
297-
ARG(0), ARG(1), ARG(2), ARG(3))),
298-
CALL_FACTORY_ENTRY("thrust::scatter", CALL(MapNames::getDpctNamespace() + "scatter", ARG("oneapi::dpl::execution::seq"),
299-
ARG(0), ARG(1), ARG(2), ARG(3))))
300-
)
301-
)
302-
)
78+
thrustFactory("thrust::scatter",
79+
{{5,PolicyState::HasPolicy,4,MapNames::getDpctNamespace()+"scatter", HelperFeatureEnum::DplExtrasAlgorithm_scatter },
80+
{4,PolicyState::NoPolicy ,4,MapNames::getDpctNamespace()+"scatter", HelperFeatureEnum::DplExtrasAlgorithm_scatter }}),
30381

30482
// thrust::unique_by_key_copy
30583
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::DplExtrasAlgorithm_unique_copy,

0 commit comments

Comments
 (0)