EXSWHTEC-272 - Implement tests for warp shfl_up and shfl_down functions #193

nives-vukovic · 2023-03-03T12:57:27Z

Add tests for warp __shfl_up function
Add tests for warp __shfl_down function

Change-Id: I66f0c09e9c7405ec7430b1883e0e89542fdb87a0

Change-Id: I212b82b1b3a78a368b85ea64e338371a34b405f9

Change-Id: Ib455f72b5be77e1a81137d15c07ea41161b16a3e

Change-Id: Ief96e274f4143e80ceb3e40f04d38ae217777583

Change-Id: I9c03cde09b42c8e3726153c2a177359efc8d6d29

- Add tests for warp __shfl_up function - Add tests for warp __shfl_down function

scchan · 2023-04-18T20:50:03Z

catch/unit/warp/warp_shfl_down.cc

+  }
+
+  const auto grid = cg::this_grid();
+  T var = static_cast<T>(grid.thread_rank() % warpSize);


var is going to be a small integer. One concern here is that shfl is going to be shuffling zeros most of the time. For larger data types, the higher order bits will always be zeros. The test could be improved by using better quality data.

We should avoid generating the input data and the output data (in the validate function) on the fly because this would reduce the usefulness of this unit test. We need to have separation of concern in input/expected output data generation and in testing of the actual functionality.

@scchan Generating input data has been added (more detailed explanation is added in the comment below). Input data is not generated on the fly anymore.

scchan · 2023-04-18T20:51:21Z

catch/unit/warp/warp_shfl_down.cc

+
+  const auto grid = cg::this_grid();
+  T var = static_cast<T>(grid.thread_rank() % warpSize);
+  out[grid.thread_rank()] = __shfl_down(var, delta, width);


We need a version of this test that has a non-uniform delta within a single wrap.

@scchan Test has been changed to have non-uniform delta in each warp that has random value from 0 to width.

scchan · 2023-04-18T21:14:24Z

catch/include/cpu_grid.h

+  unsigned int thread_count_;
+};
+
+inline dim3 GenerateThreadDimensions() {


Needs comments on how the Generate*Dimensions() functions work.

@scchan Added basic comments, if more extensive comment is required it can be added. It is mostly self-explanatory, it uses Catch2 GENERATE_COPY to generate different dimensions for blocks of threads to cover a range of dimensions, some depending on warp size, including dimensions that are smaller than one warp size or not a multiple of warp size, and some arbitrary values that have been randomly chosen to make unit testing more robust.

scchan · 2023-04-18T23:23:18Z

catch/unit/warp/warp_common.hh

+  const auto block_rank = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x;
+  const auto idx = block_rank * warps_per_block + block.thread_rank() / warpSize;
+
+  return !(active_masks[idx] & (static_cast<uint64_t>(1) << warp.thread_rank()));


I suggest implementing a generic bitmap and use the grid thread rank as the index to retrieve the active bit.

@scchan We agree that is also a possible approach, but this current approach is more in line with the current implementation.

scchan · 2023-04-18T23:25:48Z

catch/unit/warp/warp_common.hh

+                                                warps_in_grid * sizeof(uint64_t));
+    active_masks_.resize(warps_in_grid);
+    std::generate(active_masks_.begin(), active_masks_.end(),
+                  [] { return GenerateRandomInteger(0ul, std::numeric_limits<uint64_t>().max()); });


We should avoid using only random values as input because that hurts reproducibility.

We should avoid using only random values as input because that hurts reproducibility.

Agree that not only random patterns need to be checked. Do you have particular patterns in mind though?

@scchan @b-sumner Test has been expanded to have a version that uses random inputs and random active_masks, and version that uses predefined active masks and inputs (5 different patterns for active masks have been chosen and thread id has been used as input data, changed from warp id to include larger input values).

scchan · 2023-04-18T23:27:45Z

catch/unit/warp/warp_common.hh

