[SYCLomatic-test] test load and store headers - #680
Conversation
Sorry, something went wrong.
yihwang-nv
left a comment
There was a problem hiding this comment.
Same as 1784
Sorry, something went wrong.
|
Completed testing with pass results with PR : oneapi-src/SYCLomatic#1819 |
Sorry, something went wrong.
| @@ -0,0 +1,199 @@ | |||
| // ====------ onedpl_test_group_load.cpp------------ *- C++ -* ----===// | |||
There was a problem hiding this comment.
| // ====------ onedpl_test_group_load.cpp------------ *- C++ -* ----===// | |
| // ====------ util_group_load_store_test.cpp------------ *- C++ -* ----===// | |
Sorry, something went wrong.
|
|
||
| template <dpct::group::load_algorithm T> | ||
| bool helper_validation_function(const int *ptr, const char *func_name) { | ||
| if constexpr (T == dpct::group::load_algorithm::BLOCK_LOAD_DIRECT) { |
There was a problem hiding this comment.
It seems like the only difference between verification for "direct" and "striped" load algorithms is the name. This makes sense because we should "unmap" any "map" via the store operation as long as they match.
However, it doesn't make sense to just have a branch and repeat the code. You can branch for the name, and run the same verification, or just embed the info you want in func_name.
Sorry, something went wrong.
There was a problem hiding this comment.
Yes that makes sense, thanks
Sorry, something went wrong.
| return true; | ||
| } | ||
|
|
||
| bool subgroup_helper_validation_function(const int *ptr, const uint32_t *sg_sz, |
There was a problem hiding this comment.
looks like this has the same verification as well (as it should), but it can be consolidated into a single routine.
Sorry, something went wrong.
|
Can we make some of the kernels launch multiple work-groups to test for issues mentioned in the review? |
Sorry, something went wrong.
| sycl::local_accessor<uint8_t, 1> tacc(sycl::range<1>(temp_storage_size), h); | ||
| sycl::accessor data_accessor_read_write(buffer, h, sycl::read_write); | ||
| h.parallel_for( | ||
| sycl::nd_range<3>(sycl::range<3>(2, 2, 64), sycl::range<3>(1, 1, 64)), |
There was a problem hiding this comment.
All that seems to be different between the single and multi work group cases seems to be the range. Could we just add an nd_range parameter to a single test function and use that in the kernel submission to unify the single work-group / multi work-group functions?
Sorry, something went wrong.
There was a problem hiding this comment.
yes will add the change
Sorry, something went wrong.
|
|
||
| sycl::host_accessor data_accessor(buffer, sycl::read_write); | ||
| const int *ptr = data_accessor.get_multi_ptr<sycl::access::decorated::yes>(); | ||
| return helper_validation_function(ptr, "test_group_load_store"); |
There was a problem hiding this comment.
Can we add if it is the single work-group or multi work-group to the test string for more descriptive debug info?
Sorry, something went wrong.
There was a problem hiding this comment.
Added, thanks
Sorry, something went wrong.
This is linked to composite load and store tests (PRs oneapi-src/SYCLomatic#1819 , oneapi-src/SYCLomatic#1784) from group_utils.hpp
Another standalone "store-only" test is at : #725 (WIP)
cc @yihanwg @zhimingwang36 @danhoeflinger @mmichel11