Skip to content
Merged
Changes from 1 commit
Commits
Show all changes
38 commits
Select commit Hold shift + click to select a range
c08ae48
block load tests
abhilash1910 Feb 1, 2024
513b103
update test
abhilash1910 Apr 3, 2024
d8b7cb7
use onedpl iterator
abhilash1910 Apr 4, 2024
be5eefe
rm duplicate function & iterator
abhilash1910 Apr 5, 2024
66ec770
Update help_function/src/onedpl_test_group_load.cpp
abhilash1910 Apr 9, 2024
146c0d1
Update help_function/src/onedpl_test_group_load.cpp
abhilash1910 Apr 9, 2024
ec5d9a8
Update help_function/src/onedpl_test_group_load.cpp
abhilash1910 Apr 9, 2024
400f48d
rm unwanted test
abhilash1910 Apr 11, 2024
5878cb7
update review comments
abhilash1910 Apr 11, 2024
022e628
Merge branch 'SYCLomatic' into block_load_test
abhilash1910 Apr 11, 2024
ae0e1d6
rebase from upstream
abhilash1910 Apr 11, 2024
d2ae0ee
update tests
abhilash1910 Apr 29, 2024
23b7023
update template param
abhilash1910 Apr 30, 2024
7c2ff7e
fix compile issues
abhilash1910 Apr 30, 2024
fe51693
update & fix all tests
abhilash1910 Apr 30, 2024
c5447c0
revert accidental deletion
abhilash1910 Apr 30, 2024
d99a828
replace striped expected output
abhilash1910 May 1, 2024
5ddb779
update striped test case
abhilash1910 May 2, 2024
ab0277a
use subgroup size
abhilash1910 May 2, 2024
300ac08
re-add to prev line
abhilash1910 May 5, 2024
9c21a36
review commits
abhilash1910 May 6, 2024
7cd159c
refactor tests
abhilash1910 May 6, 2024
2b54c4c
Merge branch 'SYCLomatic' into block_load_test
abhilash1910 May 6, 2024
24e80cc
update test
abhilash1910 May 6, 2024
0f4c9ee
fix return no code exec
abhilash1910 May 6, 2024
890d4ff
fix compile issues and pass test
abhilash1910 May 7, 2024
6ec45aa
add comments
abhilash1910 May 7, 2024
3a5c119
apply barrier for sync and update test case
abhilash1910 May 8, 2024
8b42494
review comments
abhilash1910 May 9, 2024
e79a7fc
update
abhilash1910 May 9, 2024
68618bc
rename file
abhilash1910 May 9, 2024
1905ed9
remove barrier and use output buffer
abhilash1910 May 9, 2024
cf9fdd9
rm unused variables
abhilash1910 May 9, 2024
31382c1
fix test case
abhilash1910 May 9, 2024
5a5881b
add local changes and run tests
abhilash1910 May 9, 2024
7f6ed89
format changes
abhilash1910 May 10, 2024
92650b8
revert rename
abhilash1910 May 10, 2024
1a9d205
clang format
abhilash1910 May 10, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
78 changes: 52 additions & 26 deletions help_function/src/onedpl_test_group_load.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -13,27 +13,9 @@
#include <iostream>
#include <oneapi/dpl/iterator>



bool helper_validation_function(const int* ptr, const char* func_name) {

// Used for validation of output and expected output sequences

for (int i = 0; i < 512; ++i) {
if (ptr[i] != i) {
std::cout << func_name <<" failed\n";
std::ostream_iterator<int> Iter(std::cout, ", ");
std::copy(ptr, ptr + 512, Iter);
std::cout << std::endl;
return false;
}
}

std::cout << func_name <<" pass\n";
return true;
}

