From 32ada81b801bbab2f90ac8fbba112efcb914510b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=BA=90=E6=96=87=E9=9B=A8?= <41315874+fumiama@users.noreply.github.com> Date: Wed, 25 Mar 2026 17:20:16 +0800 Subject: [PATCH] fix(examples): some platform not support double --- examples/vadd/main.cpp | 2 +- examples/vadd/main.go | 19 ++++++------ examples/vadd/main.ll | 67 ++++++++++++++++++++--------------------- examples/vadd/main.spv | Bin 924 -> 916 bytes ze/module.go | 26 ++++++++++++++-- 5 files changed, 67 insertions(+), 47 deletions(-) diff --git a/examples/vadd/main.cpp b/examples/vadd/main.cpp index c755b18..b81eae0 100644 --- a/examples/vadd/main.cpp +++ b/examples/vadd/main.cpp @@ -1,7 +1,7 @@ #include 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>(); int idx = item.get_global_id(0); diff --git a/examples/vadd/main.go b/examples/vadd/main.go index 6495f89..542d892 100644 --- a/examples/vadd/main.go +++ b/examples/vadd/main.go @@ -11,28 +11,27 @@ import ( "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 llvm-spirv -o device_func.spv device_func_0.bc -//go:generate clang++ -target spir64-unknown-unknown -S -emit-llvm -x ir device_func_0.bc -o device_func.ll +//go:generate clang++ -target spirv64-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 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 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 var kernelspv []byte const ( - X, Y, Z = 1024, 1, 1 + X, Y, Z = 64, 1, 1 N = X * Y * Z - bufsz = N * unsafe.Sizeof(float64(0)) + bufsz = N * unsafe.Sizeof(float32(0)) ) func main() { - floatbuf := make([]float64, 2*N) + floatbuf := make([]float32, 2*N) for i := range floatbuf { - floatbuf[i] = rand.Float64() + floatbuf[i] = rand.Float32() } gpus, err := ze.InitGPUDrivers() @@ -88,7 +87,7 @@ func main() { } 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(zev2, floatbuf[N:]) diff --git a/examples/vadd/main.ll b/examples/vadd/main.ll index 68cf7ad..3dcc254 100644 --- a/examples/vadd/main.ll +++ b/examples/vadd/main.ll @@ -1,21 +1,21 @@ ; ModuleID = 'device_kern.bc' 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 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) -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 { - %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32, !noalias !11 +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 !10 %4 = icmp ult i64 %3, 2147483648 tail call void @llvm.assume(i1 %4) - %5 = getelementptr inbounds nuw double, ptr addrspace(1) %1, i64 %3 - %6 = load double, ptr addrspace(1) %5, align 8 - %7 = getelementptr inbounds nuw double, ptr addrspace(1) %0, i64 %3 - %8 = load double, ptr addrspace(1) %7, align 8 - %9 = fadd double %8, %6 - store double %9, ptr addrspace(1) %7, align 8 + %5 = getelementptr inbounds nuw float, ptr addrspace(1) %1, i64 %3 + %6 = load float, ptr addrspace(1) %5, align 4 + %7 = getelementptr inbounds nuw float, ptr addrspace(1) %0, i64 %3 + %8 = load float, ptr addrspace(1) %7, align 4 + %9 = fadd float %8, %6 + store float %9, ptr addrspace(1) %7, align 4 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 #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) } -!llvm.dependent-libraries = !{!0} -!llvm.module.flags = !{!1, !2, !3} -!opencl.spir.version = !{!4} -!spirv.Source = !{!5} -!llvm.ident = !{!6} -!sycl-esimd-split-status = !{!7} +!llvm.linker.options = !{!0, !1} +!llvm.module.flags = !{!2, !3, !4} +!opencl.spir.version = !{!5} +!spirv.Source = !{!6} +!llvm.ident = !{!7} +!sycl-esimd-split-status = !{!8} -!0 = !{!"libcpmt"} -!1 = !{i32 1, !"wchar_size", i32 2} -!2 = !{i32 1, !"sycl-device", i32 1} -!3 = !{i32 7, !"frame-pointer", i32 2} -!4 = !{i32 1, i32 2} -!5 = !{i32 4, i32 100000} -!6 = !{!"clang version 21.0.0git (https://github.com/intel/llvm d5f649b706f63b5c74e1929bc95db8de91085560)"} -!7 = !{i8 0} -!8 = !{!9} -!9 = !{!"fp64", i32 6} -!10 = !{} -!11 = !{!12, !14, !16} -!12 = distinct !{!12, !13, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv: argument 0"} -!13 = distinct !{!13, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv"} -!14 = distinct !{!14, !15, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v: argument 0"} -!15 = distinct !{!15, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v"} -!16 = distinct !{!16, !17, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv: argument 0"} -!17 = distinct !{!17, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv"} +!0 = !{!"-llibcpmt"} +!1 = !{!"/alternatename:_Avx2WmemEnabled=_Avx2WmemEnabledWeakValue"} +!2 = !{i32 1, !"wchar_size", i32 2} +!3 = !{i32 1, !"sycl-device", i32 1} +!4 = !{i32 7, !"frame-pointer", i32 2} +!5 = !{i32 1, i32 2} +!6 = !{i32 4, i32 100000} +!7 = !{!"clang version 21.0.0git (https://github.com/intel/llvm d5f649b706f63b5c74e1929bc95db8de91085560)"} +!8 = !{i8 0} +!9 = !{} +!10 = !{!11, !13, !15} +!11 = distinct !{!11, !12, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv: argument 0"} +!12 = distinct !{!12, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEE8initSizeEv"} +!13 = distinct !{!13, !14, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v: argument 0"} +!14 = distinct !{!14, !"_ZN7__spirv22initGlobalInvocationIdILi1EN4sycl3_V12idILi1EEEEET0_v"} +!15 = distinct !{!15, !16, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv: argument 0"} +!16 = distinct !{!16, !"_ZNK4sycl3_V17nd_itemILi1EE13get_global_idEv"} diff --git a/examples/vadd/main.spv b/examples/vadd/main.spv index 400381eb3e532db5dffefd61441a55c434a7e17c..51c5691b5ccb8cb354ab0137c4196762729fa818 100644 GIT binary patch delta 34 ncmbQkK81aP!A6S{jEoAC8JV6lvP^!-qzxoRnbm=$Ept2ovgrv6 delta 36 pcmbQjK8JmR0VCH&&GU?m4wHW}K4;{Z{E$f-Nb)kP14&cncmU1_3K;+Z diff --git a/ze/module.go b/ze/module.go index be9d416..914d5ce 100644 --- a/ze/module.go +++ b/ze/module.go @@ -1,7 +1,9 @@ package ze import ( + "errors" "runtime" + "strings" "github.com/fumiama/gozel" ) @@ -13,14 +15,34 @@ type ModuleHandle gozel.ZeModuleHandle func (h ContextHandle) ModuleCreate(hDevice gozel.ZeDeviceHandle, data []byte) ( ModuleHandle, error, ) { - var m gozel.ZeModuleHandle + var ( + m gozel.ZeModuleHandle + lg gozel.ZeModuleBuildLogHandle + ) _, err := gozel.ZeModuleCreate(gozel.ZeContextHandle(h), hDevice, &gozel.ZeModuleDesc{ Stype: gozel.ZE_STRUCTURE_TYPE_MODULE_DESC, Format: gozel.ZE_MODULE_FORMAT_IL_SPIRV, Inputsize: uintptr(len(data)), Pinputmodule: &data[0], - }, &m, nil) + }, &m, &lg) 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 }