+  return dist(GetRandomGenerator());
+}
+
+inline uint64_t get_predicate_mask(unsigned int test_case, unsigned int warp_size) {


This is not being used anywhere so what is this for?

@scchan All warp related PRs have a common warp_base branch, this is the function that is used in ballot, all and any tests.

… changes

…n_tests

…tests

rakesroy · 2024-02-20T07:30:57Z

PR has been merged into develop branch via commit 26a5250.

…ns (#193) Change-Id: I3013d16f48ad5f607ee0f252b497fde24c7b9164

- #154 - #438 - #425 - #424 - #423 - #365 - #356 - #279 - #274 - #190 - #189 - #188 - #156 - #49 - #439 - #437 - #436 - #435 - #193 Change-Id: I2529d0baf0f8d47d6215863321720cde2b1a846c

gargrahul and others added 30 commits October 26, 2022 03:59

SWDEV-355313 - Move catch tests and samples

cea96af

Change-Id: I66f0c09e9c7405ec7430b1883e0e89542fdb87a0

SWDEV-355313 - Add README

909e7e4

Change-Id: I212b82b1b3a78a368b85ea64e338371a34b405f9

SWDEV-355313 - Update amd-staging branch

094b9af

Change-Id: Ib455f72b5be77e1a81137d15c07ea41161b16a3e

SWDEV-355313 - Update README

9daa6d0

Change-Id: Ief96e274f4143e80ceb3e40f04d38ae217777583

SWDEV-355313 - Update latest code

c49043e

Change-Id: I9c03cde09b42c8e3726153c2a177359efc8d6d29

Migrate basic Cooperative Groups tests and integrate to catch

edf514a

Refactor basic Cooperative Groups tests

5610a48

Rename tiled partition related files and fix minor bug

c455740

Add LaunchCooperativeKernal and LaunchCooperativeKernelMultiDevice tests

82fc666

Refactor hipCGThreadBlockTileType to use common function

ef4fa46

Merge remote-tracking branch 'origin/develop' into hipCoopGroups_wip

9c0f995

Fix updated file not added during merge

b28aa60

Add coalesced_group type tests

cc32117

Add coalesced_group shuffle_up and shuffle_down tests

a177e26

Add coalesced_group shuffle tests - test fails

cdeadcf

Merge remote-tracking branch 'upstream/develop' into cg_base_dino

7b84ac1

Implement common code for cooperative group tests

d414bce

Fixed compilation errror in cooperative_groups_common.hh

609fae5

Implement busy wait device function

8cfb58b

Add thread and block dimensions generators

5cf02ca

Move cpu_grid.h and supporting functions to catch/include

fc11bf9

Use warp_size from properties in grid/block dims generators

18f2450

Fix condition for warp size 32 on AMD

65a1e57

Fix cpu_grid.h for warp function tests

c01665f

Add missing include into cpu_grid.h

e41e642

Merge remote-tracking branch 'origin/develop' into warp_common

e0e35e9

Add common functions and definitions for warp functions

d4291ae

Remove unnecessary memset

1c154be

Cleanup leftover cooperative groups files

2aed190

EXSWHTEC-272 - Implement tests for warp shfl_up and shfl_down functions

2499e31

- Add tests for warp __shfl_up function - Add tests for warp __shfl_down function

nives-vukovic added 2 commits March 3, 2023 14:29

Add memory reset after allocation

583e30a

Merge branch 'warp_common' into warp_shfl_up_down_tests

6838eb2

nives-vukovic marked this pull request as ready for review March 3, 2023 16:04

chrispaquot requested review from yxsamliu and b-sumner March 9, 2023 03:17

EXSWHTEC-272 - Fix doxygen comments

64b983a

searlmc1 requested a review from scchan April 13, 2023 18:40

scchan requested changes Apr 18, 2023

View reviewed changes

scchan reviewed Apr 18, 2023

View reviewed changes

nives-vukovic added 5 commits May 3, 2023 12:58

Expand Warp Test to include random and predefined test version

792358c

Merge branch 'warp_common' into warp_shfl_up_down_tests

a180c2d

EXSWHTEC-272 - Modify warp shfl up and down tests according to common…

de7a5bc

… changes

Add comments for block and grid dimensions generate functions

fb1615d

Merge branch 'warp_common' into warp_shfl_up_down_tests

d19342d

nives-vukovic requested a review from scchan May 23, 2023 13:45

mangupta mentioned this pull request Jun 29, 2023

SWDEV-396533 - Add test for intrinsic shfl API #313

Open

rakesroy and others added 5 commits July 11, 2023 17:11

Merge branch 'develop' into warp_shfl_up_down_tests

d4313b1

Reduce common code for warp tests

22eb41a

Merge branch 'warp_common' into warp_shfl_up_down_tests

8a07cb5

EXSWHTEC-272 - Create separate warp shfl common code

9da7394

Merge remote-tracking branch 'upstream/develop' into warp_shfl_up_dow…

a0a1e66

…n_tests

mirza-halilcevic mentioned this pull request Sep 29, 2023

EXSWHTEC-334 - Extend tests for warp shlf_up and shfl_down functions to support half-precision types #419

Closed

nives-vukovic and others added 2 commits December 8, 2023 08:59

Merge remote-tracking branch 'origin/develop' into warp_shfl_up_down_…

e8178f7

…tests

Merge branch 'develop' into warp_shfl_up_down_tests

4d90638

rakesroy closed this Feb 20, 2024

rocm-ci pushed a commit that referenced this pull request Feb 26, 2024

EXSWHTEC-272 - Implement tests for warp shfl_up and shfl_down functio…

26a5250

…ns (#193) Change-Id: I3013d16f48ad5f607ee0f252b497fde24c7b9164

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

EXSWHTEC-272 - Implement tests for warp shfl_up and shfl_down functions #193

EXSWHTEC-272 - Implement tests for warp shfl_up and shfl_down functions #193

nives-vukovic commented Mar 3, 2023

scchan Apr 18, 2023

scchan Apr 18, 2023

nives-vukovic May 3, 2023 •

edited

Loading

scchan Apr 18, 2023

nives-vukovic May 3, 2023

scchan Apr 18, 2023

nives-vukovic May 3, 2023

scchan Apr 18, 2023

nives-vukovic May 3, 2023 •

edited

Loading

scchan Apr 18, 2023

b-sumner Apr 19, 2023

nives-vukovic May 3, 2023 •

edited

Loading

scchan Apr 18, 2023

nives-vukovic May 3, 2023

rakesroy commented Feb 20, 2024

EXSWHTEC-272 - Implement tests for warp shfl_up and shfl_down functions #193

EXSWHTEC-272 - Implement tests for warp shfl_up and shfl_down functions #193

Conversation

nives-vukovic commented Mar 3, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

nives-vukovic May 3, 2023 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

nives-vukovic May 3, 2023 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

nives-vukovic May 3, 2023 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

rakesroy commented Feb 20, 2024

nives-vukovic May 3, 2023 •

edited

Loading

nives-vukovic May 3, 2023 •

edited

Loading

nives-vukovic May 3, 2023 •

edited

Loading