2025-05-30 12:20.44: New job: test ahrefs/ocannl https://github.com/ahrefs/ocannl.git#refs/heads/master (032faccf8dc533dfc393bed0a591a35ed988bb71) (linux-x86_64:(lint-fmt))Base: ocaml/opam:debian-12-ocaml-4.08@sha256:48fa4a7216c3973bb95572cf5dca98cbbcefe90f288f552e7ac70a8ccd438aa7ocamlformat 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 032faccfcat > Dockerfile <<'END-OF-DOCKERFILE'FROM ocaml/opam:debian-12-ocaml-4.08@sha256:48fa4a7216c3973bb95572cf5dca98cbbcefe90f288f552e7ac70a8ccd438aa7USER 1000:1000RUN 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 -uRUN opam depext -i duneWORKDIR /srcRUN opam depext -i ocamlformat=0.27.0COPY --chown=1000:1000 . /src/RUN opam exec -- dune build @fmt --ignore-promoted-rules || (echo "dune build @fmt failed"; exit 2)END-OF-DOCKERFILEdocker build .END-REPRO-BLOCK2025-05-30 12:20.44: 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-30 12:20.44: 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-30 12:20.44: Waiting for resource in pool OCluster2025-05-30 12:20.45: Waiting for worker…2025-05-30 12:41.32: Got resource from pool OClusterBuilding on asteria.caelum.ci.devAll commits already cachedHEAD is now at 032faccf Untested / broken: Claude's first pass at adding BF16, FP8, uint16, int32(from ocaml/opam:debian-12-ocaml-4.08@sha256:48fa4a7216c3973bb95572cf5dca98cbbcefe90f288f552e7ac70a8ccd438aa7)2025-05-30 12:51.00 ---> saved as "d1b97f3f32fc7cff4791d73e3fff398d19cc5b0541c709028ff05a921e22d2c8"/: (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] Initialiseddefault (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 environment2025-05-30 12:52.01 ---> saved as "0ed37ff9180ea5331ab17e3106ce3fc10bf21b69021c66107f159798cee036fc"/: (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 packagesThe 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.2Done.# Run eval $(opam env) to update the current shell environment2025-05-30 12:52.55 ---> saved as "dc3dc92082cbd6b1d2902dc06db6b555f449a0b1c6f054b5888f3ddab10a5d9e"/: (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 packagesThe 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.0Done.<><> 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 environment2025-05-30 12:53.56 ---> saved as "5e3a5282f655843df8e02a2497598f2918acc59345107286c51ff04d386b492b"/src: (copy (src .) (dst /src/))2025-05-30 12:53.57 ---> saved as "c375d33d40a1ccf514cac27fd3d008b3624d519ab7f8b6ea098973e44c2823f3"/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/testFile "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/testFile "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.mlindex b23f86d..1a6d615 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_precisionend@@ -210,71 +210,36 @@ let%track3_sexp link_compiled ~merge_buffer ~runner_label ctx_arrays (code : prodescription = "executes " ^ code.name ^ " on " ^ runner_label;work;} )-(*-let typ_of_prec = function- | Ops.Byte_prec _ -> "unsigned char"- | Ops.Uint16_prec _ -> "unsigned short"- | Ops.Int32_prec _ -> "int"- | Ops.Half_prec _ -> "_Float16"- | Ops.Bfloat16_prec _ -> "unsigned short" (* Stored as uint16, emulated as float *)- | Ops.Fp8_prec _ -> "unsigned char" (* Stored as uint8, emulated as float *)- | Ops.Single_prec _ -> "float"- | Ops.Double_prec _ -> "double"- | Ops.Void_prec -> "void"--(* Helper functions for bfloat16 and fp8 conversions *)-let extra_declarations =- [- "/* Emulation functions for special float types */";- "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);";- "}";- "";- "/* Simplified FP8 E5M2 format emulation */";- "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;";- " float result = (1.0f + mant * 0.25f) * powf(2.0f, (float)exp - 15.0f);";- " return sign ? -result : 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);";- " int exp = (int)floorf(log2f(f)) + 15;";- " if (exp < 0) return 0;";- " if (exp > 31) return sign ? 0xFF : 0x7F;";- " float mant = f / powf(2.0f, (float)exp - 15.0f) - 1.0f;";- " unsigned int mant_bits = (unsigned int)(mant * 4.0f + 0.5f);";- " if (mant_bits > 3) mant_bits = 3;";- " return (unsigned char)((sign << 7) | ((exp & 0x1F) << 2) | (mant_bits & 0x3));";- "}";- ]--let convert_precision ~from ~to_ =- match (from, to_) with- | p1, p2 when Ops.equal_prec p1 p2 -> ("", "")- | Ops.Bfloat16_prec _, Ops.Single_prec _ -> ("bfloat16_to_float(", ")")- | Ops.Bfloat16_prec _, Ops.Double_prec _ -> ("(double)bfloat16_to_float(", ")")- | Ops.Single_prec _, Ops.Bfloat16_prec _ -> ("float_to_bfloat16(", ")")- | Ops.Double_prec _, Ops.Bfloat16_prec _ -> ("float_to_bfloat16((float)", ")")- | Ops.Fp8_prec _, Ops.Single_prec _ -> ("fp8_to_float(", ")")- | Ops.Fp8_prec _, Ops.Double_prec _ -> ("(double)fp8_to_float(", ")")- | Ops.Single_prec _, Ops.Fp8_prec _ -> ("float_to_fp8(", ")")- | Ops.Double_prec _, Ops.Fp8_prec _ -> ("float_to_fp8((float)", ")")- | Ops.Bfloat16_prec _, _ -> ("(float)bfloat16_to_float(", ")") (* Convert via float *)- | _, Ops.Bfloat16_prec _ -> ("float_to_bfloat16((float)", ")")- | Ops.Fp8_prec _, _ -> ("(float)fp8_to_float(", ")") (* Convert via float *)- | _, Ops.Fp8_prec _ -> ("float_to_fp8((float)", ")")- | _ -> Ops.c_convert_precision ~from ~to_-*)\ No newline at end of file+(* let typ_of_prec = function | Ops.Byte_prec _ -> "unsigned char" | Ops.Uint16_prec _ -> "unsigned+ short" | Ops.Int32_prec _ -> "int" | Ops.Half_prec _ -> "_Float16" | Ops.Bfloat16_prec _ ->+ "unsigned short" (* Stored as uint16, emulated as float *) | Ops.Fp8_prec _ -> "unsigned char" (*+ Stored as uint8, emulated as float *) | Ops.Single_prec _ -> "float" | Ops.Double_prec _ ->+ "double" | Ops.Void_prec -> "void"++ (* Helper functions for bfloat16 and fp8 conversions *) let extra_declarations = [ "/* Emulation+ functions for special float types */"; "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);"; "}"; ""; "/* Simplified FP8 E5M2 format emulation */"; "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;"; "+ float result = (1.0f + mant * 0.25f) * powf(2.0f, (float)exp - 15.0f);"; " return sign ? -result+ : 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);"; " int exp =+ (int)floorf(log2f(f)) + 15;"; " if (exp < 0) return 0;"; " if (exp > 31) return sign ? 0xFF :+ 0x7F;"; " float mant = f / powf(2.0f, (float)exp - 15.0f) - 1.0f;"; " unsigned int mant_bits =+ (unsigned int)(mant * 4.0f + 0.5f);"; " if (mant_bits > 3) mant_bits = 3;"; " return (unsigned+ char)((sign << 7) | ((exp & 0x1F) << 2) | (mant_bits & 0x3));"; "}"; ]++ let convert_precision ~from ~to_ = match (from, to_) with | p1, p2 when Ops.equal_prec p1 p2 ->+ ("", "") | Ops.Bfloat16_prec _, Ops.Single_prec _ -> ("bfloat16_to_float(", ")") |+ Ops.Bfloat16_prec _, Ops.Double_prec _ -> ("(double)bfloat16_to_float(", ")") | Ops.Single_prec+ _, Ops.Bfloat16_prec _ -> ("float_to_bfloat16(", ")") | Ops.Double_prec _, Ops.Bfloat16_prec _ ->+ ("float_to_bfloat16((float)", ")") | Ops.Fp8_prec _, Ops.Single_prec _ -> ("fp8_to_float(", ")")+ | Ops.Fp8_prec _, Ops.Double_prec _ -> ("(double)fp8_to_float(", ")") | Ops.Single_prec _,+ Ops.Fp8_prec _ -> ("float_to_fp8(", ")") | Ops.Double_prec _, Ops.Fp8_prec _ ->+ ("float_to_fp8((float)", ")") | Ops.Bfloat16_prec _, _ -> ("(float)bfloat16_to_float(", ")") (*+ Convert via float *) | _, Ops.Bfloat16_prec _ -> ("float_to_bfloat16((float)", ")") |+ Ops.Fp8_prec _, _ -> ("(float)fp8_to_float(", ")") (* Convert via float *) | _, Ops.Fp8_prec _ ->+ ("float_to_fp8((float)", ")") | _ -> Ops.c_convert_precision ~from ~to_ *)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.mlindex 78f368f..0959936 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 Bfloat16let fp8 = Fp8_prec Fp8let single = Single_prec Singlelet 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+ | _ -> falselet 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 " ^ slet 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@@ -101,7 +104,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 _, _ ->+ falselet prec_in_bytes = function| Void_prec -> 0@@ -135,14 +147,14 @@ let promote_prec p1 p2 =| Void_prec, Void_prec -> Void_preclet 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 -> doubletype 'r map_prec = { f : 'ocaml 'elt_t. ('ocaml, 'elt_t) precision -> 'r }@@ -159,15 +171,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"@@ -384,14 +396,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, _ -> ("(", " <", ")")@@ -460,7 +475,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"+ inmatch (op, prec) with| Identity, _ -> ("", "")| Relu, (Byte_prec _ | Uint16_prec _ | Int32_prec _ | Fp8_prec _) -> ("fmax(0, ", ")")@@ -499,9 +518,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_ =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.mlindex 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.mlindex 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 = structfun 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 ->groupFile "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.mlindex 89e3f46..dcb44eb 100644--- a/_build/default/arrayjit/lib/c_syntax.ml+++ b/_build/default/arrayjit/lib/.formatted/c_syntax.ml@@ -185,10 +185,14 @@ 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.Bfloat16_prec _ ->(* For CC backend, convert to float for computation *)let open PPrint inFile "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.mlindex 8bcfc75..a894c29 100644--- a/_build/default/arrayjit/lib/ndarray.ml+++ b/_build/default/arrayjit/lib/.formatted/ndarray.ml@@ -25,8 +25,8 @@ type byte_nd = (char, Ops.uint8_elt) bigarraytype 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) bigarraytype double_nd = (float, Ops.float64_elt) bigarray@@ -39,14 +39,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]@@ -69,8 +69,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@@ -153,18 +153,27 @@ 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 Int.of_float values strict (* TODO: proper bfloat16 conversion *)+ | Ops.Bfloat16, Constant_fill { values; strict } ->+ constant_fill_f Int.of_float values strict (* TODO: proper bfloat16 conversion *)| Ops.Bfloat16, Range_over_offsets ->- init_bigarray_of_prec prec dims ~f:(fun idcs -> indices_to_offset ~dims ~idcs) (* TODO: proper bfloat16 conversion *)- | Ops.Bfloat16, Standard_uniform -> init_bigarray_of_prec prec dims ~f:(fun _ -> Random.int 65536) (* TODO: proper bfloat16 conversion *)- | Ops.Fp8, Constant_fill { values; strict } -> constant_fill_f (Fn.compose Char.of_int_exn Int.of_float) values strict (* TODO: proper fp8 conversion *)+ init_bigarray_of_prec prec dims ~f:(fun idcs -> indices_to_offset ~dims ~idcs)+ (* TODO: proper bfloat16 conversion *)+ | Ops.Bfloat16, Standard_uniform ->+ init_bigarray_of_prec prec dims ~f:(fun _ -> Random.int 65536)+ (* TODO: proper bfloat16 conversion *)+ | Ops.Fp8, Constant_fill { values; strict } ->+ constant_fill_f+ (Fn.compose Char.of_int_exn Int.of_float)+ values strict (* TODO: proper fp8 conversion *)| Ops.Fp8, Range_over_offsets ->init_bigarray_of_prec prec dims ~f:(fun idcs ->Char.of_int_exn @@ indices_to_offset ~dims ~idcs)@@ -255,8 +264,9 @@ let set_from_float arr idx v =| Uint16_nd arr -> A.set arr idx @@ Int.of_float v| Int32_nd arr -> A.set arr idx @@ Int32.of_float v| Half_nd arr -> A.set arr idx v- | Bfloat16_nd arr -> A.set arr idx @@ Int.of_float v (* TODO: proper bfloat16 conversion *)- | Fp8_nd arr -> A.set arr idx @@ Char.of_int_exn @@ Int.of_float v (* TODO: proper fp8 conversion *)+ | Bfloat16_nd arr -> A.set arr idx @@ Int.of_float v (* TODO: proper bfloat16 conversion *)+ | Fp8_nd arr ->+ A.set arr idx @@ Char.of_int_exn @@ Int.of_float v (* TODO: proper fp8 conversion *)| Single_nd arr -> A.set arr idx v| Double_nd arr -> A.set arr idx v@@ -266,8 +276,8 @@ let fill_from_float arr v =| Uint16_nd arr -> A.fill arr @@ Int.of_float v| Int32_nd arr -> A.fill arr @@ Int32.of_float v| Half_nd arr -> A.fill arr v- | Bfloat16_nd arr -> A.fill arr @@ Int.of_float v (* TODO: proper bfloat16 conversion *)- | Fp8_nd arr -> A.fill arr @@ Char.of_int_exn @@ Int.of_float v (* TODO: proper fp8 conversion *)+ | Bfloat16_nd arr -> A.fill arr @@ Int.of_float v (* TODO: proper bfloat16 conversion *)+ | Fp8_nd arr -> A.fill arr @@ Char.of_int_exn @@ Int.of_float v (* TODO: proper fp8 conversion *)| Single_nd arr -> A.fill arr v| Double_nd arr -> A.fill arr v@@ -310,7 +320,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)@@ -319,11 +329,17 @@ let reset_bigarray (init_op : Ops.init_op) (type o b) (prec : (o, b) Ops.precisi| Ops.Half, Range_over_offsets ->set_bigarray arr ~f:(fun idcs -> Float.of_int @@ indices_to_offset ~dims ~idcs)| 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 Int.of_float values strict (* TODO: proper bfloat16 conversion *)+ | Ops.Bfloat16, Constant_fill { values; strict } ->+ constant_set_f Int.of_float values strict (* TODO: proper bfloat16 conversion *)| Ops.Bfloat16, Range_over_offsets ->- set_bigarray arr ~f:(fun idcs -> indices_to_offset ~dims ~idcs) (* TODO: proper bfloat16 conversion *)- | Ops.Bfloat16, Standard_uniform -> set_bigarray arr ~f:(fun _ -> Random.int 65536) (* TODO: proper bfloat16 conversion *)- | Ops.Fp8, Constant_fill { values; strict } -> constant_set_f (Fn.compose Char.of_int_exn Int.of_float) values strict (* TODO: proper fp8 conversion *)+ set_bigarray arr ~f:(fun idcs -> indices_to_offset ~dims ~idcs)+ (* TODO: proper bfloat16 conversion *)+ | Ops.Bfloat16, Standard_uniform ->+ set_bigarray arr ~f:(fun _ -> Random.int 65536) (* TODO: proper bfloat16 conversion *)+ | Ops.Fp8, Constant_fill { values; strict } ->+ constant_set_f+ (Fn.compose Char.of_int_exn Int.of_float)+ values strict (* TODO: proper fp8 conversion *)| Ops.Fp8, Range_over_offsets ->set_bigarray arr ~f:(fun idcs -> Char.of_int_exn @@ indices_to_offset ~dims ~idcs)| Ops.Fp8, Standard_uniform -> set_bigarray arr ~f:(fun _ -> Rand.Lib.char ())@@ -363,8 +379,14 @@ 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 @@ Float.of_int v) arr (* TODO: proper bfloat16 conversion *)- | Fp8_nd arr -> fold_bigarray ~init ~f:(fun accu idx c -> f accu idx @@ Float.of_int @@ Char.to_int c) arr (* TODO: proper fp8 conversion *)+ | Bfloat16_nd arr ->+ fold_bigarray ~init+ ~f:(fun accu idx v -> f accu idx @@ Float.of_int v)+ arr (* TODO: proper bfloat16 conversion *)+ | Fp8_nd arr ->+ fold_bigarray ~init+ ~f:(fun accu idx c -> f accu idx @@ Float.of_int @@ Char.to_int c)+ arr (* TODO: proper fp8 conversion *)| Single_nd arr -> fold_bigarray ~init ~f arr| Double_nd arr -> fold_bigarray ~init ~f arr@@ -380,8 +402,8 @@ let get_as_float arr idx =| Uint16_nd arr -> Float.of_int @@ A.get arr idx| Int32_nd arr -> Int32.to_float @@ A.get arr idx| Half_nd arr -> A.get arr idx- | Bfloat16_nd arr -> Float.of_int @@ A.get arr idx (* TODO: proper bfloat16 conversion *)- | Fp8_nd arr -> Float.of_int @@ Char.to_int @@ A.get arr idx (* TODO: proper fp8 conversion *)+ | Bfloat16_nd arr -> Float.of_int @@ A.get arr idx (* TODO: proper bfloat16 conversion *)+ | Fp8_nd arr -> Float.of_int @@ Char.to_int @@ A.get arr idx (* TODO: proper fp8 conversion *)| Single_nd arr -> A.get arr idx| Double_nd arr -> A.get arr idxdune 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 22025-05-30 12:53.58: Job failed: Failed: Build failed