Skip to content

feat(ze): add event support & vadd demo & refactor #5

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
fumiama merged 5 commits into from
Mar 28, 2026
Merged

feat(ze): add event support & vadd demo & refactor #5

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
e51faf7
feat(ze): add event support & vadd demo
fumiama Mar 28, 2026
51f4143
refactor: move generated files into gozel
fumiama Mar 28, 2026
b1800b3
chore: make lint happy
fumiama Mar 28, 2026
f17cc76
chore: run go generate
github-actions[bot] Mar 28, 2026
50b622f
fix: mispalced api.go
fumiama Mar 28, 2026
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
3 changes: 2 additions & 1 deletion README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -190,7 +190,8 @@ Contributions of all kinds are welcome. Some particularly impactful areas:

## License

This project is licensed under the [GNU Affero General Public License v3.0](LICENSE).
- This project is generally licensed under the [GNU Affero General Public License v3.0](LICENSE).
- The files in [gozel](gozel) folder follows their original license, which is [MIT](https://github.com/oneapi-src/level-zero/blob/master/LICENSE).

---

Expand Down
7 changes: 5 additions & 2 deletions cmd/gen/api.go
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,11 +1,14 @@
package main

import "os"
import (
"os"
"path"
)

var apif *os.File

func init() {
f, err := os.Create("api.go")
f, err := os.Create(path.Join("gozel", "api.go"))
if err != nil {
panic(err)
}
Expand Down
3 changes: 2 additions & 1 deletion cmd/gen/scan.go
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@ import (
"fmt"
"io"
"os"
"path"
"strconv"
"strings"
"unicode"
Expand Down Expand Up @@ -134,7 +135,7 @@ func scanHeader(name string, scan *bufio.Scanner) {
}
fmt.Println(infh(name), "scanning region", region)
k := fmt.Sprint(name, "_", region)
f, err := os.Create(fmt.Sprint(k, ".go"))
f, err := os.Create(path.Join("gozel", fmt.Sprint(k, ".go")))
if err != nil {
panic(fmt.Sprintf("%s L%d: cannot create region %s, err: %v", name, ln, region, err))
}
Expand Down
2 changes: 1 addition & 1 deletion examples/vadd/main.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@
extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>))
void vector_add(float* a, float* b) {
auto item = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
int idx = item.get_global_id(0);
int idx = item.get_global_linear_id();

a[idx] += b[idx];
}
26 changes: 13 additions & 13 deletions examples/vadd/main.go
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,14 +4,15 @@ package main
import (
_ "embed"
"fmt"
"math"
"math/rand"
"os"
"strconv"
"strings"
"time"
"unsafe"

"github.com/fumiama/gozel"
"github.com/fumiama/gozel/gozel"
"github.com/fumiama/gozel/ze"
)

Expand Down Expand Up @@ -69,8 +70,7 @@ func main() {
fmt.Printf("%-28s (%d, %d, %d)\n", "Max Group Count (X, Y, Z):", cprop.Maxgroupcountx, cprop.Maxgroupcounty, cprop.Maxgroupcountz)
fmt.Printf("%-28s %d\n", "Max Total Group Size:", cprop.Maxtotalgroupsize)
fmt.Printf("%-28s %d\n", "Max Shared Local Memory:", cprop.Maxsharedlocalmemory)
fmt.Printf("%-28s %d\n", "Num Subgroup Sizes:", cprop.Numsubgroupsizes)
fmt.Printf("%-28s %v\n", "Subgroup Sizes:", cprop.Subgroupsizes[:])
fmt.Printf("%-28s %v\n", "Subgroup Sizes:", cprop.Subgroupsizes[:cprop.Numsubgroupsizes])

var (
X, Y, Z = uintptr(cprop.Maxgroupsizex), uintptr(1), uintptr(1)
Expand All @@ -84,7 +84,7 @@ func main() {
fmt.Printf("%-28s %d\n", "Total Elements (N):", N)
fmt.Printf("%-28s %d MiB\n", "Buffer Size:", bufsz/1024/1024)

q, err := ctx.CommandQueueCreate(dev)
q, err := ctx.CommandQueueCreate(dev, gozel.ZE_COMMAND_QUEUE_MODE_DEFAULT)
if err != nil {
panic(err)
}
Expand Down Expand Up @@ -135,11 +135,11 @@ func main() {
}
defer krn.Destroy()

err = krn.SetArgumentValue(0, unsafe.Sizeof(uintptr(0)), unsafe.Pointer(&dbufV1))
err = krn.SetArgumentValue(0, &dbufV1)
if err != nil {
panic(err)
}
err = krn.SetArgumentValue(1, unsafe.Sizeof(uintptr(0)), unsafe.Pointer(&dbufV2))
err = krn.SetArgumentValue(1, &dbufV2)
if err != nil {
panic(err)
}
Expand All @@ -154,16 +154,16 @@ func main() {
}
defer lstpre.Destroy()

err = lstpre.AppendMemoryCopy(dbufV1, hbufV1, bufsz)
err = lstpre.AppendMemoryCopy(dbufV1, hbufV1, bufsz, 0)
if err != nil {
panic(err)
}
err = lstpre.AppendMemoryCopy(dbufV2, hbufV2, bufsz)
err = lstpre.AppendMemoryCopy(dbufV2, hbufV2, bufsz, 0)
if err != nil {
panic(err)
}

err = lstpre.AppendBarrier()
err = lstpre.AppendBarrier(0)
if err != nil {
panic(err)
}
Expand All @@ -181,12 +181,12 @@ func main() {

err = lstcalc.AppendLaunchKernel(krn, &gozel.ZeGroupCount{
Groupcountx: uint32(groupCount), Groupcounty: 1, Groupcountz: 1,
})
}, 0)
if err != nil {
panic(err)
}

err = lstcalc.AppendBarrier()
err = lstcalc.AppendBarrier(0)
if err != nil {
panic(err)
}
Expand All @@ -202,7 +202,7 @@ func main() {
}
defer lstpost.Destroy()

err = lstpost.AppendMemoryCopy(hbufV1, dbufV1, bufsz)
err = lstpost.AppendMemoryCopy(hbufV1, dbufV1, bufsz, 0)
if err != nil {
panic(err)
}
Expand All @@ -217,7 +217,7 @@ func main() {
if err != nil {
panic(err)
}
err = q.Synchronize()
err = q.Synchronize(math.MaxUint64)
if err != nil {
panic(err)
}
Expand Down
34 changes: 22 additions & 12 deletions examples/vadd/main.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,18 +4,21 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:
target triple = "spirv64-unknown-unknown"

@__spirv_BuiltInGlobalInvocationId = external local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
@__spirv_BuiltInGlobalOffset = external local_unnamed_addr addrspace(1) constant <3 x i64>, align 32

; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite, inaccessiblemem: write)
define spir_kernel void @__sycl_kernel_vector_add(ptr addrspace(1) noundef align 4 captures(none) %0, ptr addrspace(1) noundef readonly align 4 captures(none) %1) local_unnamed_addr #0 !kernel_arg_buffer_location !6 !sycl_fixed_targets !7 !sycl_kernel_omit_args !8 {
%3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32, !noalias !9
%4 = icmp ult i64 %3, 2147483648
tail call void @llvm.assume(i1 %4)
%5 = getelementptr inbounds float, ptr addrspace(1) %1, i64 %3
%6 = load float, ptr addrspace(1) %5, align 4, !tbaa !16
%7 = getelementptr inbounds float, ptr addrspace(1) %0, i64 %3
%8 = load float, ptr addrspace(1) %7, align 4, !tbaa !16
%9 = fadd float %8, %6
store float %9, ptr addrspace(1) %7, align 4, !tbaa !16
%4 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalOffset, align 32, !noalias !16
%5 = sub i64 %3, %4
%6 = icmp ult i64 %5, 2147483648
tail call void @llvm.assume(i1 %6)
%7 = getelementptr inbounds float, ptr addrspace(1) %1, i64 %5
%8 = load float, ptr addrspace(1) %7, align 4, !tbaa !23
%9 = getelementptr inbounds float, ptr addrspace(1) %0, i64 %5
%10 = load float, ptr addrspace(1) %9, align 4, !tbaa !23
%11 = fadd float %10, %8
store float %11, ptr addrspace(1) %9, align 4, !tbaa !23
ret void
}

Expand Down Expand Up @@ -48,7 +51,14 @@ attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessib
!13 = distinct !{!13, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v"}
!14 = distinct !{!14, !15, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv: argument 0"}
!15 = distinct !{!15, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv"}
!16 = !{!17, !17, i64 0}
!17 = !{!"float", !18, i64 0}
!18 = !{!"omnipotent char", !19, i64 0}
!19 = !{!"Simple C++ TBAA"}
!16 = !{!17, !19, !21}
!17 = distinct !{!17, !18, !"_ZN7__spirv23InitSizesSTGlobalOffsetILi1EN4sycl3_V12idILi1EEEE8initSizeEv: argument 0"}
!18 = distinct !{!18, !"_ZN7__spirv23InitSizesSTGlobalOffsetILi1EN4sycl3_V12idILi1EEEE8initSizeEv"}
!19 = distinct !{!19, !20, !"_ZN7__spirv16initGlobalOffsetILi1EN4sycl3_V12idILi1EEEEET0_v: argument 0"}
!20 = distinct !{!20, !"_ZN7__spirv16initGlobalOffsetILi1EN4sycl3_V12idILi1EEEEET0_v"}
!21 = distinct !{!21, !22, !"_ZNK4sycl3_V17nd_itemILi1EE10get_offsetEv: argument 0"}
!22 = distinct !{!22, !"_ZNK4sycl3_V17nd_itemILi1EE10get_offsetEv"}
!23 = !{!24, !24, i64 0}
!24 = !{!"float", !25, i64 0}
!25 = !{!"omnipotent char", !26, i64 0}
!26 = !{!"Simple C++ TBAA"}
Binary file modified examples/vadd/main.spv
View file
Open in desktop
Binary file not shown.
1 change: 1 addition & 0 deletions examples/vadd_event/.gitignore
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1 @@
/device*
9 changes: 9 additions & 0 deletions examples/vadd_event/main.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,9 @@
#include <sycl/sycl.hpp>

extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>))
void vector_add(float* a, float* b) {
auto item = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
int idx = item.get_global_linear_id();

a[idx] += b[idx];
}
Loading
Loading