2025-05-31 10:11.36: New job: test ahrefs/ocannl https://github.com/ahrefs/ocannl.git#refs/heads/master (d49bc965f5fe7c717cd2c17c3b7f5b0a883d11fb) (linux-x86_64:(lint-fmt)) Base: ocaml/opam:debian-12-ocaml-4.08@sha256:48fa4a7216c3973bb95572cf5dca98cbbcefe90f288f552e7ac70a8ccd438aa7 ocamlformat version: version 0.27.0 (from opam) To reproduce locally: git clone --recursive "https://github.com/ahrefs/ocannl.git" -b "master" && cd "ocannl" && git reset --hard d49bc965 cat > Dockerfile <<'END-OF-DOCKERFILE' FROM ocaml/opam:debian-12-ocaml-4.08@sha256:48fa4a7216c3973bb95572cf5dca98cbbcefe90f288f552e7ac70a8ccd438aa7 USER 1000:1000 RUN cd ~/opam-repository && (git cat-file -e 35eb2f107a989a2d623b0bbe170696398fcb9b1e || git fetch origin master) && git reset -q --hard 35eb2f107a989a2d623b0bbe170696398fcb9b1e && git log --no-decorate -n1 --oneline && opam update -u RUN opam depext -i dune WORKDIR /src RUN opam depext -i ocamlformat=0.27.0 COPY --chown=1000:1000 . /src/ RUN opam exec -- dune build @fmt --ignore-promoted-rules || (echo "dune build @fmt failed"; exit 2) END-OF-DOCKERFILE docker build . END-REPRO-BLOCK 2025-05-31 10:11.36: Using cache hint "ahrefs/ocannl-ocaml/opam:debian-12-ocaml-4.08@sha256:48fa4a7216c3973bb95572cf5dca98cbbcefe90f288f552e7ac70a8ccd438aa7-debian-12-4.08_opam-2.3-ocamlformat-35eb2f107a989a2d623b0bbe170696398fcb9b1e" 2025-05-31 10:11.36: Using OBuilder spec: ((from ocaml/opam:debian-12-ocaml-4.08@sha256:48fa4a7216c3973bb95572cf5dca98cbbcefe90f288f552e7ac70a8ccd438aa7) (user (uid 1000) (gid 1000)) (run (cache (opam-archives (target /home/opam/.opam/download-cache))) (network host) (shell "cd ~/opam-repository && (git cat-file -e 35eb2f107a989a2d623b0bbe170696398fcb9b1e || git fetch origin master) && git reset -q --hard 35eb2f107a989a2d623b0bbe170696398fcb9b1e && git log --no-decorate -n1 --oneline && opam update -u")) (run (cache (opam-archives (target /home/opam/.opam/download-cache))) (network host) (shell "opam depext -i dune")) (workdir /src) (run (cache (opam-archives (target /home/opam/.opam/download-cache))) (network host) (shell "opam depext -i ocamlformat=0.27.0")) (copy (src .) (dst /src/)) (run (shell "opam exec -- dune build @fmt --ignore-promoted-rules || (echo \"dune build @fmt failed\"; exit 2)")) ) 2025-05-31 10:11.36: Waiting for resource in pool OCluster 2025-05-31 10:11.36: Waiting for worker… 2025-05-31 10:11.37: Got resource from pool OCluster Building on asteria.caelum.ci.dev All commits already cached HEAD is now at d49bc965 Fixes #204 and #319: fix emulation of FP8, BF16 via single prec floats in Pure_C_config numerics (from ocaml/opam:debian-12-ocaml-4.08@sha256:48fa4a7216c3973bb95572cf5dca98cbbcefe90f288f552e7ac70a8ccd438aa7) 2025-05-31 10:11.38 ---> using "d1b97f3f32fc7cff4791d73e3fff398d19cc5b0541c709028ff05a921e22d2c8" from cache /: (user (uid 1000) (gid 1000)) /: (run (cache (opam-archives (target /home/opam/.opam/download-cache))) (network host) (shell "cd ~/opam-repository && (git cat-file -e 35eb2f107a989a2d623b0bbe170696398fcb9b1e || git fetch origin master) && git reset -q --hard 35eb2f107a989a2d623b0bbe170696398fcb9b1e && git log --no-decorate -n1 --oneline && opam update -u")) 35eb2f107a Merge pull request #27838 from maiste/release-dune-3.18.2 <><> Updating package repositories ><><><><><><><><><><><><><><><><><><><><><><> [default] Initialised default (at git+file:///home/opam/opam-repository): [INFO] opam 2.1 and 2.2 include many performance and security improvements over 2.0; please consider upgrading (https://opam.ocaml.org/doc/Install.html) Everything as up-to-date as possible (run with --verbose to show unavailable upgrades). However, you may "opam upgrade" these packages explicitly, which will ask permission to downgrade or uninstall the conflicting packages. Nothing to do. # Run eval $(opam env) to update the current shell environment 2025-05-31 10:11.38 ---> using "0ed37ff9180ea5331ab17e3106ce3fc10bf21b69021c66107f159798cee036fc" from cache /: (run (cache (opam-archives (target /home/opam/.opam/download-cache))) (network host) (shell "opam depext -i dune")) # Detecting depexts using vars: arch=x86_64, os=linux, os-distribution=debian, os-family=debian # No extra OS packages requirements found. # All required OS packages found. # Now letting opam install the packages The following actions will be performed: - install dune 3.18.2 <><> Gathering sources ><><><><><><><><><><><><><><><><><><><><><><><><><><><><> [dune.3.18.2] found in cache <><> Processing actions <><><><><><><><><><><><><><><><><><><><><><><><><><><><> -> installed dune.3.18.2 Done. # Run eval $(opam env) to update the current shell environment 2025-05-31 10:11.38 ---> using "dc3dc92082cbd6b1d2902dc06db6b555f449a0b1c6f054b5888f3ddab10a5d9e" from cache /: (workdir /src) /src: (run (cache (opam-archives (target /home/opam/.opam/download-cache))) (network host) (shell "opam depext -i ocamlformat=0.27.0")) # Detecting depexts using vars: arch=x86_64, os=linux, os-distribution=debian, os-family=debian # No extra OS packages requirements found. # All required OS packages found. # Now letting opam install the packages The following actions will be performed: - install sexplib0 v0.14.0 [required by base] - install dune-build-info 3.18.2 [required by ocamlformat-lib] - install cmdliner 1.3.0 [required by ocamlformat] - install menhirLib 20240715 [required by ocamlformat-lib] - install menhirCST 20240715 [required by menhir] - install ocamlbuild 0.16.1 [required by fpath, astring, uuseg] - install menhirSdk 20240715 [required by ocamlformat-lib] - install either 1.0.0 [required by ocamlformat-lib] - install ocaml-version 4.0.0 [required by ocamlformat-lib] - install camlp-streams 5.0.1 [required by ocamlformat-lib] - install csexp 1.5.2 [required by ocamlformat] - install seq base [required by re] - install fix 20250428 [required by ocamlformat-lib] - install ocamlfind 1.9.8 [required by ocp-indent, astring, fpath, uuseg] - install menhir 20240715 [required by ocamlformat-lib] - install dune-configurator 3.18.2 [required by base] - install re 1.11.0 [required by ocamlformat] - install topkg 1.0.8 [required by fpath, astring, uuseg] - install base-bytes base [required by ocp-indent] - install base v0.14.3 [required by ocamlformat-lib] - install uutf 1.0.4 [required by ocamlformat-lib] - install astring 0.8.5 [required by ocamlformat-lib] - install ocp-indent 1.8.1 [required by ocamlformat-lib] - install stdio v0.14.0 [required by ocamlformat-lib] - install uucp 15.0.0 [required by uuseg] - install fpath 0.7.3 [required by ocamlformat-lib] - install uuseg 15.0.0 [required by ocamlformat-lib] - install ocamlformat-lib 0.27.0 [required by ocamlformat] - install ocamlformat 0.27.0 ===== 29 to install ===== <><> Gathering sources ><><><><><><><><><><><><><><><><><><><><><><><><><><><><> [astring.0.8.5] found in cache [base.v0.14.3] found in cache [camlp-streams.5.0.1] found in cache [cmdliner.1.3.0] found in cache [csexp.1.5.2] found in cache [dune-build-info.3.18.2] found in cache [dune-configurator.3.18.2] found in cache [either.1.0.0] found in cache [fix.20250428] found in cache [fpath.0.7.3] found in cache [menhir.20240715] found in cache [menhirCST.20240715] found in cache [menhirLib.20240715] found in cache [menhirSdk.20240715] found in cache [ocaml-version.4.0.0] found in cache [ocamlbuild.0.16.1] found in cache [ocamlfind.1.9.8] found in cache [ocamlformat.0.27.0] found in cache [ocamlformat-lib.0.27.0] found in cache [ocp-indent.1.8.1] found in cache [re.1.11.0] found in cache [sexplib0.v0.14.0] found in cache [stdio.v0.14.0] found in cache [topkg.1.0.8] found in cache [uucp.15.0.0] found in cache [uuseg.15.0.0] found in cache [uutf.1.0.4] found in cache <><> Processing actions <><><><><><><><><><><><><><><><><><><><><><><><><><><><> -> installed seq.base -> installed camlp-streams.5.0.1 -> installed csexp.1.5.2 -> installed cmdliner.1.3.0 -> installed either.1.0.0 -> installed fix.20250428 -> installed menhirCST.20240715 -> installed menhirLib.20240715 -> installed menhirSdk.20240715 -> installed ocaml-version.4.0.0 -> installed re.1.11.0 -> installed sexplib0.v0.14.0 -> installed dune-build-info.3.18.2 -> installed dune-configurator.3.18.2 -> installed ocamlfind.1.9.8 -> installed base-bytes.base -> installed ocamlbuild.0.16.1 -> installed ocp-indent.1.8.1 -> installed base.v0.14.3 -> installed topkg.1.0.8 -> installed stdio.v0.14.0 -> installed uutf.1.0.4 -> installed astring.0.8.5 -> installed fpath.0.7.3 -> installed menhir.20240715 -> installed uucp.15.0.0 -> installed uuseg.15.0.0 -> installed ocamlformat-lib.0.27.0 -> installed ocamlformat.0.27.0 Done. <><> ocp-indent.1.8.1 installed successfully ><><><><><><><><><><><><><><><><><> => This package requires additional configuration for use in editors. Install package 'user-setup', or manually: * for Emacs, add these lines to ~/.emacs: (add-to-list 'load-path "/home/opam/.opam/4.08/share/emacs/site-lisp") (require 'ocp-indent) * for Vim, add this line to ~/.vimrc: set rtp^="/home/opam/.opam/4.08/share/ocp-indent/vim" # Run eval $(opam env) to update the current shell environment 2025-05-31 10:11.38 ---> using "5e3a5282f655843df8e02a2497598f2918acc59345107286c51ff04d386b492b" from cache /src: (copy (src .) (dst /src/)) 2025-05-31 10:11.38 ---> saved as "d76f0196446a323c8c6a42182660ccd90a75b6192cb8dda9317dd8f25ae05524" /src: (run (shell "opam exec -- dune build @fmt --ignore-promoted-rules || (echo \"dune build @fmt failed\"; exit 2)")) File "arrayjit/bin/dune", line 6, characters 21-34: 6 | (pps ppx_minidebug ppx_sexp_conv)) ^^^^^^^^^^^^^ Error: Library "ppx_sexp_conv" not found. -> required by _build/default/arrayjit/bin/read_config.exe -> required by %{dep:../../arrayjit/bin/read_config.exe} at test/dune:25 -> required by _build/default/test/config/ocannl_backend.txt -> required by %{read:config/ocannl_backend.txt} at test/dune:44 -> required by Computing directory contents of _build/default/test File "arrayjit/bin/dune", line 6, characters 7-20: 6 | (pps ppx_minidebug ppx_sexp_conv)) ^^^^^^^^^^^^^ Error: Library "ppx_minidebug" not found. -> required by _build/default/arrayjit/bin/.merlin-conf/exe-read_config -> required by _build/default/arrayjit/bin/read_config.exe -> required by %{dep:../../arrayjit/bin/read_config.exe} at test/dune:25 -> required by _build/default/test/config/ocannl_backend.txt -> required by %{read:config/ocannl_backend.txt} at test/dune:44 -> required by Computing directory contents of _build/default/test File "arrayjit/test/dune", line 1, characters 0-0: diff --git a/_build/default/arrayjit/test/dune b/_build/default/arrayjit/test/.formatted/dune index 77087ce..ce14b1f 100644 --- a/_build/default/arrayjit/test/dune +++ b/_build/default/arrayjit/test/.formatted/dune @@ -9,7 +9,8 @@ (target test_numerical_types.output) (deps test_numerical_types.exe ocannl_config) (action - (with-stdout-to %{target} + (with-stdout-to + %{target} (run %{deps})))) (rule File "arrayjit/test/test_numerical_types.ml", line 1, characters 0-0: diff --git a/_build/default/arrayjit/test/test_numerical_types.ml b/_build/default/arrayjit/test/.formatted/test_numerical_types.ml index 1a5aa7c..2aeecca 100644 --- a/_build/default/arrayjit/test/test_numerical_types.ml +++ b/_build/default/arrayjit/test/.formatted/test_numerical_types.ml @@ -4,44 +4,46 @@ module Ops = Ir.Ops let test_bfloat16_conversions () = Stdio.printf "Testing BFloat16 conversions:\n"; - + (* Test some specific values *) - let test_values = [0.0; 1.0; -1.0; 3.14159; 1e-3; 1e3; Float.infinity; Float.neg_infinity] in - + let test_values = [ 0.0; 1.0; -1.0; 3.14159; 1e-3; 1e3; Float.infinity; Float.neg_infinity ] in + List.iter test_values ~f:(fun orig -> - let bf16 = Ndarray.float_to_bfloat16 orig in - let back = Ndarray.bfloat16_to_float bf16 in - Stdio.printf " %.6f -> 0x%04x -> %.6f\n" orig bf16 back); - + let bf16 = Ndarray.float_to_bfloat16 orig in + let back = Ndarray.bfloat16_to_float bf16 in + Stdio.printf " %.6f -> 0x%04x -> %.6f\n" orig bf16 back); + (* Test round-trip through ndarray *) - let arr = Ndarray.create_array ~debug:"test" Ops.bfloat16 ~dims:[|3; 2|] - (Ops.Constant_fill { values = [|1.0; 2.0; 3.14; -1.5; 0.125; 1000.0|]; strict = true }) in - + let arr = + Ndarray.create_array ~debug:"test" Ops.bfloat16 ~dims:[| 3; 2 |] + (Ops.Constant_fill { values = [| 1.0; 2.0; 3.14; -1.5; 0.125; 1000.0 |]; strict = true }) + in + Stdio.printf "\nBFloat16 array values:\n"; let flat_values = Ndarray.retrieve_flat_values arr in - Array.iteri flat_values ~f:(fun i v -> - Stdio.printf " [%d] = %.6f\n" i v) + Array.iteri flat_values ~f:(fun i v -> Stdio.printf " [%d] = %.6f\n" i v) let test_fp8_conversions () = Stdio.printf "\n\nTesting FP8 conversions:\n"; - + (* Test some specific values *) - let test_values = [0.0; 1.0; -1.0; 0.5; 2.0; 0.125; 16.0; -0.25] in - + let test_values = [ 0.0; 1.0; -1.0; 0.5; 2.0; 0.125; 16.0; -0.25 ] in + List.iter test_values ~f:(fun orig -> - let fp8 = Ndarray.float_to_fp8 orig in - let back = Ndarray.fp8_to_float fp8 in - Stdio.printf " %.6f -> 0x%02x -> %.6f\n" orig fp8 back); - + let fp8 = Ndarray.float_to_fp8 orig in + let back = Ndarray.fp8_to_float fp8 in + Stdio.printf " %.6f -> 0x%02x -> %.6f\n" orig fp8 back); + (* Test round-trip through ndarray *) - let arr = Ndarray.create_array ~debug:"test" Ops.fp8 ~dims:[|2; 2|] - (Ops.Constant_fill { values = [|1.0; 0.5; 2.0; -1.0|]; strict = true }) in - + let arr = + Ndarray.create_array ~debug:"test" Ops.fp8 ~dims:[| 2; 2 |] + (Ops.Constant_fill { values = [| 1.0; 0.5; 2.0; -1.0 |]; strict = true }) + in + Stdio.printf "\nFP8 array values:\n"; let flat_values = Ndarray.retrieve_flat_values arr in - Array.iteri flat_values ~f:(fun i v -> - Stdio.printf " [%d] = %.6f\n" i v) + Array.iteri flat_values ~f:(fun i v -> Stdio.printf " [%d] = %.6f\n" i v) let () = test_bfloat16_conversions (); - test_fp8_conversions () \ No newline at end of file + test_fp8_conversions () File "arrayjit/lib/cc_backend.ml", line 1, characters 0-0: diff --git a/_build/default/arrayjit/lib/cc_backend.ml b/_build/default/arrayjit/lib/.formatted/cc_backend.ml index 74ed2cd..b29e6cd 100644 --- a/_build/default/arrayjit/lib/cc_backend.ml +++ b/_build/default/arrayjit/lib/.formatted/cc_backend.ml @@ -100,7 +100,7 @@ struct (* Override to add our custom type and conversion support *) let typ_of_prec = typ_of_prec - let extra_declarations = extra_declarations (* Our bfloat16/fp8 conversion functions *) + let extra_declarations = extra_declarations (* Our bfloat16/fp8 conversion functions *) let convert_precision = convert_precision end @@ -209,4 +209,4 @@ let%track3_sexp link_compiled ~merge_buffer ~runner_label ctx_arrays (code : pro context_lifetime = (ctx_arrays, code); description = "executes " ^ code.name ^ " on " ^ runner_label; work; - } ) \ No newline at end of file + } ) File "arrayjit/lib/ops.ml", line 1, characters 0-0: diff --git a/_build/default/arrayjit/lib/ops.ml b/_build/default/arrayjit/lib/.formatted/ops.ml index 1dc0079..0444b38 100644 --- a/_build/default/arrayjit/lib/ops.ml +++ b/_build/default/arrayjit/lib/.formatted/ops.ml @@ -17,8 +17,8 @@ type ('ocaml, 'impl) precision = | Uint16 : (int, uint16_elt) precision | Int32 : (int32, int32_elt) precision | Half : (float, float16_elt) precision - | Bfloat16 : (int, uint16_elt) precision (* Using uint16 representation for now *) - | Fp8 : (char, uint8_elt) precision (* Using uint8 representation *) + | Bfloat16 : (int, uint16_elt) precision (* Using uint16 representation for now *) + | Fp8 : (char, uint8_elt) precision (* Using uint8 representation *) | Single : (float, float32_elt) precision | Double : (float, float64_elt) precision [@@deriving sexp_of] @@ -42,7 +42,10 @@ let bfloat16 = Bfloat16_prec Bfloat16 let fp8 = Fp8_prec Fp8 let single = Single_prec Single let double = Double_prec Double -let is_up_to_fp16 = function Half_prec _ | Byte_prec _ | Bfloat16_prec _ | Fp8_prec _ -> true | _ -> false + +let is_up_to_fp16 = function + | Half_prec _ | Byte_prec _ | Bfloat16_prec _ | Fp8_prec _ -> true + | _ -> false let sexp_of_prec = function | Void_prec -> Sexp.Atom "Void_prec" @@ -69,14 +72,14 @@ let prec_of_sexp = function | Sexp.Atom s -> invalid_arg @@ "prec_of_sexp: unknown precision " ^ s let precision_to_string (type ocaml elt_t) (prec : (ocaml, elt_t) precision) = - match prec with - | Byte -> "byte" + match prec with + | Byte -> "byte" | Uint16 -> "uint16" | Int32 -> "int32" - | Half -> "half" + | Half -> "half" | Bfloat16 -> "bfloat16" | Fp8 -> "fp8" - | Single -> "single" + | Single -> "single" | Double -> "double" let prec_string = function @@ -90,8 +93,7 @@ let prec_string = function | Single_prec _ -> "single" | Double_prec _ -> "double" -let prec_of_string s = - prec_of_sexp (Sexp.Atom (String.(capitalize @@ lowercase s) ^ "_prec")) +let prec_of_string s = prec_of_sexp (Sexp.Atom (String.(capitalize @@ lowercase s) ^ "_prec")) let equal_prec p1 p2 = match (p1, p2) with @@ -104,7 +106,16 @@ let equal_prec p1 p2 = | Fp8_prec _, Fp8_prec _ -> true | Single_prec _, Single_prec _ -> true | Double_prec _, Double_prec _ -> true - | Void_prec, _ | Byte_prec _, _ | Uint16_prec _, _ | Int32_prec _, _ | Half_prec _, _ | Bfloat16_prec _, _ | Fp8_prec _, _ | Single_prec _, _ | Double_prec _, _ -> false + | Void_prec, _ + | Byte_prec _, _ + | Uint16_prec _, _ + | Int32_prec _, _ + | Half_prec _, _ + | Bfloat16_prec _, _ + | Fp8_prec _, _ + | Single_prec _, _ + | Double_prec _, _ -> + false let prec_in_bytes = function | Void_prec -> 0 @@ -138,14 +149,14 @@ let promote_prec p1 p2 = | Void_prec, Void_prec -> Void_prec let pack_prec (type ocaml elt_t) (prec : (ocaml, elt_t) precision) = - match prec with - | Byte -> byte + match prec with + | Byte -> byte | Uint16 -> uint16 | Int32 -> int32 - | Half -> half + | Half -> half | Bfloat16 -> bfloat16 | Fp8 -> fp8 - | Single -> single + | Single -> single | Double -> double type 'r map_prec = { f : 'ocaml 'elt_t. ('ocaml, 'elt_t) precision -> 'r } @@ -162,15 +173,15 @@ let map_prec ?default { f } = function | Single_prec Single -> f Single | Double_prec Double -> f Double (* FIXME: this is a hack to get the code to compile. *) - | _ -> invalid_arg "map_prec: unknown precision" + | _ -> invalid_arg "map_prec: unknown precision" let c_typ_of_prec = function | Byte_prec _ -> "unsigned char" | Uint16_prec _ -> "unsigned short" | Int32_prec _ -> "int" | Half_prec _ -> "_Float16" - | Bfloat16_prec _ -> "unsigned short" (* Bfloat16 represented as uint16 *) - | Fp8_prec _ -> "unsigned char" (* FP8 represented as uint8 *) + | Bfloat16_prec _ -> "unsigned short" (* Bfloat16 represented as uint16 *) + | Fp8_prec _ -> "unsigned char" (* FP8 represented as uint8 *) | Single_prec _ -> "float" | Double_prec _ -> "double" | Void_prec -> "void" @@ -387,14 +398,17 @@ let binop_c_syntax prec v = | ToPowOf, _ -> ("powf(", ",", ")") | Relu_gate, (Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> ("(", " > 0 ?", " : 0)") | Relu_gate, _ -> ("(", " > 0.0 ?", " : 0.0)") - | Satur01_gate, (Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> ("(abs(", " ) > 0 ? 0 : (", "))") + | Satur01_gate, (Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> + ("(abs(", " ) > 0 ? 0 : (", "))") | Satur01_gate, Single_prec _ -> (* This disagrees at 0 with the semantics. *) ("(fabsf(floorf(", ")) > 0.0 ? 0.0 : (", "))") | Satur01_gate, _ -> ("(fabs(floor(", ")) > 0.0 ? 0.0 : (", "))") - | Max, (Double_prec _ | Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> ("fmax(", ",", ")") + | Max, (Double_prec _ | Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> + ("fmax(", ",", ")") | Max, _ -> ("fmaxf(", ",", ")") - | Min, (Double_prec _ | Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> ("fmin(", ",", ")") + | Min, (Double_prec _ | Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> + ("fmin(", ",", ")") | Min, _ -> ("fminf(", ",", ")") | Mod, _ -> ("(", " %", ")") | Cmplt, _ -> ("(", " <", ")") @@ -463,7 +477,11 @@ let unop_c_syntax prec op = | Double_prec _ | Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _ -> "fmax" | _ -> "fmaxf" in - let fmin () = match prec with Double_prec _ | Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _ -> "fmin" | _ -> "fminf" in + let fmin () = + match prec with + | Double_prec _ | Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _ -> "fmin" + | _ -> "fminf" + in match (op, prec) with | Identity, _ -> ("", "") | Relu, (Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> ("fmax(0, ", ")") @@ -502,9 +520,11 @@ let ternop_cd_syntax = function Where -> "where" | FMA -> "fma" let ternop_c_syntax prec op = match (op, prec) with - | Where, (Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> ("((", ") != 0 ? (", ") : (", "))") + | Where, (Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> + ("((", ") != 0 ? (", ") : (", "))") | Where, _ -> ("((", ") != 0.0 ? (", ") : (", "))") - | FMA, (Double_prec _ | Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> ("fma(", ",", ",", ")") + | FMA, (Double_prec _ | Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> + ("fma(", ",", ",", ")") | FMA, _ -> ("fmaf(", ",", ",", ")") let c_convert_precision ~from ~to_ = @@ -532,17 +552,16 @@ let c_convert_precision ~from ~to_ = (* Conversions involving BFloat16 and other types *) | Bfloat16_prec _, Half_prec _ -> ("(_Float16)bfloat16_to_float(", ")") | Half_prec _, Bfloat16_prec _ -> ("float_to_bfloat16((float)", ")") - | Bfloat16_prec _, (Byte_prec _ | Uint16_prec _ | Int32_prec _) -> + | Bfloat16_prec _, (Byte_prec _ | Uint16_prec _ | Int32_prec _) -> ("(" ^ c_typ_of_prec to_ ^ ")bfloat16_to_float(", ")") - | (Byte_prec _ | Uint16_prec _ | Int32_prec _), Bfloat16_prec _ -> + | (Byte_prec _ | Uint16_prec _ | Int32_prec _), Bfloat16_prec _ -> ("float_to_bfloat16((float)", ")") (* Conversions involving FP8 and other types *) | Fp8_prec _, Half_prec _ -> ("(_Float16)fp8_to_float(", ")") | Half_prec _, Fp8_prec _ -> ("float_to_fp8((float)", ")") - | Fp8_prec _, (Byte_prec _ | Uint16_prec _ | Int32_prec _) -> + | Fp8_prec _, (Byte_prec _ | Uint16_prec _ | Int32_prec _) -> ("(" ^ c_typ_of_prec to_ ^ ")fp8_to_float(", ")") - | (Byte_prec _ | Uint16_prec _ | Int32_prec _), Fp8_prec _ -> - ("float_to_fp8((float)", ")") + | (Byte_prec _ | Uint16_prec _ | Int32_prec _), Fp8_prec _ -> ("float_to_fp8((float)", ")") (* BFloat16 <-> FP8 conversions *) | Bfloat16_prec _, Fp8_prec _ -> ("float_to_fp8(bfloat16_to_float(", "))") | Fp8_prec _, Bfloat16_prec _ -> ("float_to_bfloat16(fp8_to_float(", "))") File "arrayjit/lib/metal_backend.ml", line 1, characters 0-0: diff --git a/_build/default/arrayjit/lib/metal_backend.ml b/_build/default/arrayjit/lib/.formatted/metal_backend.ml index db449e6..3df32b2 100644 --- a/_build/default/arrayjit/lib/metal_backend.ml +++ b/_build/default/arrayjit/lib/.formatted/metal_backend.ml @@ -444,7 +444,7 @@ end) : Ir.Backend_impl.Lowered_backend = struct | Ops.Uint16_prec _ -> "ushort" | Ops.Int32_prec _ -> "int" | Ops.Half_prec _ -> "half" - | Ops.Bfloat16_prec _ -> "bfloat" (* Metal supports bfloat16 natively *) + | Ops.Bfloat16_prec _ -> "bfloat" (* Metal supports bfloat16 natively *) | Ops.Fp8_prec _ -> invalid_arg "Metal backend does not support FP8 precision" | Ops.Single_prec _ -> "float" | Ops.Double_prec _ -> "double" @@ -455,7 +455,7 @@ end) : Ir.Backend_impl.Lowered_backend = struct | Ops.Uint16_prec _ -> "" | Ops.Int32_prec _ -> "" | Ops.Half_prec _ -> "h" - | Ops.Bfloat16_prec _ -> "bf" (* TODO: Verify actual Metal suffix for bfloat16 *) + | Ops.Bfloat16_prec _ -> "bf" (* TODO: Verify actual Metal suffix for bfloat16 *) | Ops.Fp8_prec _ -> invalid_arg "Metal backend does not support FP8 precision" | Ops.Single_prec _ -> "f" | Ops.Double_prec _ -> "" File "arrayjit/lib/cuda_backend.ml", line 1, characters 0-0: diff --git a/_build/default/arrayjit/lib/cuda_backend.ml b/_build/default/arrayjit/lib/.formatted/cuda_backend.ml index dc04c72..1395aff 100644 --- a/_build/default/arrayjit/lib/cuda_backend.ml +++ b/_build/default/arrayjit/lib/.formatted/cuda_backend.ml @@ -284,8 +284,8 @@ end) : Ir.Backend_impl.Lowered_backend = struct | Ops.Uint16_prec _ -> "unsigned short" | Ops.Int32_prec _ -> "int" | Ops.Half_prec _ -> "__half" - | Ops.Bfloat16_prec _ -> "__nv_bfloat16" (* CUDA bfloat16 type *) - | Ops.Fp8_prec _ -> "__nv_fp8_e5m2" (* CUDA FP8 type (E5M2 format) *) + | Ops.Bfloat16_prec _ -> "__nv_bfloat16" (* CUDA bfloat16 type *) + | Ops.Fp8_prec _ -> "__nv_fp8_e5m2" (* CUDA FP8 type (E5M2 format) *) | Ops.Single_prec _ -> "float" | Ops.Double_prec _ -> "double" | Ops.Void_prec -> "void" @@ -326,8 +326,8 @@ end) : Ir.Backend_impl.Lowered_backend = struct | ToPowOf, Bfloat16_prec _ -> fun v1 v2 -> group - (string "__float2bfloat16(powf(__bfloat162float(" ^^ v1 ^^ string "), __bfloat162float(" - ^^ v2 ^^ string ")))") + (string "__float2bfloat16(powf(__bfloat162float(" + ^^ v1 ^^ string "), __bfloat162float(" ^^ v2 ^^ string ")))") | Relu_gate, (Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> fun v1 v2 -> group @@ -343,15 +343,13 @@ end) : Ir.Backend_impl.Lowered_backend = struct fun v1 v2 -> group (parens - (group - (parens - (string "__bfloat162float(" ^^ v1 ^^ string ") > 0.0f")) + (group (parens (string "__bfloat162float(" ^^ v1 ^^ string ") > 0.0f")) ^^ ifflat (space ^^ string "?" ^^ space ^^ v2 ^^ space ^^ string ":" ^^ space - ^^ string "__float2bfloat16(0.0f)") + ^^ string "__float2bfloat16(0.0f)") (nest 2 (break 1 ^^ string "?" ^^ space ^^ v2 ^^ break 1 ^^ string ":" ^^ space - ^^ string "__float2bfloat16(0.0f)")))) + ^^ string "__float2bfloat16(0.0f)")))) | Satur01_gate, Byte_prec _ -> fun v1 v2 -> group File "arrayjit/lib/ndarray.ml", line 1, characters 0-0: diff --git a/_build/default/arrayjit/lib/ndarray.ml b/_build/default/arrayjit/lib/.formatted/ndarray.ml index 9bdf146..64d85ac 100644 --- a/_build/default/arrayjit/lib/ndarray.ml +++ b/_build/default/arrayjit/lib/.formatted/ndarray.ml @@ -1,5 +1,4 @@ open Base - module Lazy = Utils.Lazy (** N-dimensional arrays: a precision-handling wrapper for [Bigarray.Genarray] and its utilities. *) @@ -34,8 +33,8 @@ type byte_nd = (char, Ops.uint8_elt) bigarray type uint16_nd = (int, Ops.uint16_elt) bigarray [@@ocaml.boxed] type int32_nd = (int32, Ops.int32_elt) bigarray [@@ocaml.boxed] type half_nd = (float, Ops.float16_elt) bigarray -type bfloat16_nd = (int, Ops.uint16_elt) bigarray [@@ocaml.boxed] (* Using uint16 representation *) -type fp8_nd = (char, Ops.uint8_elt) bigarray (* Using uint8 representation *) +type bfloat16_nd = (int, Ops.uint16_elt) bigarray [@@ocaml.boxed] (* Using uint16 representation *) +type fp8_nd = (char, Ops.uint8_elt) bigarray (* Using uint8 representation *) type single_nd = (float, Ops.float32_elt) bigarray type double_nd = (float, Ops.float64_elt) bigarray @@ -48,14 +47,14 @@ let sexp_of_fp8_nd (arr : fp8_nd) = Sexp.Atom (big_ptr_to_string arr) let sexp_of_single_nd (arr : single_nd) = Sexp.Atom (big_ptr_to_string arr) let sexp_of_double_nd (arr : double_nd) = Sexp.Atom (big_ptr_to_string arr) -type t = - | Byte_nd of byte_nd +type t = + | Byte_nd of byte_nd | Uint16_nd of uint16_nd | Int32_nd of int32_nd - | Half_nd of half_nd + | Half_nd of half_nd | Bfloat16_nd of bfloat16_nd | Fp8_nd of fp8_nd - | Single_nd of single_nd + | Single_nd of single_nd | Double_nd of double_nd [@@deriving sexp_of] @@ -78,8 +77,8 @@ let precision_to_bigarray_kind (type ocaml elt_t) (prec : (ocaml, elt_t) Ops.pre | Ops.Uint16 -> Bigarray.Int16_unsigned | Ops.Int32 -> Bigarray.Int32 | Ops.Half -> Bigarray.Float16 - | Ops.Bfloat16 -> Bigarray.Int16_unsigned (* Using uint16 representation *) - | Ops.Fp8 -> Bigarray.Char (* Using uint8 representation *) + | Ops.Bfloat16 -> Bigarray.Int16_unsigned (* Using uint16 representation *) + | Ops.Fp8 -> Bigarray.Char (* Using uint8 representation *) | Ops.Single -> Bigarray.Float32 | Ops.Double -> Bigarray.Float64 @@ -162,22 +161,31 @@ let create_bigarray (type ocaml elt_t) (prec : (ocaml, elt_t) Ops.precision) ~di | Ops.Uint16, Standard_uniform -> init_bigarray_of_prec prec dims ~f:(fun _ -> Random.int 65536) | Ops.Int32, Constant_fill { values; strict } -> constant_fill_f Int32.of_float values strict | Ops.Int32, Range_over_offsets -> - init_bigarray_of_prec prec dims ~f:(fun idcs -> Int32.of_int_exn @@ indices_to_offset ~dims ~idcs) - | Ops.Int32, Standard_uniform -> init_bigarray_of_prec prec dims ~f:(fun _ -> Random.int32 Int32.max_value) + init_bigarray_of_prec prec dims ~f:(fun idcs -> + Int32.of_int_exn @@ indices_to_offset ~dims ~idcs) + | Ops.Int32, Standard_uniform -> + init_bigarray_of_prec prec dims ~f:(fun _ -> Random.int32 Int32.max_value) | Ops.Half, Constant_fill { values; strict } -> constant_fill_float values strict | Ops.Half, Range_over_offsets -> init_bigarray_of_prec prec dims ~f:(fun idcs -> Float.of_int @@ indices_to_offset ~dims ~idcs) | Ops.Half, Standard_uniform -> init_bigarray_of_prec prec dims ~f:(fun _ -> Rand.Lib.float_range 0.0 1.0) - | Ops.Bfloat16, Constant_fill { values; strict } -> constant_fill_f float_to_bfloat16 values strict + | Ops.Bfloat16, Constant_fill { values; strict } -> + constant_fill_f float_to_bfloat16 values strict | Ops.Bfloat16, Range_over_offsets -> - init_bigarray_of_prec prec dims ~f:(fun idcs -> float_to_bfloat16 @@ Float.of_int @@ indices_to_offset ~dims ~idcs) - | Ops.Bfloat16, Standard_uniform -> init_bigarray_of_prec prec dims ~f:(fun _ -> float_to_bfloat16 @@ Rand.Lib.float_range 0.0 1.0) - | Ops.Fp8, Constant_fill { values; strict } -> constant_fill_f (Fn.compose Char.of_int_exn float_to_fp8) values strict + init_bigarray_of_prec prec dims ~f:(fun idcs -> + float_to_bfloat16 @@ Float.of_int @@ indices_to_offset ~dims ~idcs) + | Ops.Bfloat16, Standard_uniform -> + init_bigarray_of_prec prec dims ~f:(fun _ -> + float_to_bfloat16 @@ Rand.Lib.float_range 0.0 1.0) + | Ops.Fp8, Constant_fill { values; strict } -> + constant_fill_f (Fn.compose Char.of_int_exn float_to_fp8) values strict | Ops.Fp8, Range_over_offsets -> init_bigarray_of_prec prec dims ~f:(fun idcs -> Char.of_int_exn @@ float_to_fp8 @@ Float.of_int @@ indices_to_offset ~dims ~idcs) - | Ops.Fp8, Standard_uniform -> init_bigarray_of_prec prec dims ~f:(fun _ -> Char.of_int_exn @@ float_to_fp8 @@ Rand.Lib.float_range 0.0 1.0) + | Ops.Fp8, Standard_uniform -> + init_bigarray_of_prec prec dims ~f:(fun _ -> + Char.of_int_exn @@ float_to_fp8 @@ Rand.Lib.float_range 0.0 1.0) | Ops.Single, Constant_fill { values; strict } -> constant_fill_float values strict | Ops.Single, Range_over_offsets -> init_bigarray_of_prec prec dims ~f:(fun idcs -> Float.of_int @@ indices_to_offset ~dims ~idcs) @@ -319,7 +327,7 @@ let reset_bigarray (init_op : Ops.init_op) (type o b) (prec : (o, b) Ops.precisi | Ops.Uint16, Constant_fill { values; strict } -> constant_set_f Int.of_float values strict | Ops.Uint16, Range_over_offsets -> set_bigarray arr ~f:(fun idcs -> indices_to_offset ~dims ~idcs) - | Ops.Uint16, Standard_uniform -> set_bigarray arr ~f:(fun _ -> Random.int 65536) (* 2^16 *) + | Ops.Uint16, Standard_uniform -> set_bigarray arr ~f:(fun _ -> Random.int 65536) (* 2^16 *) | Ops.Int32, Constant_fill { values; strict } -> constant_set_f Int32.of_float values strict | Ops.Int32, Range_over_offsets -> set_bigarray arr ~f:(fun idcs -> Int32.of_int_exn @@ indices_to_offset ~dims ~idcs) @@ -330,13 +338,17 @@ let reset_bigarray (init_op : Ops.init_op) (type o b) (prec : (o, b) Ops.precisi | Ops.Half, Standard_uniform -> set_bigarray arr ~f:(fun _ -> Rand.Lib.float_range 0.0 1.0) | Ops.Bfloat16, Constant_fill { values; strict } -> constant_set_f float_to_bfloat16 values strict | Ops.Bfloat16, Range_over_offsets -> - set_bigarray arr ~f:(fun idcs -> float_to_bfloat16 @@ Float.of_int @@ indices_to_offset ~dims ~idcs) - | Ops.Bfloat16, Standard_uniform -> set_bigarray arr ~f:(fun _ -> float_to_bfloat16 @@ Rand.Lib.float_range 0.0 1.0) - | Ops.Fp8, Constant_fill { values; strict } -> constant_set_f (Fn.compose Char.of_int_exn float_to_fp8) values strict + set_bigarray arr ~f:(fun idcs -> + float_to_bfloat16 @@ Float.of_int @@ indices_to_offset ~dims ~idcs) + | Ops.Bfloat16, Standard_uniform -> + set_bigarray arr ~f:(fun _ -> float_to_bfloat16 @@ Rand.Lib.float_range 0.0 1.0) + | Ops.Fp8, Constant_fill { values; strict } -> + constant_set_f (Fn.compose Char.of_int_exn float_to_fp8) values strict | Ops.Fp8, Range_over_offsets -> set_bigarray arr ~f:(fun idcs -> Char.of_int_exn @@ float_to_fp8 @@ Float.of_int @@ indices_to_offset ~dims ~idcs) - | Ops.Fp8, Standard_uniform -> set_bigarray arr ~f:(fun _ -> Char.of_int_exn @@ float_to_fp8 @@ Rand.Lib.float_range 0.0 1.0) + | Ops.Fp8, Standard_uniform -> + set_bigarray arr ~f:(fun _ -> Char.of_int_exn @@ float_to_fp8 @@ Rand.Lib.float_range 0.0 1.0) | Ops.Single, Constant_fill { values; strict } -> constant_set_float values strict | Ops.Single, Range_over_offsets -> set_bigarray arr ~f:(fun idcs -> Float.of_int @@ indices_to_offset ~dims ~idcs) @@ -373,8 +385,10 @@ let fold_as_float ~init ~f arr = | Uint16_nd arr -> fold_bigarray ~init ~f:(fun accu idx v -> f accu idx @@ Float.of_int v) arr | Int32_nd arr -> fold_bigarray ~init ~f:(fun accu idx v -> f accu idx @@ Int32.to_float v) arr | Half_nd arr -> fold_bigarray ~init ~f arr - | Bfloat16_nd arr -> fold_bigarray ~init ~f:(fun accu idx v -> f accu idx @@ bfloat16_to_float v) arr - | Fp8_nd arr -> fold_bigarray ~init ~f:(fun accu idx c -> f accu idx @@ fp8_to_float @@ Char.to_int c) arr + | Bfloat16_nd arr -> + fold_bigarray ~init ~f:(fun accu idx v -> f accu idx @@ bfloat16_to_float v) arr + | Fp8_nd arr -> + fold_bigarray ~init ~f:(fun accu idx c -> f accu idx @@ fp8_to_float @@ Char.to_int c) arr | Single_nd arr -> fold_bigarray ~init ~f arr | Double_nd arr -> fold_bigarray ~init ~f arr File "arrayjit/lib/c_syntax.ml", line 1, characters 0-0: diff --git a/_build/default/arrayjit/lib/c_syntax.ml b/_build/default/arrayjit/lib/.formatted/c_syntax.ml index 185cbf6..86279a1 100644 --- a/_build/default/arrayjit/lib/c_syntax.ml +++ b/_build/default/arrayjit/lib/.formatted/c_syntax.ml @@ -91,62 +91,65 @@ struct let arg_int_prefix = "const int " let extra_args = [] let includes = [ "<stdio.h>"; "<stdlib.h>"; "<string.h>"; "<math.h>" ] - let extra_declarations = [ - (* BFloat16 conversion functions *) - "static inline float bfloat16_to_float(unsigned short bf16) {"; - " unsigned int f32 = ((unsigned int)bf16) << 16;"; - " return *((float*)&f32);"; - "}"; - ""; - "static inline unsigned short float_to_bfloat16(float f) {"; - " unsigned int f32 = *((unsigned int*)&f);"; - " unsigned int rounded = f32 + 0x7FFF + ((f32 >> 16) & 1);"; - " return (unsigned short)(rounded >> 16);"; - "}"; - ""; - (* FP8 E5M2 conversion functions *) - "static inline float fp8_to_float(unsigned char fp8) {"; - " if (fp8 == 0) return 0.0f;"; - " unsigned int sign = (fp8 >> 7) & 1;"; - " unsigned int exp = (fp8 >> 2) & 0x1F;"; - " unsigned int mant = fp8 & 0x3;"; - " if (exp == 0x1F) {"; - " if (mant == 0) return sign ? -INFINITY : INFINITY;"; - " else return NAN;"; - " }"; - " if (exp == 0) {"; - " float result = ldexpf((float)mant / 4.0f, -14);"; - " if (sign) result = -result;"; - " return result;"; - " }"; - " float result = (1.0f + (float)mant * 0.25f) * ldexpf(1.0f, (int)exp - 15);"; - " if (sign) result = -result;"; - " return result;"; - "}"; - ""; - "static inline unsigned char float_to_fp8(float f) {"; - " if (f == 0.0f) return 0;"; - " unsigned int sign = (f < 0) ? 1 : 0;"; - " f = fabsf(f);"; - " if (isinf(f)) return (sign << 7) | 0x7C;"; - " if (isnan(f)) return (sign << 7) | 0x7F;"; - " int exp_val;"; - " float mant_f = frexpf(f, &exp_val);"; - " int exp = exp_val + 14;"; - " if (exp < 0) return sign << 7;"; - " if (exp > 30) return (sign << 7) | 0x7C;"; - " if (exp == 0) {"; - " float denorm_mant = f * ldexpf(1.0f, 14) * 4.0f;"; - " unsigned int mant_bits = (unsigned int)(denorm_mant + 0.5f);"; - " if (mant_bits > 3) mant_bits = 3;"; - " return (sign << 7) | mant_bits;"; - " }"; - " mant_f = (mant_f - 0.5f) * 4.0f;"; - " unsigned int mant_bits = (unsigned int)(mant_f + 0.5f);"; - " if (mant_bits > 3) mant_bits = 3;"; - " return (unsigned char)((sign << 7) | ((exp & 0x1F) << 2) | (mant_bits & 0x3));"; - "}"; - ] + + let extra_declarations = + [ + (* BFloat16 conversion functions *) + "static inline float bfloat16_to_float(unsigned short bf16) {"; + " unsigned int f32 = ((unsigned int)bf16) << 16;"; + " return *((float*)&f32);"; + "}"; + ""; + "static inline unsigned short float_to_bfloat16(float f) {"; + " unsigned int f32 = *((unsigned int*)&f);"; + " unsigned int rounded = f32 + 0x7FFF + ((f32 >> 16) & 1);"; + " return (unsigned short)(rounded >> 16);"; + "}"; + ""; + (* FP8 E5M2 conversion functions *) + "static inline float fp8_to_float(unsigned char fp8) {"; + " if (fp8 == 0) return 0.0f;"; + " unsigned int sign = (fp8 >> 7) & 1;"; + " unsigned int exp = (fp8 >> 2) & 0x1F;"; + " unsigned int mant = fp8 & 0x3;"; + " if (exp == 0x1F) {"; + " if (mant == 0) return sign ? -INFINITY : INFINITY;"; + " else return NAN;"; + " }"; + " if (exp == 0) {"; + " float result = ldexpf((float)mant / 4.0f, -14);"; + " if (sign) result = -result;"; + " return result;"; + " }"; + " float result = (1.0f + (float)mant * 0.25f) * ldexpf(1.0f, (int)exp - 15);"; + " if (sign) result = -result;"; + " return result;"; + "}"; + ""; + "static inline unsigned char float_to_fp8(float f) {"; + " if (f == 0.0f) return 0;"; + " unsigned int sign = (f < 0) ? 1 : 0;"; + " f = fabsf(f);"; + " if (isinf(f)) return (sign << 7) | 0x7C;"; + " if (isnan(f)) return (sign << 7) | 0x7F;"; + " int exp_val;"; + " float mant_f = frexpf(f, &exp_val);"; + " int exp = exp_val + 14;"; + " if (exp < 0) return sign << 7;"; + " if (exp > 30) return (sign << 7) | 0x7C;"; + " if (exp == 0) {"; + " float denorm_mant = f * ldexpf(1.0f, 14) * 4.0f;"; + " unsigned int mant_bits = (unsigned int)(denorm_mant + 0.5f);"; + " if (mant_bits > 3) mant_bits = 3;"; + " return (sign << 7) | mant_bits;"; + " }"; + " mant_f = (mant_f - 0.5f) * 4.0f;"; + " unsigned int mant_bits = (unsigned int)(mant_f + 0.5f);"; + " if (mant_bits > 3) mant_bits = 3;"; + " return (unsigned char)((sign << 7) | ((exp & 0x1F) << 2) | (mant_bits & 0x3));"; + "}"; + ] + let typ_of_prec = Ops.c_typ_of_prec let float_log_style = if Input.full_printf_support then "%g" else "%de-3" @@ -221,28 +224,34 @@ struct | Ops.Bfloat16_prec _ -> (* For BFloat16, perform operations in float precision *) let float_v1 = PPrint.(string "bfloat16_to_float(" ^^ v1 ^^ string ")") in - let float_v2 = PPrint.(string "bfloat16_to_float(" ^^ v2 ^^ string ")") in + let float_v2 = PPrint.(string "bfloat16_to_float(" ^^ v2 ^^ string ")") in let float_v3 = PPrint.(string "bfloat16_to_float(" ^^ v3 ^^ string ")") in let op_prefix, op_infix1, op_infix2, op_suffix = Ops.ternop_c_syntax Ops.single op in - let float_result = PPrint.( - group (string op_prefix ^^ float_v1 ^^ string op_infix1 - ^^ ifflat (space ^^ float_v2) (nest 2 (break 1 ^^ float_v2)) - ^^ string op_infix2 - ^^ ifflat (space ^^ float_v3) (nest 2 (break 1 ^^ float_v3)) - ^^ string op_suffix)) in + let float_result = + PPrint.( + group + (string op_prefix ^^ float_v1 ^^ string op_infix1 + ^^ ifflat (space ^^ float_v2) (nest 2 (break 1 ^^ float_v2)) + ^^ string op_infix2 + ^^ ifflat (space ^^ float_v3) (nest 2 (break 1 ^^ float_v3)) + ^^ string op_suffix)) + in PPrint.(string "float_to_bfloat16(" ^^ float_result ^^ string ")") | Ops.Fp8_prec _ -> (* For FP8, perform operations in float precision *) let float_v1 = PPrint.(string "fp8_to_float(" ^^ v1 ^^ string ")") in - let float_v2 = PPrint.(string "fp8_to_float(" ^^ v2 ^^ string ")") in + let float_v2 = PPrint.(string "fp8_to_float(" ^^ v2 ^^ string ")") in let float_v3 = PPrint.(string "fp8_to_float(" ^^ v3 ^^ string ")") in let op_prefix, op_infix1, op_infix2, op_suffix = Ops.ternop_c_syntax Ops.single op in - let float_result = PPrint.( - group (string op_prefix ^^ float_v1 ^^ string op_infix1 - ^^ ifflat (space ^^ float_v2) (nest 2 (break 1 ^^ float_v2)) - ^^ string op_infix2 - ^^ ifflat (space ^^ float_v3) (nest 2 (break 1 ^^ float_v3)) - ^^ string op_suffix)) in + let float_result = + PPrint.( + group + (string op_prefix ^^ float_v1 ^^ string op_infix1 + ^^ ifflat (space ^^ float_v2) (nest 2 (break 1 ^^ float_v2)) + ^^ string op_infix2 + ^^ ifflat (space ^^ float_v3) (nest 2 (break 1 ^^ float_v3)) + ^^ string op_suffix)) + in PPrint.(string "float_to_fp8(" ^^ float_result ^^ string ")") | _ -> let op_prefix, op_infix1, op_infix2, op_suffix = Ops.ternop_c_syntax prec op in @@ -268,18 +277,23 @@ struct ^^ string " < 1.0f")) ^^ ifflat (space ^^ string "?" ^^ space ^^ v2 ^^ space ^^ string ":" ^^ space - ^^ string "(" ^^ string (typ_of_prec prec) ^^ string ")0") + ^^ string "(" + ^^ string (typ_of_prec prec) + ^^ string ")0") (nest 2 (break 1 ^^ string "?" ^^ space ^^ v2 ^^ break 1 ^^ string ":" ^^ space - ^^ string "(" ^^ string (typ_of_prec prec) ^^ string ")0")))) + ^^ string "(" + ^^ string (typ_of_prec prec) + ^^ string ")0")))) | Ops.Fp8_prec _ -> let open PPrint in group (parens (group (parens - (string "fp8_to_float(" ^^ v1 ^^ string ") > 0.0f && fp8_to_float(" - ^^ v1 ^^ string ") < 1.0f")) + (string "fp8_to_float(" ^^ v1 + ^^ string ") > 0.0f && fp8_to_float(" + ^^ v1 ^^ string ") < 1.0f")) ^^ ifflat (space ^^ string "?" ^^ space ^^ v2 ^^ space ^^ string ":" ^^ space ^^ string "float_to_fp8(0.0f)") @@ -292,8 +306,9 @@ struct (parens (group (parens - (string "bfloat16_to_float(" ^^ v1 ^^ string ") > 0.0f && bfloat16_to_float(" - ^^ v1 ^^ string ") < 1.0f")) + (string "bfloat16_to_float(" ^^ v1 + ^^ string ") > 0.0f && bfloat16_to_float(" + ^^ v1 ^^ string ") < 1.0f")) ^^ ifflat (space ^^ string "?" ^^ space ^^ v2 ^^ space ^^ string ":" ^^ space ^^ string "float_to_bfloat16(0.0f)") @@ -334,32 +349,37 @@ struct (break 1 ^^ string "?" ^^ space ^^ v2 ^^ break 1 ^^ string ":" ^^ space ^^ string "0.0")))) | Ops.Void_prec -> invalid_arg "Pure_C_config.binop_syntax: Satur01_gate on Void_prec") - | _ -> + | _ -> ( match prec with - | Ops.Bfloat16_prec _ -> + | Ops.Bfloat16_prec _ -> ( (* For BFloat16, perform all operations in float precision *) let float_v1 = PPrint.(string "bfloat16_to_float(" ^^ v1 ^^ string ")") in let float_v2 = PPrint.(string "bfloat16_to_float(" ^^ v2 ^^ string ")") in let op_prefix, op_infix, op_suffix = Ops.binop_c_syntax Ops.single op in - let float_result = PPrint.( - group (string op_prefix ^^ float_v1 ^^ string op_infix - ^^ ifflat (space ^^ float_v2) (nest 2 (break 1 ^^ float_v2)) - ^^ string op_suffix)) in + let float_result = + PPrint.( + group + (string op_prefix ^^ float_v1 ^^ string op_infix + ^^ ifflat (space ^^ float_v2) (nest 2 (break 1 ^^ float_v2)) + ^^ string op_suffix)) + in (* For comparison operations, return float result (0.0 or 1.0) converted to BFloat16 *) - (match op with + match op with | Ops.Cmplt | Ops.Cmpeq | Ops.Cmpne | Ops.Or | Ops.And -> - PPrint.(string "float_to_bfloat16(" ^^ float_result ^^ string ")") - | _ -> - PPrint.(string "float_to_bfloat16(" ^^ float_result ^^ string ")")) + PPrint.(string "float_to_bfloat16(" ^^ float_result ^^ string ")") + | _ -> PPrint.(string "float_to_bfloat16(" ^^ float_result ^^ string ")")) | Ops.Fp8_prec _ -> (* For FP8, perform all operations in float precision *) let float_v1 = PPrint.(string "fp8_to_float(" ^^ v1 ^^ string ")") in let float_v2 = PPrint.(string "fp8_to_float(" ^^ v2 ^^ string ")") in let op_prefix, op_infix, op_suffix = Ops.binop_c_syntax Ops.single op in - let float_result = PPrint.( - group (string op_prefix ^^ float_v1 ^^ string op_infix - ^^ ifflat (space ^^ float_v2) (nest 2 (break 1 ^^ float_v2)) - ^^ string op_suffix)) in + let float_result = + PPrint.( + group + (string op_prefix ^^ float_v1 ^^ string op_infix + ^^ ifflat (space ^^ float_v2) (nest 2 (break 1 ^^ float_v2)) + ^^ string op_suffix)) + in PPrint.(string "float_to_fp8(" ^^ float_result ^^ string ")") | _ -> let op_prefix, op_infix, op_suffix = Ops.binop_c_syntax prec op in @@ -367,7 +387,7 @@ struct group (string op_prefix ^^ v1 ^^ string op_infix ^^ ifflat (space ^^ v2) (nest 2 (break 1 ^^ v2)) - ^^ string op_suffix) + ^^ string op_suffix)) let unop_syntax prec op v = match prec with dune build @fmt failed "/usr/bin/env" "bash" "-c" "opam exec -- dune build @fmt --ignore-promoted-rules || (echo "dune build @fmt failed"; exit 2)" failed with exit status 2 2025-05-31 10:11.40: Job failed: Failed: Build failed