Skip to content

Commit 20a1056

Browse files
gonidelisgonidelis
committed
Add run_to_run and not_guaranteed api tests
1 parent e5832ef commit 20a1056

File tree

1 file changed

+170
-0
lines changed

1 file changed

+170
-0
lines changed

‎cub/test/catch2_test_device_segmented_reduce_env_api.cu‎

Lines changed: 170 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -267,6 +267,176 @@ C2H_TEST("cub::DeviceSegmentedReduce::ArgMax env-based API", "[segmented_reduce]
267267
REQUIRE(h_out[2].value == 9);
268268
}
269269

270+
C2H_TEST("cub::DeviceSegmentedReduce::Min accepts run_to_run determinism requirements", "[segmented_reduce][env]")
271+
{
272+
int num_segments = 3;
273+
thrust::device_vector<int> d_offsets = {0, 3, 3, 7};
274+
auto d_offsets_it = thrust::raw_pointer_cast(d_offsets.data());
275+
thrust::device_vector<int> d_in{8, 6, 7, 5, 3, 0, 9};
276+
thrust::device_vector<int> d_out(3);
277+
278+
auto env = cuda::execution::require(cuda::execution::determinism::run_to_run);
279+
280+
auto error =
281+
cub::DeviceSegmentedReduce::Min(d_in.begin(), d_out.begin(), num_segments, d_offsets_it, d_offsets_it + 1, env);
282+
thrust::device_vector<int> expected{6, std::numeric_limits<int>::max(), 0};
283+
284+
REQUIRE(d_out == expected);
285+
REQUIRE(error == cudaSuccess);
286+
}
287+
288+
C2H_TEST("cub::DeviceSegmentedReduce::Min accepts not_guaranteed determinism requirements", "[segmented_reduce][env]")
289+
{
290+
int num_segments = 3;
291+
thrust::device_vector<int> d_offsets = {0, 3, 3, 7};
292+
auto d_offsets_it = thrust::raw_pointer_cast(d_offsets.data());
293+
thrust::device_vector<int> d_in{8, 6, 7, 5, 3, 0, 9};
294+
thrust::device_vector<int> d_out(3);
295+
296+
auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed);
297+
298+
auto error =
299+
cub::DeviceSegmentedReduce::Min(d_in.begin(), d_out.begin(), num_segments, d_offsets_it, d_offsets_it + 1, env);
300+
thrust::device_vector<int> expected{6, std::numeric_limits<int>::max(), 0};
301+
302+
REQUIRE(d_out == expected);
303+
REQUIRE(error == cudaSuccess);
304+
}
305+
306+
C2H_TEST("cub::DeviceSegmentedReduce::Max accepts run_to_run determinism requirements", "[segmented_reduce][env]")
307+
{
308+
int num_segments = 3;
309+
thrust::device_vector<int> d_offsets = {0, 3, 3, 7};
310+
auto d_offsets_it = thrust::raw_pointer_cast(d_offsets.data());
311+
thrust::device_vector<int> d_in{8, 6, 7, 5, 3, 0, 9};
312+
thrust::device_vector<int> d_out(3);
313+
314+
auto env = cuda::execution::require(cuda::execution::determinism::run_to_run);
315+
316+
auto error =
317+
cub::DeviceSegmentedReduce::Max(d_in.begin(), d_out.begin(), num_segments, d_offsets_it, d_offsets_it + 1, env);
318+
thrust::device_vector<int> expected{8, std::numeric_limits<int>::lowest(), 9};
319+
320+
REQUIRE(d_out == expected);
321+
REQUIRE(error == cudaSuccess);
322+
}
323+
324+
C2H_TEST("cub::DeviceSegmentedReduce::Max accepts not_guaranteed determinism requirements", "[segmented_reduce][env]")
325+
{
326+
int num_segments = 3;
327+
thrust::device_vector<int> d_offsets = {0, 3, 3, 7};
328+
auto d_offsets_it = thrust::raw_pointer_cast(d_offsets.data());
329+
thrust::device_vector<int> d_in{8, 6, 7, 5, 3, 0, 9};
330+
thrust::device_vector<int> d_out(3);
331+
332+
auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed);
333+
334+
auto error =
335+
cub::DeviceSegmentedReduce::Max(d_in.begin(), d_out.begin(), num_segments, d_offsets_it, d_offsets_it + 1, env);
336+
thrust::device_vector<int> expected{8, std::numeric_limits<int>::lowest(), 9};
337+
338+
REQUIRE(d_out == expected);
339+
REQUIRE(error == cudaSuccess);
340+
}
341+
342+
C2H_TEST("cub::DeviceSegmentedReduce::ArgMin accepts run_to_run determinism requirements", "[segmented_reduce][env]")
343+
{
344+
int num_segments = 3;
345+
thrust::device_vector<int> d_offsets = {0, 3, 3, 7};
346+
auto d_offsets_it = thrust::raw_pointer_cast(d_offsets.data());
347+
thrust::device_vector<int> d_in{8, 6, 7, 5, 3, 0, 9};
348+
thrust::device_vector<cub::KeyValuePair<int, int>> d_out(3);
349+
350+
auto env = cuda::execution::require(cuda::execution::determinism::run_to_run);
351+
352+
auto error =
353+
cub::DeviceSegmentedReduce::ArgMin(d_in.begin(), d_out.begin(), num_segments, d_offsets_it, d_offsets_it + 1, env);
354+
355+
REQUIRE(error == cudaSuccess);
356+
357+
thrust::host_vector<cub::KeyValuePair<int, int>> h_out(d_out);
358+
REQUIRE(h_out[0].key == 1);
359+
REQUIRE(h_out[0].value == 6);
360+
REQUIRE(h_out[1].key == 1);
361+
REQUIRE(h_out[1].value == std::numeric_limits<int>::max());
362+
REQUIRE(h_out[2].key == 2);
363+
REQUIRE(h_out[2].value == 0);
364+
}
365+
366+
C2H_TEST("cub::DeviceSegmentedReduce::ArgMin accepts not_guaranteed determinism requirements",
367+
"[segmented_reduce][env]")
368+
{
369+
int num_segments = 3;
370+
thrust::device_vector<int> d_offsets = {0, 3, 3, 7};
371+
auto d_offsets_it = thrust::raw_pointer_cast(d_offsets.data());
372+
thrust::device_vector<int> d_in{8, 6, 7, 5, 3, 0, 9};
373+
thrust::device_vector<cub::KeyValuePair<int, int>> d_out(3);
374+
375+
auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed);
376+
377+
auto error =
378+
cub::DeviceSegmentedReduce::ArgMin(d_in.begin(), d_out.begin(), num_segments, d_offsets_it, d_offsets_it + 1, env);
379+
380+
REQUIRE(error == cudaSuccess);
381+
382+
thrust::host_vector<cub::KeyValuePair<int, int>> h_out(d_out);
383+
REQUIRE(h_out[0].key == 1);
384+
REQUIRE(h_out[0].value == 6);
385+
REQUIRE(h_out[1].key == 1);
386+
REQUIRE(h_out[1].value == std::numeric_limits<int>::max());
387+
REQUIRE(h_out[2].key == 2);
388+
REQUIRE(h_out[2].value == 0);
389+
}
390+
391+
C2H_TEST("cub::DeviceSegmentedReduce::ArgMax accepts run_to_run determinism requirements", "[segmented_reduce][env]")
392+
{
393+
int num_segments = 3;
394+
thrust::device_vector<int> d_offsets = {0, 3, 3, 7};
395+
auto d_offsets_it = thrust::raw_pointer_cast(d_offsets.data());
396+
thrust::device_vector<int> d_in{8, 6, 7, 5, 3, 0, 9};
397+
thrust::device_vector<cub::KeyValuePair<int, int>> d_out(3);
398+
399+
auto env = cuda::execution::require(cuda::execution::determinism::run_to_run);
400+
401+
auto error =
402+
cub::DeviceSegmentedReduce::ArgMax(d_in.begin(), d_out.begin(), num_segments, d_offsets_it, d_offsets_it + 1, env);
403+
404+
REQUIRE(error == cudaSuccess);
405+
406+
thrust::host_vector<cub::KeyValuePair<int, int>> h_out(d_out);
407+
REQUIRE(h_out[0].key == 0);
408+
REQUIRE(h_out[0].value == 8);
409+
REQUIRE(h_out[1].key == 1);
410+
REQUIRE(h_out[1].value == std::numeric_limits<int>::lowest());
411+
REQUIRE(h_out[2].key == 3);
412+
REQUIRE(h_out[2].value == 9);
413+
}
414+
415+
C2H_TEST("cub::DeviceSegmentedReduce::ArgMax accepts not_guaranteed determinism requirements",
416+
"[segmented_reduce][env]")
417+
{
418+
int num_segments = 3;
419+
thrust::device_vector<int> d_offsets = {0, 3, 3, 7};
420+
auto d_offsets_it = thrust::raw_pointer_cast(d_offsets.data());
421+
thrust::device_vector<int> d_in{8, 6, 7, 5, 3, 0, 9};
422+
thrust::device_vector<cub::KeyValuePair<int, int>> d_out(3);
423+
424+
auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed);
425+
426+
auto error =
427+
cub::DeviceSegmentedReduce::ArgMax(d_in.begin(), d_out.begin(), num_segments, d_offsets_it, d_offsets_it + 1, env);
428+
429+
REQUIRE(error == cudaSuccess);
430+
431+
thrust::host_vector<cub::KeyValuePair<int, int>> h_out(d_out);
432+
REQUIRE(h_out[0].key == 0);
433+
REQUIRE(h_out[0].value == 8);
434+
REQUIRE(h_out[1].key == 1);
435+
REQUIRE(h_out[1].value == std::numeric_limits<int>::lowest());
436+
REQUIRE(h_out[2].key == 3);
437+
REQUIRE(h_out[2].value == 9);
438+
}
439+
270440
C2H_TEST("cub::DeviceSegmentedReduce::Reduce accepts run_to_run determinism requirements", "[segmented_reduce][env]")
271441
{
272442
// example-begin segmented-reduce-reduce-env-determinism

0 commit comments

Comments
 (0)