[SYCLomatic-test] test load and store headers #680
Conversation
abhilash1910
changed the title
abhilash1910
changed the title
|
Completed testing with pass results with PR : oneapi-src/SYCLomatic#1819 |
| @@ -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++ -* ----===// | |
|
|
||
| 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.
There was a problem hiding this comment.
Yes that makes sense, thanks
| 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.
|
Can we make some of the kernels launch multiple work-groups to test for issues mentioned in the review? |
| 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?
There was a problem hiding this comment.
yes will add the change
|
|
||
| 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?
There was a problem hiding this comment.
Added, thanks
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