bool test_load_blocked_striped(dpct::group::load_algorithm T) {
Comment thread
abhilash1910 marked this conversation as resolved.
Outdated
Comment thread
danhoeflinger marked this conversation as resolved.
Outdated
// Tests dpct::group::load_algorithm::BLOCK_LOAD_DIRECT & dpct::group::load_algorithm::BLOCK_LOAD_STRIPED
// in its entirety as API functions
sycl::queue q;
oneapi::dpl::counting_iterator<int> count_it(0);
sycl::buffer<int, 1> buffer(count_it, count_it + 512);
Expand Down Expand Up @@ -61,10 +43,22 @@ bool test_load_blocked_striped(dpct::group::load_algorithm T) {

sycl::host_accessor data_accessor(buffer, sycl::read_only);
const int *ptr = data_accessor.get_multi_ptr<sycl::access::decorated::yes>();
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Not a big deal from my side, but host_accessor overloads operator[], so we don't really need a pointer.

return helper_validation_function(ptr);
for (int i = 0; i < 512; ++i) {
if (ptr[i] != i) {
std::cout << func_name <<" failed\n";
std::ostream_iterator<int> Iter(std::cout, ", ");
std::copy(ptr, ptr + 512, Iter);
std::cout << std::endl;
return false;
}
}

std::cout <<" pass\n";
return true;
}

bool test_load_subgroup_striped_standalone() {
// Tests dpct::group::load_subgroup_striped as standalone method
sycl::queue q;
int data[512];
for (int i = 0; i < 512; i++) data[i] = i;
Expand All @@ -80,19 +74,39 @@ bool test_load_subgroup_striped_standalone() {
dpct::group::uninitialized_load_subgroup_striped<128>(item, d, thread_data);
// Write thread_data of each work item at index to the global buffer
int global_index = item.get_global_linear_id() * 4; // Each thread_data has 4 elements
for (int i = 0; i < 4; ++i) {
buffer[global_index + i] = thread_data[i];
for (int i = 0; i < 4; ++i) {
dacc[global_index + i] = thread_data[i];
}
});
});
q.wait_and_throw();

sycl::host_accessor data_accessor(buffer, sycl::read_only);
const int *ptr = data_accessor.get_multi_ptr<sycl::access::decorated::yes>();
return helper_validation_function(ptr);
int expected[512];
for (int i = 0; i < 128; i++) {
expected[4 * i + 0] = i;
expected[4 * i + 1] = 4 * i + 1;
expected[4 * i + 2] = 4 * i + 2;
expected[4 * i + 3] = 4 * i + 3;
Comment thread
danhoeflinger marked this conversation as resolved.
Outdated
}
for (int i = 0; i < 512; ++i) {
if (ptr[i] != expected[i]) {
std::cout << func_name <<" failed\n";
std::ostream_iterator<int> Iter(std::cout, ", ");
std::copy(ptr, ptr + 512, Iter);
std::cout << std::endl;
return false;
}
}

std::cout <<" pass\n";
return true;
}

bool test_load_blocked_striped_standalone(dpct::group::load_algorithm T) {
Comment thread
danhoeflinger marked this conversation as resolved.
Outdated
// Tests dpct::group::load_algorithm::BLOCK_LOAD_DIRECT & dpct::group::load_algorithm::BLOCK_LOAD_STRIPED
// as standalone methods
sycl::queue q;
Comment thread
abhilash1910 marked this conversation as resolved.
Outdated
int data[512];
for (int i = 0; i < 512; i++) data[i] = i;
Expand All @@ -109,18 +123,30 @@ bool test_load_blocked_striped_standalone(dpct::group::load_algorithm T) {
// Write thread_data of each work item at index to the global buffer
int global_index = item.get_global_linear_id() * 4; // Each thread_data has 4 elements
for (int i = 0; i < 4; ++i) {
buffer[global_index + i] = thread_data[i];
dacc[global_index + i] = thread_data[i];
}
});
});
q.wait_and_throw();

sycl::host_accessor data_accessor(buffer, sycl::read_only);
const int *ptr = data_accessor.get_multi_ptr<sycl::access::decorated::yes>();
return helper_validation_function(ptr);
for (int i = 0; i < 512; ++i) {
if (ptr[i] != i) {
std::cout << func_name <<" failed\n";
std::ostream_iterator<int> Iter(std::cout, ", ");
std::copy(ptr, ptr + 512, Iter);
std::cout << std::endl;
return false;
}
}

std::cout <<" pass\n";
return true;
}

int main() {

return !(test_load_blocked_striped(dpct::group::load_algorithm::BLOCK_LOAD_DIRECT) && test_load_blocked_striped(dpct::group::load_algorithm::BLOCK_LOAD_STRIPED) && test_load_subgroup_striped_standalone()
&& test_load_blocked_striped_standalone(dpct::group::load_algorithm::BLOCK_LOAD_STRIPED) && test_load_blocked_striped_standalone(dpct::group::load_algorithm::BLOCK_LOAD_DIRECT));
}