fumiama/gozel
1
0
Fork 0
mirror of https://github.com/fumiama/gozel.git synced 2026-06-12 04:20:28 +08:00
Code Activity

fix(examples): some platform not support double

Browse Source
This commit is contained in:
源文雨
2026-03-25 17:20:16 +08:00
parent d86e2a59d5
commit 32ada81b80
5 changed files with 67 additions and 47 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
examples/vadd/main.cpp
View File

@@ -1,7 +1,7 @@
#include <sycl/sycl.hpp> #include <sycl/sycl.hpp>
extern "C" SYCL_EXTERNAL extern "C" SYCL_EXTERNAL
void vector_add(double* a, double* b) { void vector_add(float* a, float* b) {
auto item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); auto item = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
int idx = item.get_global_id(0); int idx = item.get_global_id(0);

19
examples/vadd/main.go
View File

@@ -11,28 +11,27 @@ import (
"github.com/fumiama/gozel/ze" "github.com/fumiama/gozel/ze"
) )
//go:generate clang++ -fsycl -fsycl-device-only -fno-sycl-use-footer -faddrsig -Xclang -emit-llvm-bc main.cpp -o device_func.bc //go:generate clang++ -fsycl -fsycl-device-only -fsycl-targets=spirv64 -faddrsig -Xclang -emit-llvm-bc main.cpp -o device_func.bc
//go:generate sycl-post-link -symbols -split=auto -o device_func.table device_func.bc //go:generate sycl-post-link -symbols -split=auto -o device_func.table device_func.bc
//go:generate llvm-spirv -o device_func.spv device_func_0.bc //go:generate clang++ -target spirv64-unknown-unknown -S -emit-llvm -x ir device_func_0.bc -o device_func.ll
//go:generate clang++ -target spir64-unknown-unknown -S -emit-llvm -x ir device_func_0.bc -o device_func.ll
//go:generate go run ../../cmd/func2kernel device_func.ll device_kern.ll //go:generate go run ../../cmd/func2kernel device_func.ll device_kern.ll
//go:generate clang++ -target spir64-unknown-unknown -c -emit-llvm -x ir device_kern.ll -o device_kern.bc //go:generate clang++ -target spirv64-unknown-unknown -c -emit-llvm -x ir device_kern.ll -o device_kern.bc
//go:generate llvm-spirv -o main.spv device_kern.bc //go:generate llvm-spirv -o main.spv device_kern.bc
//go:generate clang++ -target spir64-unknown-unknown -S -emit-llvm -x ir device_kern.bc -o main.ll //go:generate clang++ -target spirv64-unknown-unknown -S -emit-llvm -x ir device_kern.bc -o main.ll
//go:embed main.spv //go:embed main.spv
var kernelspv []byte var kernelspv []byte
const ( const (
X, Y, Z = 1024, 1, 1 X, Y, Z = 64, 1, 1
N = X * Y * Z N = X * Y * Z
bufsz = N * unsafe.Sizeof(float64(0)) bufsz = N * unsafe.Sizeof(float32(0))
) )
func main() { func main() {
floatbuf := make([]float64, 2*N) floatbuf := make([]float32, 2*N)
for i := range floatbuf { for i := range floatbuf {
floatbuf[i] = rand.Float64() floatbuf[i] = rand.Float32()
} }
gpus, err := ze.InitGPUDrivers() gpus, err := ze.InitGPUDrivers()
@@ -88,7 +87,7 @@ func main() {
} }
defer ctx.MemFree(dbuf_v2) defer ctx.MemFree(dbuf_v2)
zev1, zev2 := unsafe.Slice((*float64)(hbuf_v1), N), unsafe.Slice((*float64)(hbuf_v2), N) zev1, zev2 := unsafe.Slice((*float32)(hbuf_v1), N), unsafe.Slice((*float32)(hbuf_v2), N)
copy(zev1, floatbuf[:N]) copy(zev1, floatbuf[:N])
copy(zev2, floatbuf[N:]) copy(zev2, floatbuf[N:])

67
examples/vadd/main.ll
View File

@@ -1,21 +1,21 @@
; ModuleID = 'device_kern.bc' ; ModuleID = 'device_kern.bc'
source_filename = "main.cpp" source_filename = "main.cpp"
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
target triple = "spir64-unknown-unknown" target triple = "spirv64-unknown-unknown"
@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 @__spirv_BuiltInGlobalInvocationId = 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) ; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite, inaccessiblemem: write)
define dso_local spir_kernel void @vector_add(ptr addrspace(1) noundef captures(none) %0, ptr addrspace(1) noundef readonly captures(none) %1) local_unnamed_addr #0 !sycl_used_aspects !8 !sycl_fixed_targets !10 { define spir_kernel void @vector_add(ptr addrspace(1) noundef captures(none) %0, ptr addrspace(1) noundef readonly captures(none) %1) local_unnamed_addr #0 !sycl_fixed_targets !9 {
%3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32, !noalias !11 %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32, !noalias !10
%4 = icmp ult i64 %3, 2147483648 %4 = icmp ult i64 %3, 2147483648
tail call void @llvm.assume(i1 %4) tail call void @llvm.assume(i1 %4)
%5 = getelementptr inbounds nuw double, ptr addrspace(1) %1, i64 %3 %5 = getelementptr inbounds nuw float, ptr addrspace(1) %1, i64 %3
%6 = load double, ptr addrspace(1) %5, align 8 %6 = load float, ptr addrspace(1) %5, align 4
%7 = getelementptr inbounds nuw double, ptr addrspace(1) %0, i64 %3 %7 = getelementptr inbounds nuw float, ptr addrspace(1) %0, i64 %3
%8 = load double, ptr addrspace(1) %7, align 8 %8 = load float, ptr addrspace(1) %7, align 4
%9 = fadd double %8, %6 %9 = fadd float %8, %6
store double %9, ptr addrspace(1) %7, align 8 store float %9, ptr addrspace(1) %7, align 4
ret void ret void
} }
@@ -25,28 +25,27 @@ declare void @llvm.assume(i1 noundef) #1
attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite, inaccessiblemem: write) "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-entry-point" "sycl-module-id"="main.cpp" "sycl-optlevel"="2" } attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite, inaccessiblemem: write) "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-entry-point" "sycl-module-id"="main.cpp" "sycl-optlevel"="2" }
attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) } attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) }
!llvm.dependent-libraries = !{!0} !llvm.linker.options = !{!0, !1}
!llvm.module.flags = !{!1, !2, !3} !llvm.module.flags = !{!2, !3, !4}
!opencl.spir.version = !{!4} !opencl.spir.version = !{!5}
!spirv.Source = !{!5} !spirv.Source = !{!6}
!llvm.ident = !{!6} !llvm.ident = !{!7}
!sycl-esimd-split-status = !{!7} !sycl-esimd-split-status = !{!8}
!0 = !{!"libcpmt"} !0 = !{!"-llibcpmt"}
!1 = !{i32 1, !"wchar_size", i32 2} !1 = !{!"/alternatename:_Avx2WmemEnabled=_Avx2WmemEnabledWeakValue"}
!2 = !{i32 1, !"sycl-device", i32 1} !2 = !{i32 1, !"wchar_size", i32 2}
!3 = !{i32 7, !"frame-pointer", i32 2} !3 = !{i32 1, !"sycl-device", i32 1}
!4 = !{i32 1, i32 2} !4 = !{i32 7, !"frame-pointer", i32 2}
!5 = !{i32 4, i32 100000} !5 = !{i32 1, i32 2}
!6 = !{!"clang version 21.0.0git (https://github.com/intel/llvm d5f649b706f63b5c74e1929bc95db8de91085560)"} !6 = !{i32 4, i32 100000}
!7 = !{i8 0} !7 = !{!"clang version 21.0.0git (https://github.com/intel/llvm d5f649b706f63b5c74e1929bc95db8de91085560)"}
!8 = !{!9} !8 = !{i8 0}
!9 = !{!"fp64", i32 6} !9 = !{}
!10 = !{} !10 = !{!11, !13, !15}
!11 = !{!12, !14, !16} !11 = distinct !{!11, !12, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv: argument 0"}
!12 = distinct !{!12, !13, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv: argument 0"} !12 = distinct !{!12, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv"}
!13 = distinct !{!13, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv"} !13 = distinct !{!13, !14, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v: argument 0"}
!14 = distinct !{!14, !15, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v: argument 0"} !14 = distinct !{!14, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v"}
!15 = distinct !{!15, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v"} !15 = distinct !{!15, !16, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv: argument 0"}
!16 = distinct !{!16, !17, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv: argument 0"} !16 = distinct !{!16, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv"}
!17 = distinct !{!17, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv"}

BIN
examples/vadd/main.spv
View File

Binary file not shown.

26
ze/module.go
View File

@@ -1,7 +1,9 @@
package ze package ze
import ( import (
"errors"
"runtime" "runtime"
"strings"
"github.com/fumiama/gozel" "github.com/fumiama/gozel"
) )
@@ -13,14 +15,34 @@ type ModuleHandle gozel.ZeModuleHandle
func (h ContextHandle) ModuleCreate(hDevice gozel.ZeDeviceHandle, data []byte) ( func (h ContextHandle) ModuleCreate(hDevice gozel.ZeDeviceHandle, data []byte) (
ModuleHandle, error, ModuleHandle, error,
) { ) {
var m gozel.ZeModuleHandle var (
m gozel.ZeModuleHandle
lg gozel.ZeModuleBuildLogHandle
)
_, err := gozel.ZeModuleCreate(gozel.ZeContextHandle(h), hDevice, &gozel.ZeModuleDesc{ _, err := gozel.ZeModuleCreate(gozel.ZeContextHandle(h), hDevice, &gozel.ZeModuleDesc{
Stype: gozel.ZE_STRUCTURE_TYPE_MODULE_DESC, Stype: gozel.ZE_STRUCTURE_TYPE_MODULE_DESC,
Format: gozel.ZE_MODULE_FORMAT_IL_SPIRV, Format: gozel.ZE_MODULE_FORMAT_IL_SPIRV,
Inputsize: uintptr(len(data)), Inputsize: uintptr(len(data)),
Pinputmodule: &data[0], Pinputmodule: &data[0],
}, &m, nil) }, &m, &lg)
runtime.KeepAlive(data) runtime.KeepAlive(data)
defer gozel.ZeModuleBuildLogDestroy(lg)
if err != nil {
var lgsz uintptr
_, errlg := gozel.ZeModuleBuildLogGetString(lg, &lgsz, nil)
if errlg == nil {
data := make([]byte, lgsz)
_, errlg := gozel.ZeModuleBuildLogGetString(lg, &lgsz, &data[0])
runtime.KeepAlive(data)
if errlg == nil {
sb := strings.Builder{}
sb.WriteString(err.Error())
sb.WriteString(", build log: ")
sb.Write(data)
err = errors.New(sb.String())
}
}
}
return ModuleHandle(m), err return ModuleHandle(m), err
} }