Skip to content

Commit

Permalink
add loops to get around task sequence destructor loop getting optimiz…
Browse files Browse the repository at this point in the history 
…ed away without balanced property for simple designs
  • Loading branch information
@bowenxue-intel
bowenxue-intel committed Feb 27, 2024
1 parent 8b30a48 commit 779cbce
Show file tree
Hide file tree
Showing 3 changed files with 44 additions and 34 deletions.
30 changes: 17 additions & 13 deletions sycl/test/check_device_code/task_sequence_intel_balanced.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s

// CHECK: [[TASK_SEQUENCE:%.*]] ={{.*}} call spir_func target("spirv.TaskSequenceINTEL") @_Z31__spirv_TaskSequenceCreateINTEL{{.*}}(ptr{{.*}}@_Z8arrayAdd{{.*}}, i32 -1, i32 -1, i32 2, i32 2)
// CHECK: [[TASK_SEQUENCE:%.*]] ={{.*}} call spir_func target("spirv.TaskSequenceINTEL") @_Z31__spirv_TaskSequenceCreateINTEL{{.*}}(ptr{{.*}}@_Z8arrayAdd{{.*}}, i32 -1, i32 -1, i32 0, i32 128)
// CHECK: call spir_func void @_Z30__spirv_TaskSequenceAsyncINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]], ptr addrspace(4) {{.*}}, ptr addrspace(4) {{.*}}, i32 128)
// CHECK-COUNT-1: call spir_func i32 @_Z28__spirv_TaskSequenceGetINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
// CHECK: call spir_func void @_Z32__spirv_TaskSequenceReleaseINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
Expand All @@ -10,7 +10,7 @@
using namespace sycl::ext::intel::experimental;
using namespace sycl::ext::oneapi::experimental;

constexpr int NSIZE = 128;
constexpr int kSize = 128;

