Skip to content

Commit e808d7f

Browse files
author
Mahmood Yassin
committed
adding test
1 parent dc0ffd7 commit e808d7f

File tree

1 file changed

+25
-0
lines changed

1 file changed

+25
-0
lines changed
Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
// RUN: %clang -cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -finclude-default-header -emit-cir -o - %s -fclangir | FileCheck %s --check-prefix=CIR
2+
// RUN: %clang -cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -finclude-default-header -emit-llvm -o - %s -fclangir | FileCheck %s --check-prefix=LLVM
3+
// RUN: %clang -cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -finclude-default-header -emit-llvm -o - %s | FileCheck %s --check-prefix=OG-LLVM
4+
5+
// Simple kernel using async_work_group_copy + wait_group_events
6+
7+
__kernel void test_async_copy(__global int *g_in, __local int *l_in, int size) {
8+
// int gid = get_global_id(0);
9+
10+
// Trigger async copy: global to local
11+
// event_t e_in =
12+
async_work_group_copy(
13+
l_in, // local destination
14+
g_in,// + gid * size, // global source
15+
size, // number of elements
16+
(event_t)0 // no dependency
17+
);
18+
19+
// Wait for the async operation to complete
20+
// wait_group_events(1, &e_in);
21+
}
22+
23+
// CIR: cir.call @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (!cir.ptr<!s32i, addrspace(offload_local)>, !cir.ptr<!s32i, addrspace(offload_global)>, !u64i, !cir.ptr<!void>) -> !cir.ptr<!void>
24+
// LLVM: call spir_func ptr @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}}, i64 %{{.*}}, ptr null)
25+
// OG-LLVM: call spir_func target("spirv.Event") @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(ptr addrspace(3) noundef %{{.*}}, ptr addrspace(1) noundef %{{.*}}, i64 noundef %{{.*}}, target("spirv.Event") zeroinitializer

0 commit comments

Comments
 (0)