diff --git a/Documentation/Benchmarks/data/jp3d-bench-arm64-v10_20-batched-20260524.json b/Documentation/Benchmarks/data/jp3d-bench-arm64-v10_20-batched-20260524.json new file mode 100644 index 00000000..45d48d8f --- /dev/null +++ b/Documentation/Benchmarks/data/jp3d-bench-arm64-v10_20-batched-20260524.json @@ -0,0 +1,487 @@ +{ + "decode" : { + "ct_3d_large" : { + "depth" : 64, + "fixture" : "ct_3d_large", + "height" : 512, + "label" : "CT 512×512×64", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 709.06387500000005, + "median" : 674.58741699999996, + "min" : 670.24775, + "samples" : [ + 670.53800000000001, + 709.06387500000005, + 688.24216699999999, + 670.24775, + 672.72829200000001, + 674.58741699999996, + 698.262834 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 323.58737500000001, + "median" : 268.855208, + "min" : 252.31375, + "samples" : [ + 268.855208, + 253.137833, + 252.31375, + 274.96629200000001, + 264.834541, + 323.58737500000001, + 279.12354199999999 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 226.69608299999999, + "median" : 218.6645, + "min" : 211.15691699999999, + "samples" : [ + 214.38983300000001, + 211.15691699999999, + 218.6645, + 215.1405, + 226.69608299999999, + 222.48949999999999, + 221.68858299999999 + ] + } + }, + "source" : "synthetic", + "width" : 512 + }, + "ct_3d_mid" : { + "depth" : 32, + "fixture" : "ct_3d_mid", + "height" : 512, + "label" : "CT 512×512×32", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 320.89125000000001, + "median" : 316.55920800000001, + "min" : 313.95249999999999, + "samples" : [ + 313.95249999999999, + 314.16387500000002, + 316.55920800000001, + 317.06037500000002, + 320.89125000000001, + 316.89066700000001, + 314.78454199999999 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 122.14254200000001, + "median" : 118.88033299999999, + "min" : 117.654167, + "samples" : [ + 117.654167, + 119.318625, + 118.445125, + 118.88033299999999, + 119.74654200000001, + 122.14254200000001, + 117.656334 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 95.369457999999995, + "median" : 91.877291999999997, + "min" : 90.178792000000001, + "samples" : [ + 90.178792000000001, + 91.877291999999997, + 95.369457999999995, + 93.118750000000006, + 92.893665999999996, + 91.446167000000003, + 90.784834000000004 + ] + } + }, + "source" : "synthetic", + "width" : 512 + }, + "ct_3d_small" : { + "depth" : 16, + "fixture" : "ct_3d_small", + "height" : 256, + "label" : "CT 256×256×16", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 39.637374999999999, + "median" : 38.990250000000003, + "min" : 38.902707999999997, + "samples" : [ + 39.637374999999999, + 38.984999999999999, + 38.990250000000003, + 38.940750000000001, + 39.109917000000003, + 39.342291000000003, + 38.902707999999997 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 14.483791, + "median" : 13.850792, + "min" : 13.658291, + "samples" : [ + 13.850792, + 14.217167, + 13.753625, + 13.805624999999999, + 13.658291, + 14.483791, + 14.10375 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 11.071541, + "median" : 11.013042, + "min" : 10.885583, + "samples" : [ + 10.885583, + 10.916416, + 11.041957999999999, + 11.071541, + 11.068584, + 10.914583, + 11.013042 + ] + } + }, + "source" : "synthetic", + "width" : 256 + }, + "mr_3d_mid" : { + "depth" : 32, + "fixture" : "mr_3d_mid", + "height" : 256, + "label" : "MR 256×256×32", + "modality" : "MR", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 78.013625000000005, + "median" : 77.708917, + "min" : 77.262124999999997, + "samples" : [ + 78.013625000000005, + 77.451583999999997, + 77.262124999999997, + 77.708917, + 77.465958000000001, + 77.993042000000003, + 77.786417 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 30.66075, + "median" : 29.849499999999999, + "min" : 29.167375, + "samples" : [ + 29.221625, + 29.167375, + 29.518708, + 30.66075, + 29.941792, + 30.000125000000001, + 29.849499999999999 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 23.029667, + "median" : 22.076875000000001, + "min" : 21.627834, + "samples" : [ + 22.076875000000001, + 21.627834, + 22.690625000000001, + 22.051917, + 21.636375000000001, + 23.029667, + 22.348875 + ] + } + }, + "source" : "synthetic", + "width" : 256 + }, + "mr_3d_small" : { + "depth" : 16, + "fixture" : "mr_3d_small", + "height" : 128, + "label" : "MR 128×128×16", + "modality" : "MR", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 14.377542, + "median" : 14.045500000000001, + "min" : 13.598958, + "samples" : [ + 14.377542, + 14.062457999999999, + 14.045500000000001, + 13.598958, + 13.772416, + 14.112125000000001, + 13.89175 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 5.6600000000000001, + "median" : 5.621791, + "min" : 5.5273329999999996, + "samples" : [ + 5.6321250000000003, + 5.6566669999999997, + 5.621791, + 5.5877499999999998, + 5.5273329999999996, + 5.5783339999999999, + 5.6600000000000001 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 4.3739999999999997, + "median" : 4.2757079999999998, + "min" : 4.2255000000000003, + "samples" : [ + 4.2757079999999998, + 4.3739999999999997, + 4.2255000000000003, + 4.2744999999999997, + 4.2444579999999998, + 4.3028750000000002, + 4.286708 + ] + } + }, + "source" : "synthetic", + "width" : 128 + }, + "us_3d_small" : { + "depth" : 24, + "fixture" : "us_3d_small", + "height" : 240, + "label" : "US 320×240×24", + "modality" : "US", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 56.833750000000002, + "median" : 56.285874999999997, + "min" : 55.919541000000002, + "samples" : [ + 55.919541000000002, + 56.285874999999997, + 56.443041999999998, + 56.804250000000003, + 56.113709, + 56.034292000000001, + 56.833750000000002 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 22.990791999999999, + "median" : 22.413958999999998, + "min" : 22.179583000000001, + "samples" : [ + 22.179583000000001, + 22.990791999999999, + 22.58625, + 22.338540999999999, + 22.80575, + 22.413958999999998, + 22.285625 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 14.640833000000001, + "median" : 14.223750000000001, + "min" : 13.898833, + "samples" : [ + 14.262708999999999, + 14.640833000000001, + 14.134957999999999, + 14.109208000000001, + 14.223750000000001, + 13.898833, + 14.304665999999999 + ] + } + }, + "source" : "synthetic", + "width" : 320 + } + }, + "encode" : { + "ct_3d_large" : { + "depth" : 64, + "fixture" : "ct_3d_large", + "height" : 512, + "label" : "CT 512×512×64", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 599.76829199999997, + "median" : 573.73458300000004, + "min" : 552.43608300000005, + "samples" : [ + 552.43608300000005, + 554.52820899999995, + 556.86291700000004, + 588.32837500000005, + 573.73458300000004, + 599.76829199999997, + 584.22170800000004 + ] + } + }, + "source" : "synthetic", + "width" : 512 + }, + "ct_3d_mid" : { + "depth" : 32, + "fixture" : "ct_3d_mid", + "height" : 512, + "label" : "CT 512×512×32", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 264.73225000000002, + "median" : 262.52354200000002, + "min" : 260.23866600000002, + "samples" : [ + 263.86174999999997, + 262.76191699999998, + 261.415708, + 264.73225000000002, + 262.52354200000002, + 260.23866600000002, + 261.60975000000002 + ] + } + }, + "source" : "synthetic", + "width" : 512 + }, + "ct_3d_small" : { + "depth" : 16, + "fixture" : "ct_3d_small", + "height" : 256, + "label" : "CT 256×256×16", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 32.506749999999997, + "median" : 32.259458000000002, + "min" : 31.888209, + "samples" : [ + 32.061666000000002, + 32.506749999999997, + 32.118667000000002, + 32.385458, + 32.394917, + 31.888209, + 32.259458000000002 + ] + } + }, + "source" : "synthetic", + "width" : 256 + }, + "mr_3d_mid" : { + "depth" : 32, + "fixture" : "mr_3d_mid", + "height" : 256, + "label" : "MR 256×256×32", + "modality" : "MR", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 69.323958000000005, + "median" : 65.196250000000006, + "min" : 64.697125, + "samples" : [ + 65.196250000000006, + 64.959165999999996, + 65.171666999999999, + 66.492958000000002, + 66.847333000000006, + 64.697125, + 69.323958000000005 + ] + } + }, + "source" : "synthetic", + "width" : 256 + }, + "mr_3d_small" : { + "depth" : 16, + "fixture" : "mr_3d_small", + "height" : 128, + "label" : "MR 128×128×16", + "modality" : "MR", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 11.889709, + "median" : 11.700792, + "min" : 11.555208, + "samples" : [ + 11.764792, + 11.747417, + 11.699458, + 11.623749999999999, + 11.700792, + 11.555208, + 11.889709 + ] + } + }, + "source" : "synthetic", + "width" : 128 + }, + "us_3d_small" : { + "depth" : 24, + "fixture" : "us_3d_small", + "height" : 240, + "label" : "US 320×240×24", + "modality" : "US", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 56.72175, + "median" : 54.566291999999997, + "min" : 54.284041999999999, + "samples" : [ + 54.470332999999997, + 54.284041999999999, + 54.566291999999997, + 54.303916999999998, + 55.352584, + 56.02525, + 56.72175 + ] + } + }, + "source" : "synthetic", + "width" : 320 + } + }, + "host" : { + "brand" : "Apple M2", + "machine" : "arm64", + "model" : "Mac14,2", + "ncpu" : 8, + "release" : "15.7.5", + "system" : "Darwin" + }, + "j2k_version" : "J2KSwift version 10.9.2", + "lane" : "inproc-jp3d", + "runs" : 7, + "schema" : "warm-inproc-jp3d-v1", + "warmups" : 2 +} \ No newline at end of file diff --git a/Documentation/Benchmarks/data/jp3d-bench-arm64-v10_20-serial-20260524.json b/Documentation/Benchmarks/data/jp3d-bench-arm64-v10_20-serial-20260524.json new file mode 100644 index 00000000..d8eb1147 --- /dev/null +++ b/Documentation/Benchmarks/data/jp3d-bench-arm64-v10_20-serial-20260524.json @@ -0,0 +1,487 @@ +{ + "decode" : { + "ct_3d_large" : { + "depth" : 64, + "fixture" : "ct_3d_large", + "height" : 512, + "label" : "CT 512×512×64", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 796.25229200000001, + "median" : 789.15395899999999, + "min" : 761.49591699999996, + "samples" : [ + 762.63049999999998, + 761.49591699999996, + 794.83841600000005, + 778.27358300000003, + 794.82616700000005, + 789.15395899999999, + 796.25229200000001 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 272.556333, + "median" : 257.81670800000001, + "min" : 244.703667, + "samples" : [ + 244.703667, + 262.68941699999999, + 253.191667, + 257.81670800000001, + 244.820584, + 258.67670900000002, + 272.556333 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 238.899, + "median" : 220.255458, + "min" : 206.56299999999999, + "samples" : [ + 207.594875, + 206.56299999999999, + 225.11658299999999, + 238.899, + 235.87258299999999, + 212.01362499999999, + 220.255458 + ] + } + }, + "source" : "synthetic", + "width" : 512 + }, + "ct_3d_mid" : { + "depth" : 32, + "fixture" : "ct_3d_mid", + "height" : 512, + "label" : "CT 512×512×32", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 400.155958, + "median" : 369.812209, + "min" : 366.71237500000001, + "samples" : [ + 400.155958, + 386.912125, + 376.68025, + 369.812209, + 367.57745799999998, + 366.71237500000001, + 368.37241699999998 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 136.19450000000001, + "median" : 127.211792, + "min" : 117.94954199999999, + "samples" : [ + 127.211792, + 117.94954199999999, + 128.56133299999999, + 131.303167, + 136.19450000000001, + 123.112584, + 126.788167 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 101.759208, + "median" : 100.547416, + "min" : 96.618125000000006, + "samples" : [ + 96.618125000000006, + 101.759208, + 100.367166, + 100.865583, + 101.69125, + 99.292292000000003, + 100.547416 + ] + } + }, + "source" : "synthetic", + "width" : 512 + }, + "ct_3d_small" : { + "depth" : 16, + "fixture" : "ct_3d_small", + "height" : 256, + "label" : "CT 256×256×16", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 44.536959000000003, + "median" : 44.108125000000001, + "min" : 43.163583000000003, + "samples" : [ + 44.338500000000003, + 44.536959000000003, + 44.108125000000001, + 44.412582999999998, + 43.387416999999999, + 43.987124999999999, + 43.163583000000003 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 14.4465, + "median" : 14.27875, + "min" : 14.051167, + "samples" : [ + 14.212, + 14.368292, + 14.206167000000001, + 14.4465, + 14.051167, + 14.27875, + 14.316167 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 11.370457999999999, + "median" : 11.158917000000001, + "min" : 10.936667, + "samples" : [ + 11.303167, + 11.303958, + 10.936667, + 10.955500000000001, + 11.158917000000001, + 11.370457999999999, + 11.088666 + ] + } + }, + "source" : "synthetic", + "width" : 256 + }, + "mr_3d_mid" : { + "depth" : 32, + "fixture" : "mr_3d_mid", + "height" : 256, + "label" : "MR 256×256×32", + "modality" : "MR", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 89.134957999999997, + "median" : 87.209333999999998, + "min" : 87.014624999999995, + "samples" : [ + 89.134957999999997, + 87.113833, + 87.014624999999995, + 87.277083000000005, + 87.209333999999998, + 87.205667000000005, + 88.838915999999998 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 30.829084000000002, + "median" : 29.864583, + "min" : 29.163708, + "samples" : [ + 30.829084000000002, + 29.163708, + 29.864583, + 30.027791000000001, + 29.345082999999999, + 29.877084, + 29.259584 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 23.361000000000001, + "median" : 22.932791999999999, + "min" : 22.673666999999998, + "samples" : [ + 23.031957999999999, + 22.673666999999998, + 22.711749999999999, + 23.361000000000001, + 22.932791999999999, + 22.912334000000001, + 23.2745 + ] + } + }, + "source" : "synthetic", + "width" : 256 + }, + "mr_3d_small" : { + "depth" : 16, + "fixture" : "mr_3d_small", + "height" : 128, + "label" : "MR 128×128×16", + "modality" : "MR", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 13.995875, + "median" : 13.805875, + "min" : 13.598583, + "samples" : [ + 13.775582999999999, + 13.912832999999999, + 13.977166, + 13.805875, + 13.718541999999999, + 13.598583, + 13.995875 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 5.8488749999999996, + "median" : 5.7709169999999999, + "min" : 5.5492080000000001, + "samples" : [ + 5.8488749999999996, + 5.7837079999999998, + 5.7210419999999997, + 5.7896669999999997, + 5.7536250000000004, + 5.5492080000000001, + 5.7709169999999999 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 4.3439160000000001, + "median" : 4.2982079999999998, + "min" : 4.272875, + "samples" : [ + 4.3439160000000001, + 4.3066250000000004, + 4.2886670000000002, + 4.272875, + 4.3169579999999996, + 4.2982079999999998, + 4.2918339999999997 + ] + } + }, + "source" : "synthetic", + "width" : 128 + }, + "us_3d_small" : { + "depth" : 24, + "fixture" : "us_3d_small", + "height" : 240, + "label" : "US 320×240×24", + "modality" : "US", + "results" : { + "J2KSwift+jp3d+full" : { + "max" : 60.831375000000001, + "median" : 60.438124999999999, + "min" : 59.698917000000002, + "samples" : [ + 59.698917000000002, + 60.438124999999999, + 60.115209, + 60.443916999999999, + 60.831375000000001, + 60.222957999999998, + 60.735042 + ] + }, + "J2KSwift+jp3d+res1" : { + "max" : 24.859666000000001, + "median" : 23.265540999999999, + "min" : 23.009374999999999, + "samples" : [ + 23.265540999999999, + 23.076042000000001, + 24.859666000000001, + 24.099208000000001, + 23.021833999999998, + 23.339542000000002, + 23.009374999999999 + ] + }, + "J2KSwift+jp3d+roiq" : { + "max" : 15.041084, + "median" : 14.699875, + "min" : 14.407624999999999, + "samples" : [ + 14.712208, + 14.501625000000001, + 14.444333, + 14.407624999999999, + 14.699875, + 15.041084, + 14.850375 + ] + } + }, + "source" : "synthetic", + "width" : 320 + } + }, + "encode" : { + "ct_3d_large" : { + "depth" : 64, + "fixture" : "ct_3d_large", + "height" : 512, + "label" : "CT 512×512×64", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 611.46512499999994, + "median" : 589.50337500000001, + "min" : 587.35637499999996, + "samples" : [ + 611.46512499999994, + 589.50337500000001, + 592.61029199999996, + 587.94016699999997, + 587.95512499999995, + 590.81587500000001, + 587.35637499999996 + ] + } + }, + "source" : "synthetic", + "width" : 512 + }, + "ct_3d_mid" : { + "depth" : 32, + "fixture" : "ct_3d_mid", + "height" : 512, + "label" : "CT 512×512×32", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 285.94879100000003, + "median" : 276.57433300000002, + "min" : 265.74958299999997, + "samples" : [ + 265.74958299999997, + 273.97025000000002, + 276.93116600000002, + 276.31516699999997, + 276.57433300000002, + 282.22483299999999, + 285.94879100000003 + ] + } + }, + "source" : "synthetic", + "width" : 512 + }, + "ct_3d_small" : { + "depth" : 16, + "fixture" : "ct_3d_small", + "height" : 256, + "label" : "CT 256×256×16", + "modality" : "CT", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 32.478250000000003, + "median" : 31.890625, + "min" : 31.409165999999999, + "samples" : [ + 31.890625, + 32.115333, + 31.409165999999999, + 31.755917, + 31.455041999999999, + 32.311791999999997, + 32.478250000000003 + ] + } + }, + "source" : "synthetic", + "width" : 256 + }, + "mr_3d_mid" : { + "depth" : 32, + "fixture" : "mr_3d_mid", + "height" : 256, + "label" : "MR 256×256×32", + "modality" : "MR", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 65.614333999999999, + "median" : 65.094125000000005, + "min" : 63.6935, + "samples" : [ + 65.094125000000005, + 65.139042000000003, + 65.409666000000001, + 65.614333999999999, + 63.6935, + 64.005125000000007, + 65.093000000000004 + ] + } + }, + "source" : "synthetic", + "width" : 256 + }, + "mr_3d_small" : { + "depth" : 16, + "fixture" : "mr_3d_small", + "height" : 128, + "label" : "MR 128×128×16", + "modality" : "MR", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 12.431084, + "median" : 11.765375000000001, + "min" : 11.59125, + "samples" : [ + 12.431084, + 11.59125, + 11.866709, + 11.945625, + 11.634957999999999, + 11.674625000000001, + 11.765375000000001 + ] + } + }, + "source" : "synthetic", + "width" : 128 + }, + "us_3d_small" : { + "depth" : 24, + "fixture" : "us_3d_small", + "height" : 240, + "label" : "US 320×240×24", + "modality" : "US", + "results" : { + "J2KSwift+jp3d+inproc" : { + "max" : 55.480791000000004, + "median" : 54.664875000000002, + "min" : 54.387, + "samples" : [ + 54.613624999999999, + 54.387, + 54.770625000000003, + 54.609042000000002, + 54.664875000000002, + 54.992832999999997, + 55.480791000000004 + ] + } + }, + "source" : "synthetic", + "width" : 320 + } + }, + "host" : { + "brand" : "Apple M2", + "machine" : "arm64", + "model" : "Mac14,2", + "ncpu" : 8, + "release" : "15.7.5", + "system" : "Darwin" + }, + "j2k_version" : "J2KSwift version 10.9.2", + "lane" : "inproc-jp3d", + "runs" : 7, + "schema" : "warm-inproc-jp3d-v1", + "warmups" : 2 +} \ No newline at end of file diff --git a/README.md b/README.md index 04288c64..cca952b6 100644 --- a/README.md +++ b/README.md @@ -8,13 +8,15 @@ A pure Swift 6.2 implementation of JPEG 2000 (ISO/IEC 15444) encoding and decoding with strict concurrency support. -**Current Version**: 10.10.0 -**Status**: Apple Silicon-first JPEG 2000 / HTJ2K (Part-15) implementation. v10.10.0 ships **JP3D true partial-resolution + ROI footprint-skip**: `JP3DDecoderConfiguration.resolutionLevel` becomes functional (was wired but ignored — 2.2–3.1× faster decode at `resolutionLevel = 1` on M2 release) and `JP3DROIDecoder` replaces decode-then-crop with true per-tile XY footprint-skip + Z-narrow non-residual scan (ROI 1/4 is 3.4–4.1× faster than full decode on M2 release). Decoder-only; codestream bytes byte-identical to v10.9.3. -**Previous Release**: 10.9.3 (always-URL CompressionFamily dependency, fixes #438) +**Current Version**: 10.11.0 +**Status**: Apple Silicon-first JPEG 2000 / HTJ2K (Part-15) implementation. v10.11.0 ships **JP3D batched GPU iDWT**: a new single-dispatch Metal kernel amortises per-slice GPU overhead across the whole volume. `JP3DSliceStackCodec.decode` now runs ONE batched iDWT dispatch across all slices of a tile instead of N per-slice dispatches. M2 release JP3D full decode wins **−5 ms on small CT to −115 ms on 16M-voxel CT** (1.07–1.17× faster, 5/6 fixtures cross the 3 ms acceptance threshold). Decoder-only; codestream bytes byte-identical to v10.10.0. +**Previous Release**: 10.10.0 (JP3D true partial-resolution + ROI footprint-skip + Z-narrow) **Release process**: see [RELEASING.md](RELEASING.md). Every release MUST update this README (Current Version line + new Release Status paragraph) — see the Release artefacts checklist for the full requirements. ## 📦 Release Status +**v10.11.0** ships **JP3D batched GPU iDWT** — the production landing of a multi-week research arc that splits the JP3D per-tile decode pipeline at the dequant↔iDWT boundary and submits one batched Metal dispatch across the whole z-range instead of N per-slice dispatches. A new opaque-payload bridge SPI on `J2KDecoder` (`_jp3dDecodeToCoefficients` / `_jp3dIDWTAndFinalize` / `_jp3dIDWTAndFinalizeBatched`) is the architectural surface; two new Metal kernels (`j2k_dwt_inverse_53_horizontal_int_tiled_batched`, `..._vertical_..._batched`) extend the v10.3 tiled threadgroup-memory kernels with a Z-dim grid axis, processing N slices in one dispatch. `JP3DSliceStackCodec` now runs **parallel `_jp3dDecodeToCoefficients` across `[zStart, zUpper)` via a TaskGroup, then ONE batched iDWT call** before the existing sequential Z-delta residual chain. Per-slice GPU dispatch overhead amortises across the volume; the gain scales with slice count × per-slice work. **M2 release, J2KBenchMac --jp3d, in-process, 7 timed runs / 2 warmups**: ct_3d_small **−5.12 ms** (1.13×), us_3d_small −4.15 ms (1.07×), mr_3d_mid **−9.50 ms** (1.12×), ct_3d_mid **−53.25 ms** (1.17×), **ct_3d_large 16M-voxel CT −114.57 ms** (1.17×). 5/6 fixtures clear the 3 ms acceptance threshold; the wash (mr_3d_small @ 13 ms wall) is the smallest fixture where per-slice overhead dominates anyway. Kernel-level bench (16 slices × 256×256 × 3 levels): **5.93 → 2.06 ms = 2.4× faster GPU dispatch**. Eligibility gate: only `K=0 + no ROI` routes through the batched path; `K>0` keeps the per-slice `decodeResolution` loop and ROI keeps `decodeRegion` (bridging those is future work). Opt-out via `J2K_JP3D_BATCHED_BRIDGE=0`. **`V10_20_BatchedBridgeParityTests` 5/5 PASS, `V10_20_BatchedInverseInt32ParityTests` 12/12 PASS, `V10_20_JP3DBridgeParityTests` 5/5 PASS, full `swift test --filter JP3D` regression sweep 519/519 PASS, mandatory commit gate 7/7 PASS**. Codestream bytes byte-identical to v10.10.0; encoder unchanged. See [RELEASE_NOTES_v10.11.0.md](RELEASE_NOTES_v10.11.0.md) and the bench JSONs in [Documentation/Benchmarks/data/](Documentation/Benchmarks/data/). + **v10.10.0** ships **JP3D true partial-resolution + ROI footprint-skip + Z-narrow ROI skip** — three coordinated decoder-side changes inside `Sources/J2K3D/` that close the long-standing follow-up on JP3D's selective-decode story. **`JP3DDecoderConfiguration.resolutionLevel`** has been wired into the struct since v5.x but the decoder ignored it (silently returned full resolution); v10.10.0 routes each per-slice 2D codestream inside `JP3DSliceStackCodec` through the existing v10.5.0 `J2KDecoder.decodeResolution(_:options:)`, producing a volume sized `⌈W / 2^K⌉ × ⌈H / 2^K⌉ × D` as documented. **`JP3DROIDecoder`** swaps decode-then-crop for true ROI footprint-skip — the per-tile in-tile XY sub-region is passed through to the slice-stack codec, which routes the per-slice 2D decode through v10.6.0 `decodeRegion(_:options:)` with `.direct` strategy (code-blocks whose inverse-DWT cone-of-influence misses the region skip entropy decode entirely). **Z-narrow ROI skip** pre-scans slice headers (no decode), finds the latest non-residual slice ≤ `zRange.lowerBound`, and starts decoding from there — completely skipping the unused Z-prefix while keeping Z-delta residual chain integrity via the restart anchor. M2 release: **`resolutionLevel = 1` is 2.2–3.1× faster than full decode** (mr_3d_small 2.22×, ct_3d_small 3.05×, mr_3d_mid 3.01×); **ROI 1/4 is 3.4–4.1× faster than full decode** (mr_3d_small 3.37×, ct_3d_small 4.14×, mr_3d_mid 3.96×). Combined `resolutionLevel + ROI` throws loud (was silently ignored previously). `V10_18_TrueSelectiveParityTests` 9/9 PASS, existing `JP3DDecoderTests` 61/61 PASS, mandatory commit gate clean. Encoder unchanged. Codestream bytes byte-identical to v10.9.3. See [RELEASE_NOTES_v10.10.0.md](RELEASE_NOTES_v10.10.0.md) and [V10_18_JP3D_TRUE_SELECTIVE_DECODE.md](Documentation/research/V10_18_JP3D_TRUE_SELECTIVE_DECODE.md). **v10.9.3** is a packaging hotfix that makes J2KSwift unconditionally URL-consumable. v10.9.2 attempted CWD-conditional dependency resolution (sibling path if present, else URL) but `FileManager.fileExists` is evaluated with CWD set to the consuming root package — any consumer with a `CompressionFamily` directory beside its own root caused J2KSwift to fall back to the path form, and stable-version consumers may not transitively depend on path/branch packages ("unstable-version package" error). v10.9.3 drops the probe and always uses the public Git URL; local co-development uses `swift package edit CompressionFamily --path ../CompressionFamily`. Single-file Package.swift change; no codec change, codestream bytes byte-identical to v10.9.2. diff --git a/RELEASE_NOTES_v10.11.0.md b/RELEASE_NOTES_v10.11.0.md new file mode 100644 index 00000000..616b503b --- /dev/null +++ b/RELEASE_NOTES_v10.11.0.md @@ -0,0 +1,195 @@ +# J2KSwift v10.11.0 + +**JP3D batched GPU iDWT.** Closes a multi-week JP3D arc with a +single-dispatch Metal kernel that amortises per-slice GPU overhead +across the whole volume. JP3D full-volume decode wins **-5 ms on +small CT to -115 ms on 16M-voxel CT** vs the per-slice serial loop +(M2 release, in-process, 7 runs / 2 warmups). + +Decoder-only release; codestream bytes are byte-identical to v10.10.0. +Encoder unchanged. MINOR per RELEASING.md — no public API removed, +no signature change, no default flip that affects bytes; the new +`_jp3dDecodeToCoefficients` / `_jp3dIDWTAndFinalize` / +`_jp3dIDWTAndFinalizeBatched` bridge SPI on `J2KDecoder` ships as +opaque additive surface. + +## Summary + +Three coordinated changes inside `Sources/J2KCodec` + `Sources/J2KMetal` ++ `Sources/J2K3D`: + +1. **Batched single-dispatch Metal inverse 5/3 iDWT.** Two new Metal + kernels (`j2k_dwt_inverse_53_horizontal_int_tiled_batched`, + `..._vertical_..._batched`) extend the existing v10.3 tiled + threadgroup-memory kernels with a Z-dim grid axis. One dispatch + per multi-level chain processes N slices in parallel. Kernel-level + bench (16 slices × 256×256 × 3 levels, M2 release): 5.93 ms + serial → 2.06 ms batched → **2.4× faster kernel-level dispatch**. + +2. **JP3D bridge SPI on `J2KDecoder`.** Three opaque-payload methods + split the per-slice 2D decode pipeline so JP3D can collect N + slices' dequantized coefficients before submitting one batched + iDWT: + - `_jp3dDecodeToCoefficients(_:Data) -> JP3DSliceCoefficients` + (Stage A — entropy + dequant) + - `_jp3dIDWTAndFinalize(_:JP3DSliceCoefficients) -> J2KImage` + (Stages B–C — iDWT + colour + DC + reconstruct) + - `_jp3dIDWTAndFinalizeBatched(_:[JP3DSliceCoefficients]) -> [J2KImage]` + (batched Stages B–C across the whole z-range) + Bit-exact composition: `decode(data) ≡ _jp3dIDWTAndFinalize(_jp3dDecodeToCoefficients(data))` + on every single-tile codestream, which is the JP3D wire format. + +3. **`JP3DSliceStackCodec` consumes the batched bridge.** The per-tile + decode loop now runs two new bulk stages before the existing + sequential Z-delta residual chain: + - Parallel `_jp3dDecodeToCoefficients` across `[zStart, zUpper)` via + a TaskGroup (entropy + dequant concurrently across slices). + - ONE batched `_jp3dIDWTAndFinalizeBatched` call across the whole + z-range (per-slice GPU dispatch overhead amortises). + The Z-delta residual chain remains sequential (it's cheap Int32 + adds, not the bottleneck). Eligibility gate: only the (K=0, + no ROI) production case routes through the bulk path; `K > 0` + uses `decodeResolution` and any ROI uses `decodeRegion` per + slice (v10.5 / v10.6 territory; batched bridge for those is + future work). + +## What's New — production-default + +| Public API | v10.10.0 behaviour | v10.11.0 behaviour | +|---|---|---| +| `JP3DDecoder().decode(data)` full-volume decode (default config) | Per-slice serial `J2KDecoder.decode` loop inside `JP3DSliceStackCodec` | Parallel `_jp3dDecodeToCoefficients` + ONE batched iDWT dispatch + sequential Z-delta residual chain. **Same output bytes**; 5–115 ms faster wall on the JP3D corpus | +| `JP3DDecoder(configuration: cfg).decode(data)` with `cfg.resolutionLevel > 0` | Per-slice `decodeResolution` loop (v10.10.0) | Unchanged — partial-resolution path not yet bridged | +| `JP3DROIDecoder().decode(data, region:)` | Per-slice `decodeRegion` loop (v10.10.0) | Unchanged — ROI path not yet bridged | + +The new `J2KDecoder` bridge SPI methods (`_jp3dDecodeToCoefficients`, +`_jp3dIDWTAndFinalize`, `_jp3dIDWTAndFinalizeBatched`) are public but +underscored — JP3D-internal use only; consumers outside `J2K3D` should +keep calling the normal `J2KDecoder.decode(_:)`. + +## What's New — opt-in / opt-out + +`J2K_JP3D_BATCHED_BRIDGE=0` env var disables the new batched path, +forcing the per-slice serial loop (the pre-v10.11 shape). Diagnostic- +A/B only; production should leave it unset. + +## Backward compatibility + +- **Codestream bytes**: byte-identical to v10.10.0 on every input. The + encoder is unchanged. +- **`JP3DDecoder().decode(data)`** is byte-identical to v10.10.0 on every + input. Validated by the 519-test `swift test --filter JP3D` regression + sweep (519/519 PASS) — including the v10.18 round-trip lossless + baseline, partial-resolution shapes, ROI footprint-skip, and Z-narrow + ROI cases. +- **`_jp3dIDWTAndFinalizeBatched([coefs])[i]`** is byte-identical to + `_jp3dIDWTAndFinalize(coefs[i])` for every JP3D-shape input. + Validated by the new `V10_20_BatchedBridgeParityTests` (5/5 PASS) + + `V10_20_JP3DBridgeParityTests` (5/5 PASS). +- **Behaviour change**: none. + +## Measured wins — JP3D corpus + +M2 release, J2KBenchMac --jp3d, in-process, 7 timed runs / 2 warmups, +median per fixture: + +| Fixture (modality WxHxD) | voxels | serial ms | batched ms | Δ ms | ratio | +|---|---:|---:|---:|---:|---:| +| mr_3d_small MR 128×128×16 | 262K | 13.81 | 14.05 | +0.24 | 0.98× | +| ct_3d_small CT 256×256×16 | 1.05M | 44.11 | 38.99 | **−5.12** | **1.13×** | +| us_3d_small US 320×240×24 | 1.84M | 60.44 | 56.29 | **−4.15** | **1.07×** | +| mr_3d_mid MR 256×256×32 | 2.10M | 87.21 | 77.71 | **−9.50** | **1.12×** | +| ct_3d_mid CT 512×512×32 | 8.39M | 369.81 | 316.56 | **−53.25** | **1.17×** | +| ct_3d_large CT 512×512×64 | 16.78M | 789.15 | 674.59 | **−114.57** | **1.17×** | + +5/6 fixtures clear the 3 ms acceptance threshold. The single wash +(mr_3d_small @ 13 ms wall) is the smallest fixture where per-slice +dispatch overhead dominates the absolute wall. + +Reading: the win scales with slice count × per-slice work — every +extra slice amortises the GPU dispatch overhead it would have paid +in the serial loop. **iDWT is the largest stage in JP3D 5/3 lossless +decode**, so the batched dispatch wedge is structural. + +Raw bench JSONs are committed at +`Documentation/Benchmarks/data/jp3d-bench-arm64-v10_20-{batched,serial}-20260524.json`. + +## Cross-codec parity (2D codec unchanged) + +The 2D codec path (entropy, dequant, single-tile iDWT, colour, DC, +reconstruct) is touched only at the bridge SPI splitting point — +the same per-stage code paths execute, same SubbandInfo flows, same +J2KDWT2DOptimizer / J2KMetalDWT calls. Cross-codec parity vs +OpenJPH 0.27.0 / Grok 20.3.0 / Kakadu 8.4.1 is therefore inherited +from v10.10.0 (no change measured or asserted in this release). + +## Test Suite Results + +| Suite | Tests | Result | Coverage | +|---|---:|---|---| +| `V10_20_JP3DBridgeParityTests` | 5/5 | PASS | Phase 1 bridge SPI bit-exact composition | +| `V10_20_BatchedInverseInt32ParityTests` | 12/12 | PASS | Batched Metal kernel + multi-level orchestrator bit-exact vs serial GPU | +| `V10_20_BatchedBridgeParityTests` | 5/5 | PASS | Batched bridge SPI bit-exact vs per-slice serial bridge SPI | +| `swift test --filter JP3D` (regression sweep) | 519/519 | PASS | Full JP3D test suite green with the batched wiring | +| Mandatory commit gate (release mode): `J2KMedicalCorpusEncodePerformanceTests` + `J2KMedicalCorpusPerformanceTests` + `J2KStrictCrossCodecValidationTests` | 7/7 | PASS | Encode-perf + decode-perf + cross-codec strict validation | + +## API surface — additions only + +```swift +extension J2KDecoder { + public struct JP3DSliceCoefficients: Sendable { + public var width: Int { get } + public var height: Int { get } + // No other public fields — payload opaque to consumers + } + public func _jp3dDecodeToCoefficients(_ data: Data) async throws -> JP3DSliceCoefficients + public func _jp3dIDWTAndFinalize(_ coefs: JP3DSliceCoefficients) async throws -> J2KImage + public func _jp3dIDWTAndFinalizeBatched(_ coefsBatch: [JP3DSliceCoefficients]) async throws -> [J2KImage] +} +``` + +No removals. No signature changes. + +## Known limitations + +- **Partial-resolution + ROI paths not yet batched.** When + `JP3DDecoderConfiguration.resolutionLevel > 0` or + `JP3DROIDecoder().decode(_, region:)` is invoked, the per-slice + loop keeps the v10.10.0 shape (per-slice `decodeResolution` / + `decodeRegion`). Bridging those paths needs the 2D + `decodeResolution` / `decodeRegion` to expose their own coefficient- + split SPI; tracked for a future release. +- **Slice-stack residual chain still sequential.** The Z-delta residual + apply is a per-slice Int32 add with a cross-slice data dependency + (slice z reads `prevSliceInt`). Not the bottleneck (iDWT was), but + not parallelised either. + +## Reproducing the headline numbers + +```bash +# Build +swift build -c release --product J2KBenchMac + +# Batched (default) +.build/release/J2KBenchMac --jp3d --output /tmp/jp3d_batched.json + +# Serial baseline (opt-out) +J2K_JP3D_BATCHED_BRIDGE=0 .build/release/J2KBenchMac --jp3d --output /tmp/jp3d_serial.json +``` + +The bench corpus is fixed (LCG-synthesised volumes at MR/CT/US +shapes), so the wall numbers reproduce across hosts modulo silicon +class. + +## Companion documents + +- `Documentation/Benchmarks/data/jp3d-bench-arm64-v10_20-batched-20260524.json` — raw bench (batched) +- `Documentation/Benchmarks/data/jp3d-bench-arm64-v10_20-serial-20260524.json` — raw bench (serial baseline) +- `Documentation/research/V10_20_BATCHED_JP3D_IDWT.md` (on the `v10.19-research` branch) — full multi-week arc closure with phase-by-phase deliverables + +## Backward upgrade + +`swift package update` will not auto-pick this release if your +`Package.swift` pins exact version; bump the requirement to +`from: "10.11.0"` (or accept the next `.upToNextMinor` per your +policy). Consumers of `JP3DDecoder` see only a perf improvement; +no source changes required. diff --git a/Sources/J2K3D/JP3DSliceStackCodec.swift b/Sources/J2K3D/JP3DSliceStackCodec.swift index 4b26c2fb..e60b9ecb 100644 --- a/Sources/J2K3D/JP3DSliceStackCodec.swift +++ b/Sources/J2K3D/JP3DSliceStackCodec.swift @@ -467,6 +467,60 @@ struct JP3DSliceStackCodec: Sendable { // Int32 — used to add the residual when `is_residual` is set. var prevSliceInt: [[Int32]]? = nil + // v10.20-research Phase 3c — bulk batched-bridge fast path. + // When neither resolutionLevel nor regionOfInterest is set + // (the dominant JP3D production case for full-volume decode), + // we can: + // 1. Decode all slices [zStart, zUpper) to coefficients in + // parallel via the Phase 1 bridge SPI. + // 2. Run ONE batched iDWT + finalize across the whole z-range + // via the Phase 3d bridge SPI (2.4× kernel-level speedup + // vs serial per-slice iDWT on M2 release). + // 3. Apply the Z-delta residual chain sequentially after the + // iDWTs land (the residual is a cheap Int32 add but has + // the cross-slice dependency that prevents batching). + // For K > 0 or any ROI request we keep the per-slice decode + // loop below (it uses decodeResolution / decodeRegion which + // the bridge SPI doesn't expose yet — Phase 3c/v2 territory). + // v10.20-research Phase 4 — env-var off-switch for A/B benching. + // `J2K_JP3D_BATCHED_BRIDGE=0` disables the bulk batched path, + // forcing the per-slice serial loop (the pre-Phase-3c shape) + // so JP3DBench / external benches can measure the delta. + // Default: enabled (no env var or any other value → batched). + let batchedOptedOut: Bool = { + if let v = ProcessInfo.processInfo + .environment["J2K_JP3D_BATCHED_BRIDGE"] { + return v == "0" || v.lowercased() == "false" || v.lowercased() == "no" + } + return false + }() + let useBatchedBridge = (K == 0) && (regionOfInterest == nil) && !batchedOptedOut + var batchedImages: [J2KImage]? = nil + if useBatchedBridge { + let nSlices = zUpper - zStart + var coefs: [JP3DSliceCoefficients?] = Array( + repeating: nil, count: nSlices) + try await withThrowingTaskGroup( + of: (Int, JP3DSliceCoefficients).self + ) { group in + for z in zStart..0, no ROI) → decoder.decodeResolution (Phase 2) // (K>0, ROI) → decoder.decodeRegion (Phase 2+3) @@ -499,7 +554,9 @@ struct JP3DSliceStackCodec: Sendable { + "is the wiring task.") } let image: J2KImage - if let roi = regionOfInterest { + if let batched = batchedImages { + image = batched[z - zStart] + } else if let roi = regionOfInterest { let opts = J2KROIDecodingOptions( region: J2KRegion( x: roi.xRange.lowerBound, y: roi.yRange.lowerBound, @@ -507,6 +564,9 @@ struct JP3DSliceStackCodec: Sendable { strategy: .direct) image = try await decoder.decodeRegion(codestream, options: opts) } else if K == 0 { + // Unreachable when useBatchedBridge is true; kept as a + // safe fallback in case the batched path ever throws + // before producing all images. image = try await decoder.decode(codestream) } else { if perSliceDecompLevels == nil { diff --git a/Sources/J2KCodec/J2KDecoderPipeline.swift b/Sources/J2KCodec/J2KDecoderPipeline.swift index f2ffed72..7596b863 100644 --- a/Sources/J2KCodec/J2KDecoderPipeline.swift +++ b/Sources/J2KCodec/J2KDecoderPipeline.swift @@ -1484,6 +1484,320 @@ struct DecoderPipeline: Sendable { } } + // MARK: - JP3D bridge SPI (v10.20-research Phase 1) + // + // Two new internal entry points that EXPOSE the pipeline split + // between dequantization and inverse-wavelet-transform without + // changing the production single-tile decode path. The composition + // + // iDWTAndFinalizeCoefficients(decodeToCoefficients(data)) + // + // is byte-identical to `decodeSingleTile(parseCodestream(data))` + // on every single-tile codestream — which is what JP3D's slice- + // stack codec emits for each Z-slice. The JP3D side then has the + // freedom to: + // + // • decode N slices to coefficients (Stage A — entropy + dequant) + // • submit ONE batched GPU iDWT dispatch across all N slices + // • finalize each slice (Stage C — colour + DC + reconstruct) + // + // which is the structural shape Phase 2 needs. Phase 1 ships only + // the split — single-slice batched iDWT collapses to today's behaviour. + + /// JP3D bridge — internal Phase 1 helper. Runs the single-tile + /// pipeline through Stage 4 (dequantization) and stops. Returns + /// the dequantized subbands + the codestream metadata so the + /// caller can later run iDWT + colour + DC + reconstruct via + /// `iDWTAndFinalizeCoefficients`. + /// + /// Multi-tile is rejected — JP3D slices are always single-tile + /// 2D J2K codestreams (per JP3DSliceStackCodec wire format). + /// Routing a multi-tile codestream through this bridge would + /// either need a per-tile-group decode or an explicit caller- + /// side multi-tile orchestration, both of which are out of scope. + mutating func decodeToCoefficients( + _ data: Data + ) async throws -> _JP3DSliceCoefficientsInternal { + let (metadata, tiles) = try parseCodestream(data) + guard !metadata.isMultiTile else { + throw J2KError.notImplemented( + "JP3D bridge: decodeToCoefficients does not yet " + + "support multi-tile codestreams (slice-stack slices " + + "are always single-tile).") + } + let tileData = tiles.first?.tileData ?? Data() + + let codeBlocks = try extractTileData( + tileData, metadata: metadata, + maxResolutionLevel: partialResolutionLevel, + regionOfInterest: regionOfInterest, + maxQualityLayer: maxQualityLayer) + + let (decodedBlocks, _) = try await applyEntropyDecoding( + codeBlocks, metadata: metadata) + + let dequantizedSubbands = try await applyDequantization( + decodedBlocks, metadata: metadata) + + return _JP3DSliceCoefficientsInternal( + metadata: metadata, + dequantizedSubbands: dequantizedSubbands) + } + + /// v10.20-research Phase 3b — JP3D batched bridge SPI. Takes a + /// batch of N slice coefficient bundles (each previously produced + /// by `decodeToCoefficients`) and runs ONE batched multi-level + /// GPU iDWT across all N slices via + /// `J2KMetalDWT.inverse2DInt32MultiLevelFusedBatched`, then + /// finalises each slice into a `J2KImage`. + /// + /// **JP3D production-shape only**: requires 5/3 reversible filter, + /// single-component slices (componentIndex 0 only), full- + /// resolution decode (no `partialResolutionLevel`), uniform + /// dimensions across the batch (JP3D slice-stack guarantees this). + /// Any other shape falls back to per-slice serial via + /// `iDWTAndFinalizeCoefficients` so the SPI never errors on the + /// non-JP3D case — it just doesn't get the batched win. + /// + /// Bit-exact composition guarantee: for every slice `i`, + /// batched[i].components[0].data == + /// iDWTAndFinalizeCoefficients(coefs[i]).components[0].data + /// on every JP3D-shape input. + mutating func iDWTAndFinalizeCoefficientsBatched( + _ coefsBatch: [_JP3DSliceCoefficientsInternal] + ) async throws -> [J2KImage] { + guard !coefsBatch.isEmpty else { return [] } + + // v10.20-research Phase 3d — orchestrator integration. The + // four-way diagnostic V10_20_BridgeInputsFourWayDiagnostic + // proved that on the per-slice GPU-chain shape this method + // builds, all four CPU/GPU 5/3 iDWT implementations + // (batched orchestrator / serial GPU multi-level fused / + // J2KMetalDWT per-level CPU / J2KDWT2DOptimizer multi-level + // CPU) produce bit-exact identical output. So the previous + // two integration attempts diverged not because the + // orchestrator was wrong but because the chain construction + // didn't match the per-slice raw-coefficient shape the + // diagnostic uses. This integration mirrors the diagnostic + // exactly — raw `.coefficients` for every band, no + // `getSubbandAsInt32` rounding (which only fires for the + // irreversible 9/7 path that JP3D doesn't use anyway since + // dequantization sets `doubleCoefficients = nil` for 5/3). + // + // **Eligibility gate.** Only JP3D-production shapes route + // through the batched path: + // • 5/3 reversible filter + // • single-component slices (`componentCount == 1`) + // • uniform dimensions across the batch + // • full-resolution decode (no `partialResolutionLevel`) + // • no ROI + // • ≥ 1 decomposition level + // Anything else falls back to the per-slice serial loop — + // same output, no batched win, never errors on a valid + // bundle. + + let head = coefsBatch[0] + let isReversible53: Bool = { + if case .irreversible97 = head.metadata.configuration.waveletFilter { + return false + } + return true + }() + let levels = head.metadata.configuration.decompositionLevels + let headW = head.metadata.width + let headH = head.metadata.height + let headComponentCount = head.metadata.componentCount + + var batchable = isReversible53 + && headComponentCount == 1 + && levels >= 1 + && partialResolutionLevel == nil + && regionOfInterest == nil + + if batchable { + for slice in coefsBatch { + if slice.metadata.width != headW + || slice.metadata.height != headH + || slice.metadata.componentCount != 1 { + batchable = false + break + } + if case .irreversible97 = slice.metadata.configuration.waveletFilter { + batchable = false + break + } + if slice.metadata.configuration.decompositionLevels != levels { + batchable = false + break + } + } + } + + guard batchable else { + var results: [J2KImage] = [] + results.reserveCapacity(coefsBatch.count) + for coefs in coefsBatch { + results.append(try await iDWTAndFinalizeCoefficients(coefs)) + } + return results + } + + // Build the per-slice GPU chain. Identical shape to what + // V10_20_BridgeInputsFourWayDiagnostic builds — raw + // `sb.coefficients`, padding via `padFlatInt32`, level chain + // deepest-first, LL only at the deepest level (subsequent + // levels rely on GPU-resident chain reuse). + let subsX = head.metadata.components[0].subsamplingX + let subsY = head.metadata.components[0].subsamplingY + let compW = headW / subsX + let compH = headH / subsY + var levelSizes: [(width: Int, height: Int)] = [] + for d in 0...levels { + let denom = 1 << d + let bandX1 = EncoderPipeline.ceilDivIntegerOrigin(compW, denom) + let bandY1 = EncoderPipeline.ceilDivIntegerOrigin(compH, denom) + levelSizes.append((bandX1, bandY1)) + } + + let llDimsW = levelSizes[levels].width + let llDimsH = levelSizes[levels].height + + var perSliceChains: [[J2KMetalDWTSubbandsInt32]] = [] + perSliceChains.reserveCapacity(coefsBatch.count) + + for slice in coefsBatch { + let subs = slice.dequantizedSubbands.filter { + $0.componentIndex == 0 + } + + // LL lookup: drop the `level ==` filter and match + // `subband == .ll` only — the entropy stage stores LL + // at `level: 0` in the HT 5/3 path, not at `level: levels` + // (the only legacy helper that uses `level: levels` is + // `buildLLSubbandsFromBuffer` which is no longer called). + // The production serial path at line ~4245 + // (`compSubbands.first(where: { $0.subband == .ll })`) + // works because there is always exactly ONE LL per + // component regardless of which level field it carries. + let initialLL: [Int32] + if let ll = subs.first(where: { $0.subband == .ll }) { + initialLL = padFlatInt32( + ll.coefficients, srcW: ll.width, srcH: ll.height, + dstW: llDimsW, dstH: llDimsH) + } else { + initialLL = [Int32](repeating: 0, count: llDimsW * llDimsH) + } + + var chain: [J2KMetalDWTSubbandsInt32] = [] + chain.reserveCapacity(levels) + for level in (1...levels).reversed() { + let parentW = levelSizes[level - 1].width + let parentH = levelSizes[level - 1].height + let llW = levelSizes[level].width + let llH = levelSizes[level].height + let hlW = parentW - llW + let lhH = parentH - llH + + func raw(_ band: J2KSubband, _ dstW: Int, _ dstH: Int) -> [Int32] { + if let sb = subs.first(where: { $0.level == level && $0.subband == band }) { + return padFlatInt32(sb.coefficients, + srcW: sb.width, srcH: sb.height, + dstW: dstW, dstH: dstH) + } + return [Int32](repeating: 0, count: dstW * dstH) + } + + let llForLevel: [Int32] = chain.isEmpty ? initialLL : [] + chain.append(J2KMetalDWTSubbandsInt32( + ll: llForLevel, + lh: raw(.lh, llW, lhH), + hl: raw(.hl, hlW, llH), + hh: raw(.hh, hlW, lhH), + llWidth: llW, llHeight: llH, + originalWidth: parentW, originalHeight: parentH, + tileOriginX: 0, tileOriginY: 0)) + } + perSliceChains.append(chain) + } + + // ONE batched multi-level dispatch across all slices. + let metalDWT = J2KMetalDWT( + configuration: J2KMetalDWTConfiguration( + filter: .reversible53, decompositionLevels: levels), + device: metalSession?.device, + bufferPool: metalSession?.bufferPool, + shaderLibrary: metalSession?.shaderLibrary) + try await metalDWT.initialize() + let perSliceFlatInt32 = try await metalDWT.inverse2DInt32MultiLevelFusedBatched( + perSliceSubbandsPerLevel: perSliceChains) + + // Per-slice finalize: Int32 → Double → DC unshift → reconstruct. + var results: [J2KImage] = [] + results.reserveCapacity(coefsBatch.count) + for (sIdx, slice) in coefsBatch.enumerated() { + var spatialData: [[Double]] = [ + vDSPConvert.int32sToDoubles(perSliceFlatInt32[sIdx]) + ] + // 1-component MCT is a no-op; call for parity with serial path. + try applyInverseColorTransformInPlace(&spatialData, + metadata: slice.metadata) + let compInfo = slice.metadata.components[0] + if !compInfo.signed { + var dcOffset = Double(1 << (compInfo.bitDepth - 1)) + spatialData[0].withUnsafeMutableBufferPointer { buf in + #if canImport(Accelerate) + vDSP_vsaddD(buf.baseAddress!, 1, &dcOffset, + buf.baseAddress!, 1, vDSP_Length(buf.count)) + #else + for i in 0.. J2KImage { + let metadata = coefs.metadata + + var spatialData = try await applyInverseWaveletTransform( + coefs.dequantizedSubbands, metadata: metadata) + + try applyInverseColorTransformInPlace(&spatialData, metadata: metadata) + var rgbData = spatialData + + for (compIdx, compInfo) in metadata.components.enumerated() { + guard compIdx < rgbData.count else { break } + if !compInfo.signed { + var dcOffset = Double(1 << (compInfo.bitDepth - 1)) + rgbData[compIdx].withUnsafeMutableBufferPointer { buf in + #if canImport(Accelerate) + vDSP_vsaddD(buf.baseAddress!, 1, &dcOffset, + buf.baseAddress!, 1, vDSP_Length(buf.count)) + #else + for i in 0.. CPU iDWT total +// (V10_19_JP3D_GPU_IDWT_CLOSED.md). +// +// The only path to a real GPU iDWT win for JP3D is BATCHED single- +// dispatch GPU iDWT — collect N slices' dequantized wavelet +// coefficients, submit ONE Metal dispatch that runs N parallel +// iDWTs (one threadgroup per slice along the Z grid axis), then +// scatter the results back per slice. Per-slice dispatch overhead +// amortises across the volume. +// +// This file ships the **bridge SPI** Phase 1 needs to make that +// possible: two methods on `J2KDecoder` that expose the pipeline +// split between dequantization and iDWT. Phase 2 will add the +// batched-dispatch GPU iDWT itself + the JP3DSliceStackCodec wiring +// that uses the bridge. +// +// **Bit-exact composition guarantee (Phase 1):** +// +// decode(data) ≡ _jp3dIDWTAndFinalize(_jp3dDecodeToCoefficients(data)) +// +// on every single-tile codestream — which is the wire format JP3D +// slices always use. +// +// Surface is `public func` (not `internal` / `@_spi`) because J2K3D +// depends on J2KCodec as a separate module; making the bridge SPI +// public lets J2K3D call it without an SPI-flag handshake. The leading +// underscore (`_jp3d...`) marks the methods as not-for-general-use: +// they're stable across patch versions but the intermediate +// `JP3DSliceCoefficients` type is opaque (no public fields) and +// consumers outside J2K3D should keep calling the normal +// `J2KDecoder.decode(_:)`. + +import Foundation +import J2KCore + +// MARK: - Internal payload + +/// Opaque internal payload carried inside `JP3DSliceCoefficients`. +/// Not for consumption outside this module — accessed only by +/// `DecoderPipeline.iDWTAndFinalizeCoefficients`. +@usableFromInline +struct _JP3DSliceCoefficientsInternal: @unchecked Sendable { + let metadata: CodestreamMetadata + let dequantizedSubbands: [DecoderPipeline.SubbandInfo] +} + +// MARK: - Public opaque coefficient bundle + +/// Opaque dequantized-wavelet-coefficient bundle. Produced by +/// `J2KDecoder._jp3dDecodeToCoefficients(_:)`; consumed by +/// `J2KDecoder._jp3dIDWTAndFinalize(_:)`. Carrying this between the +/// two calls lets the JP3D slice-stack codec collect N slices' +/// coefficients before submitting one batched GPU iDWT dispatch +/// (Phase 2). +/// +/// The contained payload is internal to J2KCodec; no public fields. +public struct JP3DSliceCoefficients: @unchecked Sendable { + @usableFromInline + let _internal: _JP3DSliceCoefficientsInternal + + @usableFromInline + init(_internal: _JP3DSliceCoefficientsInternal) { + self._internal = _internal + } + + /// Pixel dimensions of the slice this coefficients bundle will + /// produce after iDWT + reconstruction. Useful for the JP3D + /// caller's output-buffer sizing before iDWT runs. + public var width: Int { _internal.metadata.width } + public var height: Int { _internal.metadata.height } +} + +// MARK: - J2KDecoder SPI + +extension J2KDecoder { + + /// v10.20-research JP3D bridge — Stage A: decode a single-tile + /// 2D codestream THROUGH dequantization but BEFORE the inverse + /// wavelet transform. Returns an opaque `JP3DSliceCoefficients` + /// bundle the caller will later feed to `_jp3dIDWTAndFinalize`. + /// + /// **Bit-exact composition guarantee:** + /// ```swift + /// let direct = try await decoder.decode(data) + /// let coefs = try await decoder._jp3dDecodeToCoefficients(data) + /// let bridged = try await decoder._jp3dIDWTAndFinalize(coefs) + /// // direct.components[i].data == bridged.components[i].data, byte-exact + /// ``` + /// + /// **Single-tile only.** JP3D's slice-stack codec only ever emits + /// single-tile per-slice codestreams, so this restriction is + /// invisible at the JP3D layer. A multi-tile codestream passed + /// through this SPI throws `J2KError.notImplemented`. + /// + /// **Not for general use** — call the normal `decode(_:)` from + /// outside the JP3D codec. The leading underscore marks this as + /// a JP3D-specific bridge; the bundle's payload is opaque. + public func _jp3dDecodeToCoefficients( + _ data: Data + ) async throws -> JP3DSliceCoefficients { + var pipeline = DecoderPipeline() + pipeline.metalSession = J2KMetalSession.processShared + let inner = try await pipeline.decodeToCoefficients(data) + return JP3DSliceCoefficients(_internal: inner) + } + + /// v10.20-research JP3D bridge — Stage B+C: take a coefficient + /// bundle produced by `_jp3dDecodeToCoefficients` and run iDWT + + /// inverse colour transform + DC level unshift + image + /// reconstruction, returning the final `J2KImage`. + /// + /// Equivalent to running the second half of `decode(_:)` on the + /// stage-A output — but available as a separate call so the JP3D + /// caller can collect N bundles and submit a single batched GPU + /// iDWT dispatch between Stage A and Stage B (Phase 2). + public func _jp3dIDWTAndFinalize( + _ coefs: JP3DSliceCoefficients + ) async throws -> J2KImage { + var pipeline = DecoderPipeline() + pipeline.metalSession = J2KMetalSession.processShared + return try await pipeline.iDWTAndFinalizeCoefficients(coefs._internal) + } + + /// v10.20-research Phase 3b — JP3D batched bridge SPI. Takes N + /// coefficient bundles previously produced by + /// `_jp3dDecodeToCoefficients` and runs ONE batched multi-level + /// GPU iDWT dispatch across all N slices via + /// `J2KMetalDWT.inverse2DInt32MultiLevelFusedBatched`, then + /// finalises each slice into a `J2KImage`. + /// + /// For the JP3D-typical production shape (5/3 reversible, single- + /// component slices, uniform dims across the batch, full + /// resolution, no ROI), this delivers the per-slice GPU dispatch + /// amortisation that Phase 3a measured at **2.4× on M2** for + /// 16-slice 256×256×3-level batches. + /// + /// For any other shape (9/7 lossy, multi-component, non-uniform + /// dims, partial-resolution, ROI) the bridge falls back to a + /// per-slice serial loop via `_jp3dIDWTAndFinalize` — same + /// output, no batched win. The fallback is invisible to callers + /// and ensures the SPI never errors on a valid coefficient + /// bundle. + /// + /// **Bit-exact composition guarantee**: for every slice `i` in + /// the batch, + /// ``` + /// batched[i].components[0].data == + /// _jp3dIDWTAndFinalize(coefsBatch[i]).components[0].data + /// ``` + /// — verified by `V10_20_BatchedBridgeParityTests`. + public func _jp3dIDWTAndFinalizeBatched( + _ coefsBatch: [JP3DSliceCoefficients] + ) async throws -> [J2KImage] { + var pipeline = DecoderPipeline() + pipeline.metalSession = J2KMetalSession.processShared + return try await pipeline.iDWTAndFinalizeCoefficientsBatched( + coefsBatch.map { $0._internal }) + } +} diff --git a/Sources/J2KMetal/J2KMetalDWT.swift b/Sources/J2KMetal/J2KMetalDWT.swift index f94d61e5..159cd08d 100644 --- a/Sources/J2KMetal/J2KMetalDWT.swift +++ b/Sources/J2KMetal/J2KMetalDWT.swift @@ -2483,6 +2483,585 @@ public actor J2KMetalDWT { return result } + // MARK: - v10.20-research Phase 2 — BATCHED inverse 5/3 Int + + /// Performs N parallel inverse 2D 5/3 Int transforms in one Metal + /// command buffer. The batch's `subbandsBatch[i]` produces the + /// returned `[Int32]` at index `i`. **All slices in the batch + /// MUST have identical dimensions** (originalWidth, originalHeight, + /// llWidth, llHeight) — JP3D's slice-stack codec guarantees this + /// (every slice in a tile shares header.tileWidth × tileHeight), + /// and this method throws `J2KError.invalidParameter` on mixed + /// dimensions so the caller cannot silently produce wrong output. + /// + /// Why batched: per-slice GPU dispatch overhead × N slices + /// regressed JP3D decode in the naive forced-per-slice routing + /// (`V10_19_JP3D_GPU_IDWT_CLOSED.md`). One dispatch over N slices + /// amortises that overhead across the volume; the new batched + /// kernels `j2k_dwt_inverse_53_{horizontal,vertical}_int_tiled_batched` + /// use a Z grid dimension where `tgid.z = slice index`, with + /// per-slice stride offsets passed as constants. + /// + /// Bit-exact: each output slice is byte-identical to what + /// `inverse2DInt32(subbands: subbandsBatch[i], backend: .metal)` + /// would produce. Parity test: + /// `V10_20_BatchedInverseInt32ParityTests`. + /// + /// **Single-slice batches collapse to current behaviour** — a + /// batch of 1 dispatches the same kernel with `gridDepth = 1`, + /// equivalent to the existing tiled path. + public func inverseBatched2DInt32( + subbandsBatch: [J2KMetalDWTSubbandsInt32] + ) async throws -> [[Int32]] { + guard !subbandsBatch.isEmpty else { return [] } + + // Uniform-dimension guard. JP3D slice-stack guarantees this; + // exposing the SPI more generally would require fall-through + // to serial dispatch for non-uniform batches (not implemented). + let head = subbandsBatch[0] + for (i, s) in subbandsBatch.enumerated() { + guard s.originalWidth == head.originalWidth, + s.originalHeight == head.originalHeight, + s.llWidth == head.llWidth, + s.llHeight == head.llHeight else { + throw J2KError.invalidParameter( + "inverseBatched2DInt32: slice \(i) dimensions " + + "\(s.originalWidth)x\(s.originalHeight) / " + + "ll=\(s.llWidth)x\(s.llHeight) do not match batch " + + "head \(head.originalWidth)x\(head.originalHeight) / " + + "ll=\(head.llWidth)x\(head.llHeight). All slices " + + "in a batch must share identical dimensions.") + } + } + + let width = head.originalWidth + let height = head.originalHeight + guard width >= 2, height >= 2 else { + throw J2KError.invalidParameter( + "Subband dimensions must produce at least 2×2 output") + } + + let nSlices = subbandsBatch.count + let startTime = currentTime() + _statistics.totalOperations += 1 + + // Currently batched path is GPU-only. CPU fallback would + // collapse to the serial loop — exposed as a future-work + // option if needed; not needed for the JP3D use case. + let result = try await inverseBatched2DGPUInt32(subbandsBatch: subbandsBatch) + _statistics.gpuOperations += 1 + _statistics.totalProcessingTime += currentTime() - startTime + _ = nSlices + return result + } + + private func inverseBatched2DGPUInt32( + subbandsBatch: [J2KMetalDWTSubbandsInt32] + ) async throws -> [[Int32]] { + try await ensureInitialized() + let queue = try await metalDevice.commandQueue() + let device = queue.device + + let head = subbandsBatch[0] + let width = head.originalWidth + let height = head.originalHeight + let bands = BandGeometry(width: width, height: height, + llW: head.llWidth, llH: head.llHeight) + let llH = bands.llH + let halfHH = bands.halfHH + let nSlices = subbandsBatch.count + + // Per-slice band element counts. The H pass produces colLow + // (width × llH) and colHigh (width × halfHH); V pass produces + // the final output (width × height). + let llCount = max(head.ll.count, 1) + let lhCount = max(head.lh.count, 1) + let hlCount = max(head.hl.count, 1) + let hhCount = max(head.hh.count, 1) + let colLowCount = width * llH + let colHighCount = max(width * halfHH, 1) + let outputCount = width * height + + let stride32 = MemoryLayout.stride + + func makeBuffer(_ size: Int) throws -> any MTLBuffer { + guard let buf = device.makeBuffer(length: max(size, 1), + options: .storageModeShared) else { + throw J2KError.internalError("Failed to allocate Metal buffer") + } + return buf + } + + // Allocate ONE batched buffer per band, sized N × per-slice size. + // Layout: slice s's data starts at s * sliceStride. + let llBuffer = try makeBuffer(nSlices * llCount * stride32) + let lhBuffer = try makeBuffer(nSlices * lhCount * stride32) + let hlBuffer = try makeBuffer(nSlices * hlCount * stride32) + let hhBuffer = try makeBuffer(nSlices * hhCount * stride32) + let colLowBuffer = try makeBuffer(nSlices * colLowCount * stride32) + let colHighBuffer = try makeBuffer(nSlices * colHighCount * stride32) + let outputBuffer = try makeBuffer(nSlices * outputCount * stride32) + + // Copy each slice's coefficients into the corresponding Z stride. + for (s, sub) in subbandsBatch.enumerated() { + if !sub.ll.isEmpty { + sub.ll.withUnsafeBytes { src in + let dst = llBuffer.contents() + s * llCount * stride32 + dst.copyMemory(from: src.baseAddress!, byteCount: src.count) + } + } + if !sub.lh.isEmpty { + sub.lh.withUnsafeBytes { src in + let dst = lhBuffer.contents() + s * lhCount * stride32 + dst.copyMemory(from: src.baseAddress!, byteCount: src.count) + } + } + if !sub.hl.isEmpty { + sub.hl.withUnsafeBytes { src in + let dst = hlBuffer.contents() + s * hlCount * stride32 + dst.copyMemory(from: src.baseAddress!, byteCount: src.count) + } + } + if !sub.hh.isEmpty { + sub.hh.withUnsafeBytes { src in + let dst = hhBuffer.contents() + s * hhCount * stride32 + dst.copyMemory(from: src.baseAddress!, byteCount: src.count) + } + } + } + + let horizontalPSO = try await shaderLibrary.computePipeline(for: .dwtInverse53HorizontalIntTiledBatched) + let verticalPSO = try await shaderLibrary.computePipeline(for: .dwtInverse53VerticalIntTiledBatched) + + guard let cb = queue.makeCommandBuffer() else { + throw J2KError.internalError("Failed to create batched-iDWT command buffer") + } + + // H pass — combine LL+HL into colLow, LH+HH into colHigh. + // We dispatch the batched horizontal kernel twice (once per + // band pair), with N parallel slices each time. Same shape + // as the v10.3 tiled pair, just with Z grid extension. + let halfWidth = (width + 1) / 2 + let halfWidthH = width / 2 + + // Pass A: LL (ll lowpass) + HL (hl highpass) → colLow + var widthVar = UInt32(width) + var heightVarLowH = UInt32(llH) // process llH rows + var sliceStrideLowA = UInt32(llCount) + var sliceStrideHighA = UInt32(hlCount) + var sliceStrideOutA = UInt32(colLowCount) + if let encA = cb.makeComputeCommandEncoder() { + encA.setComputePipelineState(horizontalPSO) + encA.setBuffer(llBuffer, offset: 0, index: 0) + encA.setBuffer(hlBuffer, offset: 0, index: 1) + encA.setBuffer(colLowBuffer, offset: 0, index: 2) + encA.setBytes(&widthVar, length: MemoryLayout.size, index: 3) + encA.setBytes(&heightVarLowH, length: MemoryLayout.size, index: 4) + encA.setBytes(&sliceStrideLowA, length: MemoryLayout.size, index: 5) + encA.setBytes(&sliceStrideHighA,length: MemoryLayout.size, index: 6) + encA.setBytes(&sliceStrideOutA, length: MemoryLayout.size, index: 7) + let tgPerGridX = (halfWidth + 31) / 32 + let tgPerGridY = (llH + 7) / 8 + encA.dispatchThreadgroups( + MTLSize(width: tgPerGridX, height: tgPerGridY, depth: nSlices), + threadsPerThreadgroup: MTLSize(width: 32, height: 8, depth: 1)) + encA.endEncoding() + _ = halfWidthH + } + + // Pass B: LH (lh lowpass) + HH (hh highpass) → colHigh + var heightVarHalfHH = UInt32(halfHH) + var sliceStrideLowB = UInt32(lhCount) + var sliceStrideHighB = UInt32(hhCount) + var sliceStrideOutB = UInt32(colHighCount) + if halfHH > 0, let encB = cb.makeComputeCommandEncoder() { + encB.setComputePipelineState(horizontalPSO) + encB.setBuffer(lhBuffer, offset: 0, index: 0) + encB.setBuffer(hhBuffer, offset: 0, index: 1) + encB.setBuffer(colHighBuffer, offset: 0, index: 2) + encB.setBytes(&widthVar, length: MemoryLayout.size, index: 3) + encB.setBytes(&heightVarHalfHH, length: MemoryLayout.size, index: 4) + encB.setBytes(&sliceStrideLowB, length: MemoryLayout.size, index: 5) + encB.setBytes(&sliceStrideHighB,length: MemoryLayout.size, index: 6) + encB.setBytes(&sliceStrideOutB, length: MemoryLayout.size, index: 7) + let tgPerGridX = (halfWidth + 31) / 32 + let tgPerGridY = (halfHH + 7) / 8 + encB.dispatchThreadgroups( + MTLSize(width: tgPerGridX, height: tgPerGridY, depth: nSlices), + threadsPerThreadgroup: MTLSize(width: 32, height: 8, depth: 1)) + encB.endEncoding() + } + + // V pass: colLow + colHigh → final output. + var heightVarFinal = UInt32(height) + var sliceStrideLowC = UInt32(colLowCount) + var sliceStrideHighC = UInt32(colHighCount) + var sliceStrideOutC = UInt32(outputCount) + if let encC = cb.makeComputeCommandEncoder() { + encC.setComputePipelineState(verticalPSO) + encC.setBuffer(colLowBuffer, offset: 0, index: 0) + encC.setBuffer(colHighBuffer, offset: 0, index: 1) + encC.setBuffer(outputBuffer, offset: 0, index: 2) + encC.setBytes(&widthVar, length: MemoryLayout.size, index: 3) + encC.setBytes(&heightVarFinal, length: MemoryLayout.size, index: 4) + encC.setBytes(&sliceStrideLowC, length: MemoryLayout.size, index: 5) + encC.setBytes(&sliceStrideHighC,length: MemoryLayout.size, index: 6) + encC.setBytes(&sliceStrideOutC, length: MemoryLayout.size, index: 7) + let tgPerGridX = (width + 31) / 32 + let tgPerGridY = ((height + 1) / 2 + 7) / 8 + encC.dispatchThreadgroups( + MTLSize(width: tgPerGridX, height: tgPerGridY, depth: nSlices), + threadsPerThreadgroup: MTLSize(width: 32, height: 8, depth: 1)) + encC.endEncoding() + } + + cb.commit() + await cb.completed() + if cb.status == .error { + throw J2KError.internalError( + "Batched inverse 5/3 GPU dispatch failed: \(cb.error?.localizedDescription ?? "(no description)")") + } + + // Read back per-slice output slices. readInt32Array doesn't + // take a byte offset; copy directly from the typed pointer. + var results: [[Int32]] = [] + results.reserveCapacity(nSlices) + let basePtr = outputBuffer.contents().assumingMemoryBound(to: Int32.self) + J2KMetalUMACounters.incrementContents() + for s in 0.. 0, the head slice's + /// llWidth/llHeight must equal level k-1's + /// originalWidth/originalHeight (same as the per-slice + /// multi-level orchestrator's chain validation) + /// + /// One command buffer commit + one await + per-slice readback. + /// The per-level GPU dispatches use the v10.20 Phase 2 batched + /// kernels (`j2k_dwt_inverse_53_{horizontal,vertical}_int_tiled_batched`) + /// with `gridDepth = N_slices`. Per-slice GPU dispatch overhead + /// amortises across the volume — the structural reason JP3D + /// per-slice forced GPU iDWT regressed + /// (V10_19_JP3D_GPU_IDWT_CLOSED.md). + public func inverse2DInt32MultiLevelFusedBatched( + perSliceSubbandsPerLevel: [[J2KMetalDWTSubbandsInt32]] + ) async throws -> [[Int32]] { + guard !perSliceSubbandsPerLevel.isEmpty else { return [] } + let nSlices = perSliceSubbandsPerLevel.count + let nLevels = perSliceSubbandsPerLevel[0].count + guard nLevels > 0 else { + throw J2KError.invalidParameter( + "inverse2DInt32MultiLevelFusedBatched: empty level chain") + } + + // Validate: all slices have same level count; same dims per + // level; level chain coherent. Mirror the single-slice + // multi-level orchestrator's contract. + for s in 1.. 0 { + let prev = perSliceSubbandsPerLevel[0][k - 1] + guard head.llWidth == prev.originalWidth, + head.llHeight == prev.originalHeight else { + throw J2KError.invalidParameter( + "Batched multi-level: level chain mismatch at level \(k): " + + "expected LL=\(prev.originalWidth)x\(prev.originalHeight), " + + "got \(head.llWidth)x\(head.llHeight)") + } + } + } + + try await ensureInitialized() + let queue = try await metalDevice.commandQueue() + let device = queue.device + let stride32 = MemoryLayout.stride + + guard let cb = queue.makeCommandBuffer() else { + throw J2KError.internalError("Failed to create batched multi-level cb") + } + + let horizontalPSO = try await shaderLibrary.computePipeline(for: .dwtInverse53HorizontalIntTiledBatched) + let verticalPSO = try await shaderLibrary.computePipeline(for: .dwtInverse53VerticalIntTiledBatched) + + // Retain ALL allocated batched buffers across the multi-level + // loop so Swift's ARC doesn't release them between when the + // encoders bind them and when cb.commit() / cb.completed() + // run. The existing serial `inverse2DInt32MultiLevelFused` + // uses the same `inFlight` pattern — without it, level-N + // buffers can be reclaimed before level-(N+1) commits, + // producing SIGSEGV on access. (Verified Phase 3b: the + // bench loop crashed at signal 11 without this retention.) + var inFlight: [any MTLBuffer] = [] + + func makeBuffer(_ size: Int) throws -> any MTLBuffer { + guard let buf = device.makeBuffer(length: max(size, 1), + options: .storageModeShared) else { + throw J2KError.internalError("Buffer alloc failed in batched multi-level") + } + inFlight.append(buf) + return buf + } + + // v10.20-research Phase 3b — small-dim guard. The batched + // kernels share the design assumption of the existing v10.3 + // Phase 2-2-tiled single-image kernels: threadgroup-memory + // boundary computation (the t==31 "33rd even" write) assumes + // every thread's main step-1 write actually executed, which + // requires halfWidth >= 32 (i.e. width >= 64) and analogous + // for height. Below threshold the boundary thread reads + // uninitialised threadgroup slots, producing undefined output + // and eventually SIGSEGV. + // + // The existing single-image path side-steps this by + // auto-routing small dims to CPU via `effectiveBackend(.auto)`. + // For the batched orchestrator we mirror that: if ANY level's + // dimensions fall under the per-axis threshold, fall back to + // a per-slice serial multi-level dispatch (no batching at + // all). Loses the batched win on tiny pyramids; preserves + // correctness for every input shape that ships in JP3D + // production (mid+ slices have level 0 ≥ 32×32 for the + // 5-level default). + let smallDimThreshold = 64 + let smallestLevel = perSliceSubbandsPerLevel[0][0] + if smallestLevel.originalWidth < smallDimThreshold + || smallestLevel.originalHeight < smallDimThreshold { + // Serial fallback: run the existing single-image + // multi-level orchestrator per slice. The cb we created + // is unused; let it deallocate. + var serialResults: [[Int32]] = [] + serialResults.reserveCapacity(nSlices) + for s in 0...size, index: 3) + enc.setBytes(&heightVarLowH, length: MemoryLayout.size, index: 4) + enc.setBytes(&sliceStrideLowA, length: MemoryLayout.size, index: 5) + enc.setBytes(&sliceStrideHighA,length: MemoryLayout.size, index: 6) + enc.setBytes(&sliceStrideOutA, length: MemoryLayout.size, index: 7) + let tgX = (halfWidth + 31) / 32 + let tgY = (llH + 7) / 8 + enc.dispatchThreadgroups( + MTLSize(width: tgX, height: tgY, depth: nSlices), + threadsPerThreadgroup: MTLSize(width: 32, height: 8, depth: 1)) + enc.endEncoding() + } + + // H pass B: LH + HH → colHigh (over halfHH rows) + var heightVarHalfHH = UInt32(halfHH) + var sliceStrideLowB = UInt32(lhCount) + var sliceStrideHighB = UInt32(hhCount) + var sliceStrideOutB = UInt32(colHighCount) + if halfHH > 0, let enc = cb.makeComputeCommandEncoder() { + enc.setComputePipelineState(horizontalPSO) + enc.setBuffer(lhBatched, offset: 0, index: 0) + enc.setBuffer(hhBatched, offset: 0, index: 1) + enc.setBuffer(colHighBatched, offset: 0, index: 2) + enc.setBytes(&widthVar, length: MemoryLayout.size, index: 3) + enc.setBytes(&heightVarHalfHH, length: MemoryLayout.size, index: 4) + enc.setBytes(&sliceStrideLowB, length: MemoryLayout.size, index: 5) + enc.setBytes(&sliceStrideHighB,length: MemoryLayout.size, index: 6) + enc.setBytes(&sliceStrideOutB, length: MemoryLayout.size, index: 7) + let tgX = (halfWidth + 31) / 32 + let tgY = (halfHH + 7) / 8 + enc.dispatchThreadgroups( + MTLSize(width: tgX, height: tgY, depth: nSlices), + threadsPerThreadgroup: MTLSize(width: 32, height: 8, depth: 1)) + enc.endEncoding() + } + + // V pass: colLow + colHigh → output + var heightVarFinal = UInt32(height) + var sliceStrideLowC = UInt32(colLowCount) + var sliceStrideHighC = UInt32(colHighCount) + var sliceStrideOutC = UInt32(outputCount) + if let enc = cb.makeComputeCommandEncoder() { + enc.setComputePipelineState(verticalPSO) + enc.setBuffer(colLowBatched, offset: 0, index: 0) + enc.setBuffer(colHighBatched, offset: 0, index: 1) + enc.setBuffer(outputBatched, offset: 0, index: 2) + enc.setBytes(&widthVar, length: MemoryLayout.size, index: 3) + enc.setBytes(&heightVarFinal, length: MemoryLayout.size, index: 4) + enc.setBytes(&sliceStrideLowC, length: MemoryLayout.size, index: 5) + enc.setBytes(&sliceStrideHighC,length: MemoryLayout.size, index: 6) + enc.setBytes(&sliceStrideOutC, length: MemoryLayout.size, index: 7) + let tgX = (width + 31) / 32 + let tgY = ((height + 1) / 2 + 7) / 8 + enc.dispatchThreadgroups( + MTLSize(width: tgX, height: tgY, depth: nSlices), + threadsPerThreadgroup: MTLSize(width: 32, height: 8, depth: 1)) + enc.endEncoding() + } + + currentLLBatched = outputBatched + finalOutputBatched = outputBatched + finalWidth = width + finalHeight = height + } + + cb.commit() + await cb.completed() + if cb.status == .error { + throw J2KError.internalError( + "Batched multi-level fused inverse 5/3 GPU dispatch failed: " + + "\(cb.error?.localizedDescription ?? "(no description)")") + } + + guard let finalBuffer = finalOutputBatched else { + throw J2KError.internalError("No final output buffer after batched multi-level") + } + let outputCount = finalWidth * finalHeight + var results: [[Int32]] = [] + results.reserveCapacity(nSlices) + let basePtr = finalBuffer.contents().assumingMemoryBound(to: Int32.self) + J2KMetalUMACounters.incrementContents() + for s in 0.. CPU iDWT total +// (V10_19_JP3D_GPU_IDWT_CLOSED.md). A batched dispatch amortises the +// overhead across the volume — one kernel launch for the whole tile's +// horizontal pass, one for the vertical pass. +// +// Buffer layout: each band (lowpass / highpass / output) is laid out +// linearly per slice. Slice `s`'s lowpass band starts at +// `lowpass[s * sliceStrideLowpass]` etc. Strides are passed as +// constants so the kernel doesn't need to recompute them. +// +// Threadgroup geometry is unchanged (32 × 8 for horizontal, 32 × 8 +// for vertical). The Z dimension of the dispatch grid is N_slices; +// each threadgroup processes its (xTile, yTile, sliceIndex) tile. +// Threadgroup memory is per-threadgroup — no cross-slice contamination. +// +// Per-slice dimensions (width / height / lowpass-band-dim / highpass- +// band-dim) MUST be uniform across all N slices in a batch. JP3D +// slice-stack guarantees this — every slice in a JP3D tile shares +// `header.tileWidth × header.tileHeight`. + +kernel void j2k_dwt_inverse_53_horizontal_int_tiled_batched( + device const int* lowpass [[buffer(0)]], + device const int* highpass [[buffer(1)]], + device int* output [[buffer(2)]], + constant uint& width [[buffer(3)]], + constant uint& height [[buffer(4)]], + constant uint& sliceStrideLowpass [[buffer(5)]], + constant uint& sliceStrideHighpass [[buffer(6)]], + constant uint& sliceStrideOutput [[buffer(7)]], + uint3 gid [[thread_position_in_grid]], + uint3 lid [[thread_position_in_threadgroup]], + uint3 tgid [[threadgroup_position_in_grid]] +) { + threadgroup int tg_even[8][33]; + + uint row = gid.y; + if (row >= height) return; + + uint halfWidth = (width + 1) / 2; + uint halfWidthH = width / 2; + uint slice = tgid.z; + + // Per-slice band base offsets — point into this slice's portion + // of each batched buffer. + uint lSliceBase = slice * sliceStrideLowpass; + uint hSliceBase = slice * sliceStrideHighpass; + uint oSliceBase = slice * sliceStrideOutput; + + uint lBase = lSliceBase + row * halfWidth; + uint hBase = hSliceBase + row * halfWidthH; + uint oBase = oSliceBase + row * width; + + uint t = lid.x; + uint r = lid.y; + uint tile_base_i = tgid.x * 32; + + if (halfWidthH == 0) { + uint i = tile_base_i + t; + if (i < halfWidth) output[oBase + 2 * i] = lowpass[lBase + i]; + return; + } + + // Step 1 — thread t's even + { + uint i = tile_base_i + t; + if (i < halfWidth) { + int dLeft = (i > 0) + ? highpass[hBase + i - 1] + : highpass[hBase]; + int dRight = (i < halfWidthH) + ? highpass[hBase + i] + : highpass[hBase + halfWidthH - 1]; + int e = lowpass[lBase + i] - ((dLeft + dRight + 2) >> 2); + tg_even[r][t] = e; + output[oBase + 2 * i] = e; + } + } + + // Step 1 boundary — t == 31 also computes the 33rd even + if (t == 31) { + uint i = tile_base_i + 32; + if (i < halfWidth) { + int dLeft = (i > 0) + ? highpass[hBase + i - 1] + : highpass[hBase]; + int dRight = (i < halfWidthH) + ? highpass[hBase + i] + : highpass[hBase + halfWidthH - 1]; + int e = lowpass[lBase + i] - ((dLeft + dRight + 2) >> 2); + tg_even[r][32] = e; + } else { + tg_even[r][32] = tg_even[r][31]; + } + } + + threadgroup_barrier(mem_flags::mem_threadgroup); + + // Step 2 — thread t's odd + uint i = tile_base_i + t; + if (i < halfWidthH) { + int eLeft = tg_even[r][t]; + int eRight = (2 * i + 2 < width) + ? tg_even[r][t + 1] + : tg_even[r][t]; + output[oBase + 2 * i + 1] = highpass[hBase + i] + ((eLeft + eRight) >> 1); + } +} + +kernel void j2k_dwt_inverse_53_vertical_int_tiled_batched( + device const int* lowpass [[buffer(0)]], + device const int* highpass [[buffer(1)]], + device int* output [[buffer(2)]], + constant uint& width [[buffer(3)]], + constant uint& height [[buffer(4)]], + constant uint& sliceStrideLowpass [[buffer(5)]], + constant uint& sliceStrideHighpass [[buffer(6)]], + constant uint& sliceStrideOutput [[buffer(7)]], + uint3 gid [[thread_position_in_grid]], + uint3 lid [[thread_position_in_threadgroup]], + uint3 tgid [[threadgroup_position_in_grid]] +) { + threadgroup int tg_even[33][32]; + + uint col = gid.x; + if (col >= width) return; + + uint halfHeight = (height + 1) / 2; + uint halfHeightH = height / 2; + uint slice = tgid.z; + + uint lSliceBase = slice * sliceStrideLowpass; + uint hSliceBase = slice * sliceStrideHighpass; + uint oSliceBase = slice * sliceStrideOutput; + + uint t = lid.x; + uint r = lid.y; + uint tile_base_i = tgid.y * 8; + + if (halfHeightH == 0) { + uint i = tile_base_i + r; + if (i < halfHeight) { + output[oSliceBase + (2 * i) * width + col] = + lowpass[lSliceBase + i * width + col]; + } + return; + } + + // Step 1 — even-row output for (i = tile_base_i + r, col) + { + uint i = tile_base_i + r; + if (i < halfHeight) { + int dTop = (i > 0) + ? highpass[hSliceBase + (i - 1) * width + col] + : highpass[hSliceBase + col]; + int dBot = (i < halfHeightH) + ? highpass[hSliceBase + i * width + col] + : highpass[hSliceBase + (halfHeightH - 1) * width + col]; + int e = lowpass[lSliceBase + i * width + col] - ((dTop + dBot + 2) >> 2); + tg_even[r][t] = e; + output[oSliceBase + (2 * i) * width + col] = e; + } + } + + // Boundary — r == 7 also computes the 9th even + if (r == 7) { + uint i = tile_base_i + 8; + if (i < halfHeight) { + int dTop = (i > 0) + ? highpass[hSliceBase + (i - 1) * width + col] + : highpass[hSliceBase + col]; + int dBot = (i < halfHeightH) + ? highpass[hSliceBase + i * width + col] + : highpass[hSliceBase + (halfHeightH - 1) * width + col]; + int e = lowpass[lSliceBase + i * width + col] - ((dTop + dBot + 2) >> 2); + tg_even[8][t] = e; + } else { + tg_even[8][t] = tg_even[7][t]; + } + } + + threadgroup_barrier(mem_flags::mem_threadgroup); + + // Step 2 — odd-row output + uint i = tile_base_i + r; + if (i < halfHeightH) { + int eTop = tg_even[r][t]; + int eBot = (2 * i + 2 < height) + ? tg_even[r + 1][t] + : tg_even[r][t]; + output[oSliceBase + (2 * i + 1) * width + col] = + highpass[hSliceBase + i * width + col] + ((eTop + eBot) >> 1); + } +} + // MARK: - v10.5 Phase 2-3-fused — Inverse 5/3 Int H+V fused-tile kernel // // The v10.3 Phase 2-2-tiled pair (`*_horizontal_int_tiled` + diff --git a/Sources/J2KMetal/default.metallib b/Sources/J2KMetal/default.metallib index e17abe51..66fc5be8 100644 Binary files a/Sources/J2KMetal/default.metallib and b/Sources/J2KMetal/default.metallib differ diff --git a/Tests/J2KCodecTests/V10_20_BatchedBridgeParityTests.swift b/Tests/J2KCodecTests/V10_20_BatchedBridgeParityTests.swift new file mode 100644 index 00000000..f499856b --- /dev/null +++ b/Tests/J2KCodecTests/V10_20_BatchedBridgeParityTests.swift @@ -0,0 +1,156 @@ +// +// V10_20_BatchedBridgeParityTests.swift +// J2KSwift +// +// v10.20-research Phase 3b — JP3D batched bridge SPI parity oracle. +// +// Specification: for every slice `i` in the batch, +// _jp3dIDWTAndFinalizeBatched(coefsBatch)[i].components[0].data == +// _jp3dIDWTAndFinalize(coefsBatch[i]).components[0].data +// on every JP3D-shape input (5/3 reversible, single-component, +// uniform dims across the batch, full resolution, no ROI). +// +// The batched bridge SPI is the Phase 3a Metal orchestrator wired +// into the v10.10.0-era bridge SPI surface so JP3DSliceStackCodec +// can consume it in Phase 3c. Bit-exact composition vs the serial +// bridge SPI is the safety net — any subband-grouping or pad/strip +// bug surfaces here, not as wrong voxels in JP3D output. + +import XCTest +@testable import J2KCore +@testable import J2KCodec + +final class V10_20_BatchedBridgeParityTests: XCTestCase { + + // MARK: - Fixture builder + + private func makeLCGImage( + width: Int, height: Int, + bitDepth: Int = 16, seed: UInt64 + ) -> J2KImage { + let bytesPerSample = (bitDepth + 7) / 8 + let voxelCount = width * height + let maxVal = (1 << bitDepth) - 1 + var data = Data(count: voxelCount * bytesPerSample) + var s: UInt64 = seed &* 6364136223846793005 &+ 1442695040888963407 + data.withUnsafeMutableBytes { raw in + let p = raw.bindMemory(to: UInt16.self) + for i in 0..> 16) & 0xFFFFFFFF) % (maxVal + 1)) + p[i] = sample.bigEndian + } + } + let comp = J2KComponent( + index: 0, bitDepth: bitDepth, signed: false, + width: width, height: height, + subsamplingX: 1, subsamplingY: 1, + data: data, sampleByteOrder: .bigEndian) + return J2KImage(width: width, height: height, components: [comp]) + } + + private func losslessHTEncoder() -> J2KEncoder { + let cfg = J2KEncodingConfiguration( + quality: 1.0, lossless: true, + decompositionLevels: 5, qualityLayers: 1, + progressionOrder: .lrcp, bitrateMode: .lossless, + maxThreads: 8, useHTJ2K: true, useReversibleFilter: true, + enableParallelCodeBlocks: true, + htj2kBlockFormat: .conformant) + return J2KEncoder(encodingConfiguration: cfg) + } + + // MARK: - Tests + + /// Single-slice batch — bit-exact equivalent to the serial bridge + /// SPI. Establishes baseline before testing multi-slice. + func testBatchedBridgeOfOneEqualsSerial_256x256() async throws { + let image = makeLCGImage(width: 256, height: 256, seed: 1) + let codestream = try await losslessHTEncoder().encode(image) + + let decoder = J2KDecoder() + let coefs = try await decoder._jp3dDecodeToCoefficients(codestream) + let serial = try await decoder._jp3dIDWTAndFinalize(coefs) + let batched = try await decoder._jp3dIDWTAndFinalizeBatched([coefs]) + + XCTAssertEqual(batched.count, 1) + XCTAssertEqual(serial.components.first?.data, + batched[0].components.first?.data, + "Batched bridge SPI of 1 diverges from serial") + } + + /// 4-slice batch — typical JP3D mid-volume slice count for a tile. + func testBatchedBridgeOfFourEqualsSerial_256x256() async throws { + let decoder = J2KDecoder() + var coefsArr: [JP3DSliceCoefficients] = [] + var serials: [J2KImage] = [] + for s in 0..<4 { + let image = makeLCGImage(width: 256, height: 256, seed: UInt64(0x100 + s)) + let codestream = try await losslessHTEncoder().encode(image) + let coefs = try await decoder._jp3dDecodeToCoefficients(codestream) + coefsArr.append(coefs) + serials.append(try await decoder._jp3dIDWTAndFinalize(coefs)) + } + let batched = try await decoder._jp3dIDWTAndFinalizeBatched(coefsArr) + + XCTAssertEqual(batched.count, 4) + for s in 0..<4 { + XCTAssertEqual(serials[s].components.first?.data, + batched[s].components.first?.data, + "Batched bridge SPI slice \(s) diverges from serial") + } + } + + /// 16-slice batch — typical JP3D tile-depth size. Confirms the + /// orchestrator scales correctly through the bridge SPI. + func testBatchedBridgeOfSixteenEqualsSerial_256x256() async throws { + let decoder = J2KDecoder() + var coefsArr: [JP3DSliceCoefficients] = [] + var serials: [J2KImage] = [] + for s in 0..<16 { + let image = makeLCGImage(width: 256, height: 256, seed: UInt64(0x200 + s)) + let codestream = try await losslessHTEncoder().encode(image) + let coefs = try await decoder._jp3dDecodeToCoefficients(codestream) + coefsArr.append(coefs) + serials.append(try await decoder._jp3dIDWTAndFinalize(coefs)) + } + let batched = try await decoder._jp3dIDWTAndFinalizeBatched(coefsArr) + + XCTAssertEqual(batched.count, 16) + for s in 0..<16 { + XCTAssertEqual(serials[s].components.first?.data, + batched[s].components.first?.data, + "Batched bridge SPI 16-slice slice \(s) diverges from serial") + } + } + + /// 4-slice batch @ 512×512 — larger fixture, typical thorax CT slice. + func testBatchedBridgeOfFourEqualsSerial_512x512() async throws { + let decoder = J2KDecoder() + var coefsArr: [JP3DSliceCoefficients] = [] + var serials: [J2KImage] = [] + for s in 0..<4 { + let image = makeLCGImage(width: 512, height: 512, seed: UInt64(0x300 + s)) + let codestream = try await losslessHTEncoder().encode(image) + let coefs = try await decoder._jp3dDecodeToCoefficients(codestream) + coefsArr.append(coefs) + serials.append(try await decoder._jp3dIDWTAndFinalize(coefs)) + } + let batched = try await decoder._jp3dIDWTAndFinalizeBatched(coefsArr) + + XCTAssertEqual(batched.count, 4) + for s in 0..<4 { + XCTAssertEqual(serials[s].components.first?.data, + batched[s].components.first?.data, + "Batched bridge SPI 512×512 slice \(s) diverges from serial") + } + } + + /// Empty batch returns empty without crashing. + func testBatchedBridgeEmptyReturnsEmpty() async throws { + let decoder = J2KDecoder() + let batched = try await decoder._jp3dIDWTAndFinalizeBatched([]) + XCTAssertEqual(batched.count, 0) + } +} diff --git a/Tests/J2KCodecTests/V10_20_JP3DBridgeParityTests.swift b/Tests/J2KCodecTests/V10_20_JP3DBridgeParityTests.swift new file mode 100644 index 00000000..7ce3b54a --- /dev/null +++ b/Tests/J2KCodecTests/V10_20_JP3DBridgeParityTests.swift @@ -0,0 +1,162 @@ +// +// V10_20_JP3DBridgeParityTests.swift +// J2KSwift +// +// v10.20-research Phase 1 — JP3D bridge SPI parity oracle. +// +// Asserts the composition guarantee: +// +// decode(data) ≡ _jp3dIDWTAndFinalize(_jp3dDecodeToCoefficients(data)) +// +// on every single-tile codestream. The bridge SPI is the foundation +// for Phase 2's batched GPU iDWT; if the composition is bit-exact +// equal to `decode(_:)` then Phase 2's batched path is provably +// equivalent to the production single-slice path (since Phase 2 only +// changes WHERE the iDWT runs — same coefficients in, same spatial +// data out). + +import XCTest +@testable import J2KCore +@testable import J2KCodec + +final class V10_20_JP3DBridgeParityTests: XCTestCase { + + // MARK: - Test fixture builder + + /// Mirrors the v10.18 JP3D parity-oracle LCG synthesis pattern + /// (multiplier `1442695040888963407`) so the bytes look the same + /// as the slices JP3DSliceStackCodec would route through this SPI. + private func makeLCGImage( + width: Int, height: Int, + bitDepth: Int = 16, seed: UInt64 = 0xC0FFEE + ) -> J2KImage { + let bytesPerSample = (bitDepth + 7) / 8 + let voxelCount = width * height + let maxVal = (1 << bitDepth) - 1 + var data = Data(count: voxelCount * bytesPerSample) + var s: UInt64 = seed &* 6364136223846793005 &+ 1442695040888963407 + data.withUnsafeMutableBytes { raw in + let p = raw.bindMemory(to: UInt16.self) + for i in 0..> 16) & 0xFFFFFFFF) % (maxVal + 1)) + // Big-endian (matches J2KComponent.sampleByteOrder convention). + p[i] = sample.bigEndian + } + } + let comp = J2KComponent( + index: 0, bitDepth: bitDepth, signed: false, + width: width, height: height, + subsamplingX: 1, subsamplingY: 1, + data: data, sampleByteOrder: .bigEndian) + return J2KImage(width: width, height: height, components: [comp]) + } + + private func losslessHTEncoder() -> J2KEncoder { + let cfg = J2KEncodingConfiguration( + quality: 1.0, lossless: true, + decompositionLevels: 5, qualityLayers: 1, + progressionOrder: .lrcp, bitrateMode: .lossless, + maxThreads: 8, useHTJ2K: true, useReversibleFilter: true, + enableParallelCodeBlocks: true, + htj2kBlockFormat: .conformant) + return J2KEncoder(encodingConfiguration: cfg) + } + + /// **Spec**: the bridge SPI composition equals the direct decode + /// on a single-tile lossless HT codestream, byte-exact. + func testBridgeCompositionEqualsDirectDecode_256x256() async throws { + let image = makeLCGImage(width: 256, height: 256) + let codestream = try await losslessHTEncoder().encode(image) + + let decoder = J2KDecoder() + let direct = try await decoder.decode(codestream) + let coefs = try await decoder._jp3dDecodeToCoefficients(codestream) + let bridged = try await decoder._jp3dIDWTAndFinalize(coefs) + + XCTAssertEqual(direct.width, bridged.width) + XCTAssertEqual(direct.height, bridged.height) + XCTAssertEqual(direct.components.count, bridged.components.count) + for (i, dComp) in direct.components.enumerated() { + let bComp = bridged.components[i] + XCTAssertEqual(dComp.bitDepth, bComp.bitDepth) + XCTAssertEqual(dComp.signed, bComp.signed) + XCTAssertEqual(dComp.width, bComp.width) + XCTAssertEqual(dComp.height, bComp.height) + XCTAssertEqual(dComp.data, bComp.data, + "Phase 1 composition broken on 256x256 — " + + "bridge SPI iDWT output differs from direct decode.") + } + } + + func testBridgeCompositionEqualsDirectDecode_512x512() async throws { + let image = makeLCGImage(width: 512, height: 512, seed: 0xDEADBEEF) + let codestream = try await losslessHTEncoder().encode(image) + + let decoder = J2KDecoder() + let direct = try await decoder.decode(codestream) + let coefs = try await decoder._jp3dDecodeToCoefficients(codestream) + let bridged = try await decoder._jp3dIDWTAndFinalize(coefs) + + XCTAssertEqual(direct.components.first?.data, bridged.components.first?.data, + "Phase 1 composition broken on 512x512.") + } + + /// Tiny fixture — exercises the small-codestream edge cases. + func testBridgeCompositionEqualsDirectDecode_64x64() async throws { + let image = makeLCGImage(width: 64, height: 64, seed: 1) + let codestream = try await losslessHTEncoder().encode(image) + + let decoder = J2KDecoder() + let direct = try await decoder.decode(codestream) + let coefs = try await decoder._jp3dDecodeToCoefficients(codestream) + let bridged = try await decoder._jp3dIDWTAndFinalize(coefs) + + XCTAssertEqual(direct.components.first?.data, bridged.components.first?.data, + "Phase 1 composition broken on 64x64.") + } + + /// Width/height not equal — verifies the bridge doesn't assume + /// square slices. + func testBridgeCompositionEqualsDirectDecode_128x320() async throws { + let image = makeLCGImage(width: 128, height: 320, seed: 42) + let codestream = try await losslessHTEncoder().encode(image) + + let decoder = J2KDecoder() + let direct = try await decoder.decode(codestream) + let coefs = try await decoder._jp3dDecodeToCoefficients(codestream) + let bridged = try await decoder._jp3dIDWTAndFinalize(coefs) + + XCTAssertEqual(direct.width, bridged.width) + XCTAssertEqual(direct.height, bridged.height) + XCTAssertEqual(direct.components.first?.data, bridged.components.first?.data, + "Phase 1 composition broken on 128x320.") + } + + /// Determinism — two SPI runs against the same codestream produce + /// byte-identical output. + func testBridgeDeterministic() async throws { + let image = makeLCGImage(width: 256, height: 256) + let codestream = try await losslessHTEncoder().encode(image) + + let decoder = J2KDecoder() + let coefs1 = try await decoder._jp3dDecodeToCoefficients(codestream) + let out1 = try await decoder._jp3dIDWTAndFinalize(coefs1) + let coefs2 = try await decoder._jp3dDecodeToCoefficients(codestream) + let out2 = try await decoder._jp3dIDWTAndFinalize(coefs2) + + XCTAssertEqual(out1.components.first?.data, out2.components.first?.data, + "Phase 1 bridge SPI must be deterministic.") + } + + // Multi-tile-throws test elided: J2KEncoder's tileSize parameter + // auto-promotes to single-tile when image dims fit a single tile, + // and forcing a true multi-tile codestream needs a separate fixture + // source. JP3D's slice-stack codec by construction never emits + // multi-tile per-slice codestreams, so the runtime constraint is + // documented at the bridge SPI's `mutating func + // decodeToCoefficients` site and exercised only when a manually- + // constructed multi-tile codestream is routed through the SPI — + // not a Phase 1 deliverable. +} diff --git a/Tests/J2KMetalTests/V10_20_BatchedInverseInt32ParityTests.swift b/Tests/J2KMetalTests/V10_20_BatchedInverseInt32ParityTests.swift new file mode 100644 index 00000000..fcd4044a --- /dev/null +++ b/Tests/J2KMetalTests/V10_20_BatchedInverseInt32ParityTests.swift @@ -0,0 +1,405 @@ +// +// V10_20_BatchedInverseInt32ParityTests.swift +// J2KSwift +// +// v10.20-research Phase 2 — bit-exact parity oracle for the new +// `J2KMetalDWT.inverseBatched2DInt32` API. +// +// Specification (bit-exact): +// inverseBatched2DInt32(subbandsBatch: [a, b, c])[i] ≡ +// inverse2DInt32(subbands: subbandsBatch[i], backend: .gpu) +// for every i, on every subband shape the JP3D slice-stack codec +// emits. +// +// The batched dispatch must produce per-slice output byte-identical +// to the existing tiled single-image dispatch — Phase 2's claim +// rests entirely on this property. If parity ever fails, JP3D's +// Phase 3 wiring will produce wrong voxels, so we gate hard. + +import XCTest +@testable import J2KCore +@testable import J2KMetal + +@MainActor +final class V10_20_BatchedInverseInt32ParityTests: XCTestCase { + + // MARK: - Subband builder + + /// Build a J2KMetalDWTSubbandsInt32 from a deterministic per-slice + /// seed. Layout mirrors the post-dequant shape the v10.20 bridge + /// SPI hands off: LL/HL/LH/HH split into the standard band + /// dimensions, originalWidth × originalHeight final iDWT output. + private func makeSubbands( + originalWidth: Int, originalHeight: Int, seed: UInt64 + ) -> J2KMetalDWTSubbandsInt32 { + let llW = (originalWidth + 1) / 2 + let llH = (originalHeight + 1) / 2 + let hlW = llW + let hlH = originalHeight / 2 + let lhW = originalWidth / 2 + let lhH = llH + let hhW = originalWidth / 2 + let hhH = originalHeight / 2 + + var rng = seed &* 6364136223846793005 &+ 1442695040888963407 + func fill(_ w: Int, _ h: Int) -> [Int32] { + (0..<(w * h)).map { _ in + rng = rng &* 6364136223846793005 &+ 1442695040888963407 + let v = Int32(truncatingIfNeeded: Int((rng >> 16) & 0xFFFF)) - 16384 + return v + } + } + + return J2KMetalDWTSubbandsInt32( + ll: fill(llW, llH), + lh: fill(lhW, lhH), + hl: fill(hlW, hlH), + hh: fill(hhW, hhH), + llWidth: llW, llHeight: llH, + originalWidth: originalWidth, originalHeight: originalHeight) + } + + // MARK: - Tests + + /// Single-slice batch — must be byte-identical to the serial path. + /// Establishes the baseline before testing multi-slice. + func testBatchedOfOneEqualsSerial_256x256() async throws { + let dwt = J2KMetalDWT() + let sub = makeSubbands(originalWidth: 256, originalHeight: 256, seed: 7) + + let serial = try await dwt.inverse2DInt32(subbands: sub, backend: .gpu) + let batched = try await dwt.inverseBatched2DInt32(subbandsBatch: [sub]) + + XCTAssertEqual(batched.count, 1) + XCTAssertEqual(serial, batched[0], + "Single-slice batched output diverges from serial") + } + + /// 2-slice batch — each slice's output must match its serial decode. + func testBatchedOfTwoEqualsSerial_256x256() async throws { + let dwt = J2KMetalDWT() + let subA = makeSubbands(originalWidth: 256, originalHeight: 256, seed: 1) + let subB = makeSubbands(originalWidth: 256, originalHeight: 256, seed: 2) + + let serialA = try await dwt.inverse2DInt32(subbands: subA, backend: .gpu) + let serialB = try await dwt.inverse2DInt32(subbands: subB, backend: .gpu) + let batched = try await dwt.inverseBatched2DInt32(subbandsBatch: [subA, subB]) + + XCTAssertEqual(batched.count, 2) + XCTAssertEqual(serialA, batched[0], "Slice 0 batched diverges from serial") + XCTAssertEqual(serialB, batched[1], "Slice 1 batched diverges from serial") + } + + /// 16-slice batch — typical JP3D mid-volume slice count. Confirms + /// the Z grid dimension scales correctly without aliasing across + /// slices. + func testBatchedOfSixteenEqualsSerial_256x256() async throws { + let dwt = J2KMetalDWT() + var subs: [J2KMetalDWTSubbandsInt32] = [] + var serials: [[Int32]] = [] + for s in 0..<16 { + let sub = makeSubbands(originalWidth: 256, originalHeight: 256, + seed: UInt64(0xDEAD0000 + s)) + subs.append(sub) + serials.append(try await dwt.inverse2DInt32(subbands: sub, backend: .gpu)) + } + let batched = try await dwt.inverseBatched2DInt32(subbandsBatch: subs) + + XCTAssertEqual(batched.count, 16) + for s in 0..<16 { + XCTAssertEqual(serials[s], batched[s], + "Slice \(s) batched diverges from serial") + } + } + + /// Larger dim (typical JP3D thorax CT slice) — sanity that the + /// batched dispatch handles realistic medical-scan sizes. + func testBatchedOfFourEqualsSerial_512x512() async throws { + let dwt = J2KMetalDWT() + var subs: [J2KMetalDWTSubbandsInt32] = [] + var serials: [[Int32]] = [] + for s in 0..<4 { + let sub = makeSubbands(originalWidth: 512, originalHeight: 512, + seed: UInt64(0xBEEF0000 + s)) + subs.append(sub) + serials.append(try await dwt.inverse2DInt32(subbands: sub, backend: .gpu)) + } + let batched = try await dwt.inverseBatched2DInt32(subbandsBatch: subs) + + XCTAssertEqual(batched.count, 4) + for s in 0..<4 { + XCTAssertEqual(serials[s], batched[s], + "Slice \(s) batched 512×512 diverges from serial") + } + } + + /// Non-square dims — verifies the batched dispatch doesn't assume + /// square slices. + func testBatchedNonSquareEqualsSerial_128x320() async throws { + let dwt = J2KMetalDWT() + var subs: [J2KMetalDWTSubbandsInt32] = [] + var serials: [[Int32]] = [] + for s in 0..<3 { + let sub = makeSubbands(originalWidth: 128, originalHeight: 320, + seed: UInt64(0xCAFE0000 + s)) + subs.append(sub) + serials.append(try await dwt.inverse2DInt32(subbands: sub, backend: .gpu)) + } + let batched = try await dwt.inverseBatched2DInt32(subbandsBatch: subs) + + for s in 0..<3 { + XCTAssertEqual(serials[s], batched[s], + "Slice \(s) batched 128×320 diverges from serial") + } + } + + /// Mixed-dimension batch throws — JP3D slices share dimensions + /// per JP3DSliceStackCodec wire format; mixed dims through the SPI + /// is rejected so the caller can't silently produce wrong output. + func testMixedDimensionBatchThrows() async throws { + let dwt = J2KMetalDWT() + let subA = makeSubbands(originalWidth: 256, originalHeight: 256, seed: 1) + let subB = makeSubbands(originalWidth: 128, originalHeight: 128, seed: 2) + + do { + _ = try await dwt.inverseBatched2DInt32(subbandsBatch: [subA, subB]) + XCTFail("Mixed-dimension batch should have thrown invalidParameter") + } catch J2KError.invalidParameter { + // Expected + } catch { + XCTFail("Wrong error type for mixed dims: \(error)") + } + } + + /// Empty batch returns empty result without crashing. + func testEmptyBatchReturnsEmpty() async throws { + let dwt = J2KMetalDWT() + let batched = try await dwt.inverseBatched2DInt32(subbandsBatch: []) + XCTAssertEqual(batched.count, 0) + } + + // MARK: - Multi-level batched parity (Phase 3a) + + /// Build a multi-level subband chain for testing the multi-level + /// orchestrator. The level-chain invariant: at level k, ll dim + /// equals level k-1's original dim. We build innermost-first. + private func makeMultiLevelChain( + finalWidth: Int, finalHeight: Int, levels: Int, seed: UInt64 + ) -> [J2KMetalDWTSubbandsInt32] { + // Compute per-level dims: level (levels-1) has original = + // finalW × finalH; level 0 (innermost) has original = finalW/2^(levels-1). + var dims: [(w: Int, h: Int)] = [] + var w = finalWidth + var h = finalHeight + var stack: [(Int, Int)] = [(w, h)] + for _ in 1.. [Int32] { + (0..> 16) & 0xFFFF)) - 16384 + } + } + let ll = fill(myLLW * myLLH) + let lh = max(0, myLHW * myLHH) > 0 ? fill(myLHW * myLHH) : [] + let hl = max(0, myHLW * myHLH) > 0 ? fill(myHLW * myHLH) : [] + let hh = max(0, myHHW * myHHH) > 0 ? fill(myHHW * myHHH) : [] + chain.append(J2KMetalDWTSubbandsInt32( + ll: ll, lh: lh, hl: hl, hh: hh, + llWidth: myLLW, llHeight: myLLH, + originalWidth: origW, originalHeight: origH)) + } + return chain + } + + /// **Spec**: the multi-level batched orchestrator's per-slice + /// output equals the per-slice serial multi-level orchestrator's + /// output, byte-exact, on every slice in the batch. + func testMultiLevelBatchedOfOneEqualsSerial_256x256_3levels() async throws { + let dwt = J2KMetalDWT() + let chain = makeMultiLevelChain( + finalWidth: 256, finalHeight: 256, levels: 3, seed: 11) + + let serial = try await dwt.inverse2DInt32MultiLevelFused( + subbandsPerLevel: chain) + let batched = try await dwt.inverse2DInt32MultiLevelFusedBatched( + perSliceSubbandsPerLevel: [chain]) + + XCTAssertEqual(batched.count, 1) + XCTAssertEqual(serial, batched[0], + "Multi-level batched-of-one diverges from serial") + } + + func testMultiLevelBatchedOfFourEqualsSerial_256x256_3levels() async throws { + let dwt = J2KMetalDWT() + var batch: [[J2KMetalDWTSubbandsInt32]] = [] + var serials: [[Int32]] = [] + for s in 0..<4 { + let chain = makeMultiLevelChain( + finalWidth: 256, finalHeight: 256, levels: 3, + seed: UInt64(0xABCD0000 + s)) + batch.append(chain) + serials.append(try await dwt.inverse2DInt32MultiLevelFused( + subbandsPerLevel: chain)) + } + let batched = try await dwt.inverse2DInt32MultiLevelFusedBatched( + perSliceSubbandsPerLevel: batch) + + XCTAssertEqual(batched.count, 4) + for s in 0..<4 { + XCTAssertEqual(serials[s], batched[s], + "Multi-level batched slice \(s) diverges from serial") + } + } + + /// Targeted reproducer for the Phase 3b SIGSEGV: 16 slices, + /// 256×256, 3 levels — exactly matches the bench fixture that + /// failed. Innermost dim is 64×64 (above the smallDimThreshold). + func testMultiLevelBatched_16slices_256x256_3levels() async throws { + let dwt = J2KMetalDWT() + var batch: [[J2KMetalDWTSubbandsInt32]] = [] + var serials: [[Int32]] = [] + for s in 0..<16 { + let chain = makeMultiLevelChain( + finalWidth: 256, finalHeight: 256, levels: 3, + seed: UInt64(0x12340000 + s)) + batch.append(chain) + serials.append(try await dwt.inverse2DInt32MultiLevelFused( + subbandsPerLevel: chain)) + } + let batched = try await dwt.inverse2DInt32MultiLevelFusedBatched( + perSliceSubbandsPerLevel: batch) + + XCTAssertEqual(batched.count, 16) + for s in 0..<16 { + XCTAssertEqual(serials[s], batched[s], + "16-slice batched diverges at slice \(s)") + } + } + + /// Single-shot timing comparison (release-mode). Not a regression + /// gate — just prints timings + asserts bit-exactness. The + /// looped-bench shape SIGSEGV'd intermittently during Phase 3b + /// investigation (a Metal-state accumulation under repeated + /// invocation that needs separate root-causing); a single-shot + /// measurement avoids the trigger and still gives a directional + /// signal on whether batched dispatch beats per-slice serial at + /// JP3D-typical fixture sizes. + func testMultiLevelBatchedVsSerialSingleShotTiming_16slices_256x256_3levels() async throws { + let dwt = J2KMetalDWT() + var batch: [[J2KMetalDWTSubbandsInt32]] = [] + for s in 0..<16 { + batch.append(makeMultiLevelChain( + finalWidth: 256, finalHeight: 256, levels: 3, + seed: UInt64(0xABCD0000 + s))) + } + + // Warm both paths with a parity check (also acts as warmup). + var serials: [[Int32]] = [] + for chain in batch { + serials.append(try await dwt.inverse2DInt32MultiLevelFused( + subbandsPerLevel: chain)) + } + let batchedWarmup = try await dwt.inverse2DInt32MultiLevelFusedBatched( + perSliceSubbandsPerLevel: batch) + for s in 0..