int arrayAdd(int *data1, int *data2, int N) {
int ret = 0;
Expand All @@ -23,20 +23,24 @@ int arrayAdd(int *data1, int *data2, int N) {

int main() {
sycl::queue myQueue;
int result = 0;
std::vector<int> results(kSize);
myQueue.submit([&](sycl::handler &cgh) {
sycl::buffer<int, 1> result_sycl(&result, sycl::range<1>(1));
auto result_acc = result_sycl.get_access<sycl::access::mode::write>(cgh);
cgh.single_task([=](sycl::kernel_handler kh) {
int d1[NSIZE], d2[NSIZE];
sycl::buffer buffer_results(results);
sycl::accessor results_acc(buffer_results, sycl::write_only, sycl::no_init);
cgh.single_task([=]() {
int d1[kSize], d2[kSize];
task_sequence<arrayAdd,
decltype(properties{balanced, invocation_capacity<2>,
response_capacity<2>})>
sot_object;
sot_object.async(d1, d2, NSIZE);
result_acc[0] = sot_object.get();
decltype(properties{balanced, invocation_capacity<kSize>})>
arrayAddTask;
for (int i = 0; i < kSize; i++) {
arrayAddTask.async(d1, d2, kSize);
}

for (int i = 0; i < kSize; i++) {
results_acc[i] = arrayAddTask.get();
}
});
});
myQueue.wait();
return result;
return 0;
}
30 changes: 17 additions & 13 deletions sycl/test/check_device_code/task_sequence_intel_explicit_get.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s

// CHECK: [[TASK_SEQUENCE:%.*]] ={{.*}} call spir_func target("spirv.TaskSequenceINTEL") @_Z31__spirv_TaskSequenceCreateINTEL{{.*}}(ptr{{.*}}@_Z8arrayAdd{{.*}}, i32 -1, i32 -1, i32 2, i32 2)
// CHECK: [[TASK_SEQUENCE:%.*]] ={{.*}} call spir_func target("spirv.TaskSequenceINTEL") @_Z31__spirv_TaskSequenceCreateINTEL{{.*}}(ptr{{.*}}@_Z8arrayAdd{{.*}}, i32 -1, i32 -1, i32 0, i32 128)
// CHECK: call spir_func void @_Z30__spirv_TaskSequenceAsyncINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]], ptr addrspace(4) {{.*}}, ptr addrspace(4) {{.*}}, i32 128)
// CHECK-COUNT-2: call spir_func i32 @_Z28__spirv_TaskSequenceGetINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
// CHECK: call spir_func void @_Z32__spirv_TaskSequenceReleaseINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
Expand All @@ -10,7 +10,7 @@
using namespace sycl::ext::intel::experimental;
using namespace sycl::ext::oneapi::experimental;

constexpr int NSIZE = 128;
constexpr int kSize = 128;

int arrayAdd(int *data1, int *data2, int N) {
int ret = 0;
Expand All @@ -23,19 +23,23 @@ int arrayAdd(int *data1, int *data2, int N) {

int main() {
sycl::queue myQueue;
int result = 0;
std::vector<int> results(kSize);
myQueue.submit([&](sycl::handler &cgh) {
sycl::buffer<int, 1> result_sycl(&result, sycl::range<1>(1));
auto result_acc = result_sycl.get_access<sycl::access::mode::write>(cgh);
cgh.single_task([=](sycl::kernel_handler kh) {
int d1[NSIZE], d2[NSIZE];
task_sequence<arrayAdd, decltype(properties{invocation_capacity<2>,
response_capacity<2>})>
sot_object;
sot_object.async(d1, d2, NSIZE);
result_acc[0] = sot_object.get();
sycl::buffer buffer_results(results);
sycl::accessor results_acc(buffer_results, sycl::write_only, sycl::no_init);
cgh.single_task([=]() {
int d1[kSize], d2[kSize];
task_sequence<arrayAdd, decltype(properties{invocation_capacity<kSize>})>
arrayAddTask;
for (int i = 0; i < kSize; i++) {
arrayAddTask.async(d1, d2, kSize);
}

for (int i = 0; i < kSize; i++) {
results_acc[i] = arrayAddTask.get();
}
});
});
myQueue.wait();
return result;
return 0;
}
18 changes: 10 additions & 8 deletions sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s

// CHECK: [[TASK_SEQUENCE:%.*]] ={{.*}} call spir_func target("spirv.TaskSequenceINTEL") @_Z31__spirv_TaskSequenceCreateINTEL{{.*}}(ptr{{.*}}@_Z8arrayAdd{{.*}}, i32 -1, i32 -1, i32 2, i32 2)
// CHECK: [[TASK_SEQUENCE:%.*]] ={{.*}} call spir_func target("spirv.TaskSequenceINTEL") @_Z31__spirv_TaskSequenceCreateINTEL{{.*}}(ptr{{.*}}@_Z8arrayAdd{{.*}}, i32 0, i32 1, i32 1, i32 1)
// CHECK: call spir_func void @_Z30__spirv_TaskSequenceAsyncINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]], ptr addrspace(4) {{.*}}, ptr addrspace(4) {{.*}}, i32 128)
// CHECK-COUNT-1: call spir_func i32 @_Z28__spirv_TaskSequenceGetINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
// CHECK: call spir_func void @_Z32__spirv_TaskSequenceReleaseINTEL{{.*}}(target("spirv.TaskSequenceINTEL") [[TASK_SEQUENCE]])
Expand All @@ -10,7 +10,7 @@
using namespace sycl::ext::intel::experimental;
using namespace sycl::ext::oneapi::experimental;

constexpr int NSIZE = 128;
constexpr int kSize = 128;

int arrayAdd(int *data1, int *data2, int N) {
int ret = 0;
Expand All @@ -25,12 +25,14 @@ int main() {
sycl::queue myQueue;

myQueue.submit([&](sycl::handler &cgh) {
cgh.single_task([=](sycl::kernel_handler kh) {
int d1[NSIZE], d2[NSIZE];
task_sequence<arrayAdd, decltype(properties{invocation_capacity<2>,
response_capacity<2>})>
sot_object;
sot_object.async(d1, d2, NSIZE);
cgh.single_task([=]() {
int d1[kSize], d2[kSize];
task_sequence<arrayAdd,
decltype(properties{pipelined<0>, use_stall_enable_clusters,
invocation_capacity<1>,
response_capacity<1>})>
arrayAddTask;
arrayAddTask.async(d1, d2, kSize);
});
});
myQueue.wait();
Expand Down

0 comments on commit 779cbce

Please sign in to comment.