swh:1:snp:70f530b74f5be73cfb71c212c9e3317ce44c1ebc
- HEAD
- refs/heads/Halide_unsharp
- refs/heads/abadams/aggressive_is_single_point
- refs/heads/abadams/align_strided_const_loads
- refs/heads/abadams/alloca
- refs/heads/abadams/atomic_parallel_compiled_in
- refs/heads/abadams/atomic_vector_non_recursive
- refs/heads/abadams/averaging_tree
- refs/heads/abadams/avoid_name_mangling_in_cross_module_dependencies
- refs/heads/abadams/better_absd
- refs/heads/abadams/better_codegen_for_non_const_ramps
- refs/heads/abadams/bgu_cholesky
- refs/heads/abadams/braces_around_statements
- refs/heads/abadams/cache_tighten_producer_consumer_nodes
- refs/heads/abadams/check_reorder_dups
- refs/heads/abadams/clarify_broadcast_shuffle
- refs/heads/abadams/compositing_app
- refs/heads/abadams/cond_wait_spin
- refs/heads/abadams/cse_in_unroll_split_tuples
- refs/heads/abadams/custom_cuda_context
- refs/heads/abadams/custom_cuda_context_2
- refs/heads/abadams/custom_cuda_context_3
- refs/heads/abadams/d3d12abi
- refs/heads/abadams/deflake_mullapudi_reorder
- refs/heads/abadams/delete_prepare_for_early_exit
- refs/heads/abadams/depthwise_separable_conv
- refs/heads/abadams/diagnose_boundary_condition_failure
- refs/heads/abadams/disable_onnx_app_on_mac
- refs/heads/abadams/divide_using_pavgw
- refs/heads/abadams/dont_link_to_cudart
- refs/heads/abadams/dont_reinterpret_concat
- refs/heads/abadams/early_out
- refs/heads/abadams/enable_f16c
- refs/heads/abadams/extract_concat_bits
- refs/heads/abadams/fast_integer_divide_round_to_zero
- refs/heads/abadams/faster_runtime_integer_division
- refs/heads/abadams/faster_substitute_facts
- refs/heads/abadams/faster_unroll
- refs/heads/abadams/fix-arm-seg2
- refs/heads/abadams/fix_4211
- refs/heads/abadams/fix_5323
- refs/heads/abadams/fix_5329
- refs/heads/abadams/fix_5889
- refs/heads/abadams/fix_6984
- refs/heads/abadams/fix_7229
- refs/heads/abadams/fix_7260
- refs/heads/abadams/fix_7365
- refs/heads/abadams/fix_7374
- refs/heads/abadams/fix_7504
- refs/heads/abadams/fix_7514
- refs/heads/abadams/fix_7531
- refs/heads/abadams/fix_7584
- refs/heads/abadams/fix_7584_v2
- refs/heads/abadams/fix_7742
- refs/heads/abadams/fix_7756
- refs/heads/abadams/fix_7761
- refs/heads/abadams/fix_7768
- refs/heads/abadams/fix_7786
- refs/heads/abadams/fix_7810
- refs/heads/abadams/fix_7811
- refs/heads/abadams/fix_7815
- refs/heads/abadams/fix_7867
- refs/heads/abadams/fix_7871
- refs/heads/abadams/fix_7872
- refs/heads/abadams/fix_7873
- refs/heads/abadams/fix_7888
- refs/heads/abadams/fix_7890
- refs/heads/abadams/fix_7891
- refs/heads/abadams/fix_7892
- refs/heads/abadams/fix_7893
- refs/heads/abadams/fix_7906
- refs/heads/abadams/fix_7909
- refs/heads/abadams/fix_7968
- refs/heads/abadams/fix_8038
- refs/heads/abadams/fix_8054
- refs/heads/abadams/fix_8170
- refs/heads/abadams/fix_8184
- refs/heads/abadams/fix_arm_fcvtmp
- refs/heads/abadams/fix_autoschedule_feature_transposition
- refs/heads/abadams/fix_cse_name_collisions
- refs/heads/abadams/fix_cuda_mat_mul_assert
- refs/heads/abadams/fix_deinterleave_bug
- refs/heads/abadams/fix_deinterleave_for_reinterpret
- refs/heads/abadams/fix_div_round_to_zero
- refs/heads/abadams/fix_fft_compile_time_regression
- refs/heads/abadams/fix_generate_output_snippets
- refs/heads/abadams/fix_if_nesting_condition
- refs/heads/abadams/fix_leaks_in_memoize_test
- refs/heads/abadams/fix_lgtm_warnings
- refs/heads/abadams/fix_links_to_master
- refs/heads/abadams/fix_load_of_broadcast
- refs/heads/abadams/fix_lossless_cast_of_sub
- refs/heads/abadams/fix_onnx_app
- refs/heads/abadams/fix_pointless_lower_condition
- refs/heads/abadams/fix_potential_gpu_deadlock
- refs/heads/abadams/fix_realize_condition_depends_on_tuple
- refs/heads/abadams/fix_reduce_expr_modulo_of_vector
- refs/heads/abadams/fix_riscv_vx_vi
- refs/heads/abadams/fix_round
- refs/heads/abadams/fix_stencil_chain_gpu_schedule
- refs/heads/abadams/fix_track_bounds_intervals
- refs/heads/abadams/fix_tutorial_2
- refs/heads/abadams/fix_ub_in_lower_rounding_shift_right
- refs/heads/abadams/forward_partition_methods
- refs/heads/abadams/fully_fused_depthwise_separable_conv
- refs/heads/abadams/fuzz_sliding_window
- refs/heads/abadams/gaussian_blur_app
- refs/heads/abadams/generator_infinite_default_timeout
- refs/heads/abadams/gpu_autoscheduler_parallel_random_probes
- refs/heads/abadams/include_riscv_in_readme
- refs/heads/abadams/interleave_nested_vector
- refs/heads/abadams/ir_match_by_ref
- refs/heads/abadams/lerp_plus_cast
- refs/heads/abadams/local_laplacian_code_size
- refs/heads/abadams/lower_halving_sub
- refs/heads/abadams/lower_rounding_shift_right
- refs/heads/abadams/mac-arm-fixes
- refs/heads/abadams/make_fast_inverse_test_throughput_limited
- refs/heads/abadams/makefile_serialization_support
- refs/heads/abadams/mismatched_new_delete
- refs/heads/abadams/mixed_sign_mul_shift_right
- refs/heads/abadams/mixed_width_mul_shift_right
- refs/heads/abadams/multiple_scatter
- refs/heads/abadams/mux_intrinsic
- refs/heads/abadams/name_helpers
- refs/heads/abadams/narrow_predicates
- refs/heads/abadams/nested_vectorization_compile_time_regression_fix
- refs/heads/abadams/nested_vectorization_tweaks
- refs/heads/abadams/parallel_simd_op_check
- refs/heads/abadams/per_instance_profiling
- refs/heads/abadams/precompute_shared_mem_size
- refs/heads/abadams/prefer_no_gather
- refs/heads/abadams/print_uncaught_exception
- refs/heads/abadams/promote_fixed_point_intrinsics
- refs/heads/abadams/psabdw
- refs/heads/abadams/random_pipelines
- refs/heads/abadams/rationalize_gpu_for_loop_names
- refs/heads/abadams/reenable_unscheduled_stage_warning
- refs/heads/abadams/refactor_constant_interval
- refs/heads/abadams/reinterpret_vector
- refs/heads/abadams/remove_arch_os_for_shaders
- refs/heads/abadams/remove_bad_pruning
- refs/heads/abadams/remove_parameter_self_references
- refs/heads/abadams/remove_readnone_on_functions
- refs/heads/abadams/remove_use_of_python_config_in_onnx_makefile
- refs/heads/abadams/reschedule_bgu
- refs/heads/abadams/reschedule_bilateral_grid
- refs/heads/abadams/rewrite_atomic_pass
- refs/heads/abadams/rewrite_ir_equality
- refs/heads/abadams/rounding_shift_right_use_average
- refs/heads/abadams/rungenmain_error
- refs/heads/abadams/sampling_profiler_overhead_v2
- refs/heads/abadams/scope_improvements
- refs/heads/abadams/simpler_broadcasts
- refs/heads/abadams/simplify_correlated_pyramid
- refs/heads/abadams/siotas_20
- refs/heads/abadams/sioutas_20
- refs/heads/abadams/slide_over_split_loop
- refs/heads/abadams/sorting_network_working_branch
- refs/heads/abadams/stable_topological_order
- refs/heads/abadams/string_view
- refs/heads/abadams/strip_asserts_last
- refs/heads/abadams/switch_stmt
- refs/heads/abadams/target_specific_lerp
- refs/heads/abadams/time_lowering_passes
- refs/heads/abadams/track_failedness_through_solver_lets
- refs/heads/abadams/turn_off_slp_vectorization_for_avx512
- refs/heads/abadams/tweak_unpack_buffers
- refs/heads/abadams/undo_pointless_widening
- refs/heads/abadams/unordered_blocks
- refs/heads/abadams/unsigned_demosaic
- refs/heads/abadams/update_makefile_for_llvm_19
- refs/heads/abadams/use_arm_for_runtime_triple
- refs/heads/abadams/use_pmaddubsw_for_downsample
- refs/heads/abadams/validate_gpu_schedules
- refs/heads/abadams/vector_reduce_hexagon_predicate
- refs/heads/abadams/vector_scan
- refs/heads/abadams/vst_type_fix
- refs/heads/abadams/widening_let_bug
- refs/heads/abadams/x86_avg
- refs/heads/abadams/zen4
- refs/heads/adadams/profile_allocator
- refs/heads/add_image_checks_after_bounds_inference_plus_new_rules
- refs/heads/add_outermost_to_extern
- refs/heads/add_vectorization_to_search_space
- refs/heads/aelphy/feature_cadence_changes
- refs/heads/aelphy/float_extracts
- refs/heads/align_loads_comment_fix
- refs/heads/alina-strided-store
- refs/heads/another_buffer_copy_fix
- refs/heads/arm_sve_redux
- refs/heads/ataei-block_asserts-codegen
- refs/heads/ataei-debug_info
- refs/heads/ataei-fix-pow
- refs/heads/ataei-gen_str_param
- refs/heads/ataei-implicit_lhs_vars
- refs/heads/ataei-onnx
- refs/heads/ataei-onnx_converter_update
- refs/heads/ataei-onnx_pybind
- refs/heads/ataei-resnet50_benchmarks
- refs/heads/ataei-standalone_autoscheduler
- refs/heads/ataei_lots_of_inputs
- refs/heads/auto_sched_benchmarks
- refs/heads/auto_sched_estimates
- refs/heads/auto_sched_inline
- refs/heads/auto_sched_test_notparallel
- refs/heads/autoschedule_top_down
- refs/heads/autoschedule_with_convnet
- refs/heads/autoscheduler_scalar_imageparam_fix
- refs/heads/backports/10.x
- refs/heads/backports/11.x
- refs/heads/backports/12.x
- refs/heads/backports/13.x
- refs/heads/balance_expressions
- refs/heads/bazel
- refs/heads/benchmarks
- refs/heads/blaze
- refs/heads/bounds_buffer_lets_fix
- refs/heads/bounds_correct_vs_bounds_loaded_reduced
- refs/heads/buffer_device_api_target
- refs/heads/bug_device_free
- refs/heads/bug_inline_unbounded
- refs/heads/build/fix-xcode-2
- refs/heads/build/manylinux-fixes
- refs/heads/circ_buffer
- refs/heads/cmake-no-runtime-debug-symbols
- refs/heads/cmake/asan
- refs/heads/cmake/deps-cleanup
- refs/heads/cmake/find-modules
- refs/heads/cmake/spirv
- refs/heads/cmake_wasm_features
- refs/heads/compute_at_guard_with_if_goes_on_stack
- refs/heads/compute_with_at
- refs/heads/compute_with_check
- refs/heads/compute_with_excessive_bounds
- refs/heads/compute_with_inlined
- refs/heads/compute_with_remove_is_right_level
- refs/heads/cpack/nuget
- refs/heads/ctest/wrappers
- refs/heads/cuda-constant
- refs/heads/d3d12-allocation-cache
- refs/heads/deferred_cse_after_inlining
- refs/heads/destructor_calls_deinit
- refs/heads/dg/deserialize_unmapped_objects
- refs/heads/dg/fix_vulkan_codegen_bool_conversion
- refs/heads/dg/vulkan_conform_api
- refs/heads/dg/vulkan_region_allocator_fixes
- refs/heads/dgerstmann/fix-vulkan-memory-config-init
- refs/heads/disable_acquire_release_test_vulkan
- refs/heads/distinct_wrapper_names
- refs/heads/dkg/6863_asan_fixes
- refs/heads/dkg/vulkan
- refs/heads/dpalermo_dmabuf
- refs/heads/dpalermo_dmabuf_libion
- refs/heads/dpalermo_hexagon_remote_202003
- refs/heads/dpalermo_sdk4_2_0_2
- refs/heads/ds/buffer-get-pure
- refs/heads/ds/opt-tile-size
- refs/heads/ds/tail-none
- refs/heads/ds/while
- refs/heads/dsharletg/bitwise-intrinsics
- refs/heads/dsharletg/find-vector-reduce
- refs/heads/dsharletg/jit-optimization
- refs/heads/dsharletg/memcpy-copy_from
- refs/heads/dsharletg/pattern-headroom
- refs/heads/dsharletg/refactor-host-alignment
- refs/heads/dsharletg/runtime-size
- refs/heads/dsharletg/simplify-abs
- refs/heads/dsharletg/simplify-type-bounds
- refs/heads/dsharletg/specialize-bounds
- refs/heads/dsharletg/upsample-channels
- refs/heads/empty_prefetch
- refs/heads/emscripten_vector_fix
- refs/heads/export_all-wsmoses
- refs/heads/expr_auto_sched
- refs/heads/extern_bugs
- refs/heads/extern_host_alloc
- refs/heads/factor_parallel_codegen_hack
- refs/heads/fast_sync_tsan
- refs/heads/faster_integer_division
- refs/heads/feature/apps-external
- refs/heads/feature/cmake-presets
- refs/heads/feature/convert
- refs/heads/feature/f16_interleave
- refs/heads/feature/gather_load_q7
- refs/heads/feature/llvm-codemodel
- refs/heads/feature/load_predicated
- refs/heads/feature/luma_regression
- refs/heads/feature/maintanence
- refs/heads/feature/reinterprets
- refs/heads/feature/tcm_bump_allocator
- refs/heads/feature/xtensa_fix_interleave_q8
- refs/heads/feature/xtensa_q8_tests
- refs/heads/find_intrinsics_issue
- refs/heads/find_intrinsics_widening_lets
- refs/heads/fix-floated-pure-stage
- refs/heads/fix-race-condition
- refs/heads/fix_hexagon_alignment
- refs/heads/fix_hvx_intrinsics
- refs/heads/fix_prefetch_test
- refs/heads/fix_windows_vs15_build
- refs/heads/fixed_length_vectors
- refs/heads/fixed_point_local_laplac
- refs/heads/gemmlowp
- refs/heads/generate
- refs/heads/gha/pip
- refs/heads/gpu_canon_fix
- refs/heads/halide_ir_flatbuffer
- refs/heads/hex_dma2_async
- refs/heads/hexagon_le_runtime
- refs/heads/hexagon_priority
- refs/heads/hexagon_setpriority
- refs/heads/hexagon_strided_pred_load
- refs/heads/hexagon_sysmon_markers
- refs/heads/imaging-synthesis
- refs/heads/includes_fix
- refs/heads/ios_fast_sync_fix
- refs/heads/jia-kai-fix-runtime-cuda-init
- refs/heads/kamil-openglcompute-infinity
- refs/heads/kamil/name_pthread_workers
- refs/heads/kp_bit_shift
- refs/heads/line_buffer
- refs/heads/loop_carry_not_working
- refs/heads/lower_on_huge_stack
- refs/heads/main
- refs/heads/master
- refs/heads/memoize_with_extents
- refs/heads/metal_float16
- refs/heads/metaprogrammed_simplifier_mod
- refs/heads/mohamedadaly-vmlal
- refs/heads/more_powerful_sliding
- refs/heads/new_autoschedule_with_new_simplifier_arm_worker_branch
- refs/heads/new_autoscheduler
- refs/heads/new_simplifier_rule_testing
- refs/heads/newer_ion_ioctl
- refs/heads/no_bounds_query_when_bounds_used
- refs/heads/opengl_compute_buffer_types_fix
- refs/heads/openglcompute_reuse_shared_allocations
- refs/heads/optmize_reorder
- refs/heads/par_for_opt
- refs/heads/pdb/fix_7806
- refs/heads/pdb/hexagon_remote_cmake
- refs/heads/pdb_add_libcpp_makefile_inc
- refs/heads/pdb_eliminate_interleaves_test
- refs/heads/pdb_fix_clang_build
- refs/heads/pdb_fix_install_qc
- refs/heads/pdb_fix_loop_carry
- refs/heads/pdb_fix_simd_op_check_hvx
- refs/heads/pdb_mul_div_mod_multi_thread
- refs/heads/pdb_remove_hvx_v64
- refs/heads/perform_inline_with_order
- refs/heads/pr/2572
- refs/heads/pr/2676
- refs/heads/pr/2975
- refs/heads/pr/3017
- refs/heads/pr/3081
- refs/heads/pr/3387
- refs/heads/pr/3939
- refs/heads/pr/3960
- refs/heads/pr/4380
- refs/heads/pr/4414
- refs/heads/pr/5331
- refs/heads/pr/5438
- refs/heads/pr/5455
- refs/heads/pr/5758_2
- refs/heads/predicated_vector
- refs/heads/prefetch_specialize
- refs/heads/print_schedule
- refs/heads/profile_hardware_counters
- refs/heads/random-pipelines
- refs/heads/rdom_with_pure_vars
- refs/heads/readme-fix-gcd
- refs/heads/realization_order
- refs/heads/refactor_module
- refs/heads/register_promotion
- refs/heads/release/10.x
- refs/heads/release/11.x
- refs/heads/release/12.x
- refs/heads/release/13.x
- refs/heads/release/14.x
- refs/heads/release/15.x
- refs/heads/release/16.x
- refs/heads/release/17.x
- refs/heads/release/8.x
- refs/heads/remove_max_on_fuse_factor
- refs/heads/reorder_rvar
- refs/heads/reset_unique_counter
- refs/heads/revert-3612-ataei-speedup_compiletime
- refs/heads/revert-7009-rootjalex/distribute-w_shl
- refs/heads/revert-7601-compile_hexagon_remote
- refs/heads/riscv_update
- refs/heads/rl_simplifier_rules
- refs/heads/rootjalex/add_simpl_rules
- refs/heads/rootjalex/arm-optimize
- refs/heads/rootjalex/autoscheduler_mcts
- refs/heads/rootjalex/bounds-rewriter
- refs/heads/rootjalex/bounds_synthesis
- refs/heads/rootjalex/cbounds
- refs/heads/rootjalex/cbounds_predicated
- refs/heads/rootjalex/fix-sat-overflow
- refs/heads/rootjalex/fix_estimate_issue
- refs/heads/rootjalex/fix_failed_unrolls
- refs/heads/rootjalex/gsoc_codegen
- refs/heads/rootjalex/improve_cbounds_fixed
- refs/heads/rootjalex/improve_constant_bounds
- refs/heads/rootjalex/pitchfork-arm
- refs/heads/rootjalex/reinterpret-simplify
- refs/heads/rootjalex/rts
- refs/heads/rootjalex/super_simplify_bounds
- refs/heads/rootjalex/test_cbounds_fixed
- refs/heads/rootjalex/test_constant_bounds
- refs/heads/rootjalex/trs-codegen
- refs/heads/rootjalex/trs-codegen-cross
- refs/heads/rootjalex/trs-merge
- refs/heads/rootjalex/uint32-int32-cast
- refs/heads/rootjalex/x86-hadds
- refs/heads/rootjalex/x86-optimize
- refs/heads/rootjalex/x86-optimize-test
- refs/heads/rootjalex/x86-sat
- refs/heads/rootjalex/x86-test
- refs/heads/rule_removal_experiments
- refs/heads/schedule-output-storage
- refs/heads/separate_bounds_query_entrypoint
- refs/heads/shallow
- refs/heads/shift_amount_type_change
- refs/heads/shoaibkamil/cmake-without-arm
- refs/heads/shoaibkamil/correct_memory_fences
- refs/heads/shoaibkamil/d3d-fixes
- refs/heads/shoaibkamil/deprecate_openglcompute
- refs/heads/shoaibkamil/json
- refs/heads/shoaibkamil/llvm_clone_tag
- refs/heads/shoaibkamil/minor-vcpkg-doc-change
- refs/heads/shoaibkamil/opengl_compute_tests
- refs/heads/shoaibkamil/performance_tests_as_generators
- refs/heads/shoaibkamil/rule_removal_experiments
- refs/heads/shoaibkamil/super_simplify_with_interpreter
- refs/heads/shoaibkamil/windows-arm-fix-attributes
- refs/heads/sim_shlib_addr_print
- refs/heads/simplify-nested-broadcasts
- refs/heads/simplify-vectorreduce-shuffles2
- refs/heads/simplify_mod
- refs/heads/sioutas_2020
- refs/heads/sioutas_2020_autoscheduler
- refs/heads/slomp/gpu-codegen-profiling
- refs/heads/slomp/msvc-static-analysis
- refs/heads/solve_div
- refs/heads/solve_div_master
- refs/heads/solve_div_simplifier_test
- refs/heads/sr/python-late-binding-defaults
- refs/heads/srj-aaa
- refs/heads/srj-alloc
- refs/heads/srj-alloca
- refs/heads/srj-appmake2
- refs/heads/srj-armv83a
- refs/heads/srj-aslog
- refs/heads/srj-assert
- refs/heads/srj-assoc
- refs/heads/srj-auto-multi
- refs/heads/srj-auto-multi2
- refs/heads/srj-auto_schedule_mat_mul
- refs/heads/srj-autosched
- refs/heads/srj-b2cpphide
- refs/heads/srj-barr
- refs/heads/srj-bits
- refs/heads/srj-blacklist
- refs/heads/srj-bounds
- refs/heads/srj-bufcalltype
- refs/heads/srj-bufcallwrap
- refs/heads/srj-bufcallwrap2
- refs/heads/srj-buffer
- refs/heads/srj-bv
- refs/heads/srj-classic-autotune
- refs/heads/srj-clean
- refs/heads/srj-constcall
- refs/heads/srj-crosscompile
- refs/heads/srj-ctlz
- refs/heads/srj-cvec-patch
- refs/heads/srj-dag
- refs/heads/srj-debug-to-file
- refs/heads/srj-deir
- refs/heads/srj-f16
- refs/heads/srj-fp16
- refs/heads/srj-fsch
- refs/heads/srj-fthru
- refs/heads/srj-g2
- refs/heads/srj-g3
- refs/heads/srj-gha-test-fixes
- refs/heads/srj-hidden
- refs/heads/srj-hide2
- refs/heads/srj-hvx
- refs/heads/srj-hvx-bug
- refs/heads/srj-hvx-codegen-bug
- refs/heads/srj-hvx-nocopy
- refs/heads/srj-hvxshift
- refs/heads/srj-iib
- refs/heads/srj-initshape
- refs/heads/srj-inv
- refs/heads/srj-ir
- refs/heads/srj-irmut2
- refs/heads/srj-iwyu
- refs/heads/srj-iwyu3
- refs/heads/srj-javascript_work_in_progress
- refs/heads/srj-lensblur
- refs/heads/srj-lessinc
- refs/heads/srj-llvm-loop-opt
- refs/heads/srj-mak
- refs/heads/srj-maxthreads
- refs/heads/srj-mod
- refs/heads/srj-msan
- refs/heads/srj-msan-call
- refs/heads/srj-muldivmod
- refs/heads/srj-mut
- refs/heads/srj-outputs-2
- refs/heads/srj-parse
- refs/heads/srj-pch
- refs/heads/srj-printfunc
- refs/heads/srj-pygp
- refs/heads/srj-revertbits
- refs/heads/srj-schedule-storage
- refs/heads/srj-shl-shr-2
- refs/heads/srj-sio
- refs/heads/srj-static-const
- refs/heads/srj-strided-store
- refs/heads/srj-tidyh
- refs/heads/srj-tiff
- refs/heads/srj-trace
- refs/heads/srj-tutorial
- refs/heads/srj-using
- refs/heads/srj-wasmfix
- refs/heads/srj-xor2
- refs/heads/srj/abstract-gen-without-get-output-func-KEEP
- refs/heads/srj/aligned-alloc
- refs/heads/srj/aligned-alloc-2
- refs/heads/srj/aligned-malloc-with-aligned-alloc
- refs/heads/srj/all-explicit-ctor
- refs/heads/srj/anderson-thread-info-ptr
- refs/heads/srj/aot-perf
- refs/heads/srj/apps-hamal
- refs/heads/srj/argv-signatures
- refs/heads/srj/argv-types
- refs/heads/srj/async-test
- refs/heads/srj/b2cpp-const-data
- refs/heads/srj/better-xt-dispatch
- refs/heads/srj/bfloat1
- refs/heads/srj/bp
- refs/heads/srj/build_halide_h
- refs/heads/srj/c-bool
- refs/heads/srj/cache-clear
- refs/heads/srj/clang-fmt-ignore
- refs/heads/srj/clang-tidy
- refs/heads/srj/clear-c-cache
- refs/heads/srj/cmake-asan
- refs/heads/srj/cmake-asan2
- refs/heads/srj/cmake-jit-generators
- refs/heads/srj/configure-cmake
- refs/heads/srj/cpp-generator-v2-experiment-KEEP
- refs/heads/srj/crosscompile
- refs/heads/srj/csv
- refs/heads/srj/ctad
- refs/heads/srj/debug-to-file-api
- refs/heads/srj/depr
- refs/heads/srj/deprecation
- refs/heads/srj/device-copy
- refs/heads/srj/example
- refs/heads/srj/experiment
- refs/heads/srj/experiment-6967
- refs/heads/srj/exporting
- refs/heads/srj/expr_t
- refs/heads/srj/external-tensors
- refs/heads/srj/f16-convert
- refs/heads/srj/fix-pytorch
- refs/heads/srj/fixed-rollback
- refs/heads/srj/fopen-fix
- refs/heads/srj/forward
- refs/heads/srj/forward-name
- refs/heads/srj/gen-func
- refs/heads/srj/gen-func-2
- refs/heads/srj/gen-func-3
- refs/heads/srj/gen2-1
- refs/heads/srj/gen_closure
- refs/heads/srj/generator_aot_gpu_multi_context_threaded
- refs/heads/srj/globals
- refs/heads/srj/halide-buffer-crop
- refs/heads/srj/halide-malloc-alignment
- refs/heads/srj/halide-must-use
- refs/heads/srj/halide-runtime-must-use-result
- refs/heads/srj/hang-repro
- refs/heads/srj/hannk
- refs/heads/srj/hannk-aliasing
- refs/heads/srj/hannk-error-checking
- refs/heads/srj/hannk-errors
- refs/heads/srj/hannk-inplace
- refs/heads/srj/hannk-mmap
- refs/heads/srj/hannk-tflite-27
- refs/heads/srj/hannk-verbosity
- refs/heads/srj/hdrs
- refs/heads/srj/html-becomes-viz
- refs/heads/srj/implicit-mult-widening
- refs/heads/srj/issue-7076
- refs/heads/srj/iwyu
- refs/heads/srj/iwyu-2
- refs/heads/srj/iwyu-6
- refs/heads/srj/libHANNK
- refs/heads/srj/llvm_type_of
- refs/heads/srj/maybe-unused
- refs/heads/srj/meanop
- refs/heads/srj/metadata-calling-convention
- refs/heads/srj/more-tidy
- refs/heads/srj/msan-dtf
- refs/heads/srj/multimeta
- refs/heads/srj/nanobind
- refs/heads/srj/new-rt-1
- refs/heads/srj/no-threadpool
- refs/heads/srj/no-timeout-thread
- refs/heads/srj/oglc-mutexed
- refs/heads/srj/param-map
- refs/heads/srj/pip-15.x
- refs/heads/srj/pip-cron
- refs/heads/srj/possible-uninited
- refs/heads/srj/pr-7566
- refs/heads/srj/printer-size
- refs/heads/srj/profiler-data-race
- refs/heads/srj/ptr-int-cast
- refs/heads/srj/pyapps
- refs/heads/srj/pyext-fix
- refs/heads/srj/pygen-class
- refs/heads/srj/pygen-deux
- refs/heads/srj/pygen-func
- refs/heads/srj/pygen-native-types
- refs/heads/srj/pyinstall
- refs/heads/srj/pypi-try
- refs/heads/srj/pystuff
- refs/heads/srj/python-buffer-unpack
- refs/heads/srj/python-tutorial
- refs/heads/srj/reshape
- refs/heads/srj/rt-error-smallify
- refs/heads/srj/rt-return-types
- refs/heads/srj/runtime-error-handling
- refs/heads/srj/sat-fixes-exp
- refs/heads/srj/sat-fixes-exp-2
- refs/heads/srj/shadow-field
- refs/heads/srj/snprintf
- refs/heads/srj/spirv-license
- refs/heads/srj/stat-buf-deprecations
- refs/heads/srj/static-buffer-generators
- refs/heads/srj/stmt-html
- refs/heads/srj/stringify
- refs/heads/srj/synth-gen-params
- refs/heads/srj/synth-params-python
- refs/heads/srj/test-arm_sve_redux
- refs/heads/srj/test-intrinsics-bounds
- refs/heads/srj/test8076
- refs/heads/srj/test8078
- refs/heads/srj/test8094
- refs/heads/srj/test8105a
- refs/heads/srj/test8115
- refs/heads/srj/test_tmpdir_fix
- refs/heads/srj/tidy
- refs/heads/srj/tidy-format-14
- refs/heads/srj/tidymore
- refs/heads/srj/tidymore2
- refs/heads/srj/tls
- refs/heads/srj/tls-3
- refs/heads/srj/tls-4
- refs/heads/srj/tls-ucon
- refs/heads/srj/tmp-unschedule-experiment
- refs/heads/srj/tot-fix
- refs/heads/srj/try-revert-sat
- refs/heads/srj/type-traits
- refs/heads/srj/typed-func
- refs/heads/srj/ucon-all-const
- refs/heads/srj/ucon-non-const
- refs/heads/srj/visit-warnings
- refs/heads/srj/wasm-atomic2
- refs/heads/srj/wasm-simd
- refs/heads/srj/wasm-stuff
- refs/heads/srj/wasm-threads
- refs/heads/srj/wasm-updates
- refs/heads/srj/wasm-work
- refs/heads/srj/wip
- refs/heads/srj/x-rounding
- refs/heads/srj/xbuf
- refs/heads/srj/xc+plus+size+tmp
- refs/heads/srj/xc-types
- refs/heads/srj/xt-uint-cast-test
- refs/heads/srj/xtensa-arch
- refs/heads/srj/xtensa-merge
- refs/heads/srj/xvc-experimetn
- refs/heads/srj/zlib-embed
- refs/heads/standalone_autoscheduler
- refs/heads/standalone_autoscheduler_arm_worker
- refs/heads/standalone_autoscheduler_arm_worker_amazon
- refs/heads/standalone_autoscheduler_gpu
- refs/heads/standalone_autoscheduler_hexagon
- refs/heads/sticky_task_assignments
- refs/heads/store_with
- refs/heads/store_with_solver_for_super_simplify
- refs/heads/strict_float_cse_fix
- refs/heads/super_simplify
- refs/heads/super_simplify_v2
- refs/heads/super_simplify_v3
- refs/heads/transitive_wrapper
- refs/heads/trigger-release-v16
- refs/heads/tzumao-autodiff-boundarycond
- refs/heads/tzumao-gradient-autoscheduler-bug
- refs/heads/tzumao-predicate-store-load
- refs/heads/tzumao-python-buffer
- refs/heads/tzumao_autodiff_unbounded
- refs/heads/tzumao_improve_gradient_autoscheduler
- refs/heads/tzumao_issue_4297
- refs/heads/tzumao_licm_before_BI
- refs/heads/unbounded_bugs
- refs/heads/undo_async_copy_chain_black_list
- refs/heads/use_string_literals_for_blobs
- refs/heads/users/lukas/python-pip
- refs/heads/validate_sched_error_msg
- refs/heads/var_ir_fix
- refs/heads/vksnk/async-experiment
- refs/heads/vksnk/async-multiple-producers
- refs/heads/vksnk/async-order
- refs/heads/vksnk/better-loop-carry
- refs/heads/vksnk/better-message
- refs/heads/vksnk/bound-storage
- refs/heads/vksnk/bounds-widen-right
- refs/heads/vksnk/c-print-type
- refs/heads/vksnk/c-round
- refs/heads/vksnk/check-return-result
- refs/heads/vksnk/compute-with-bug
- refs/heads/vksnk/compute_with_async
- refs/heads/vksnk/dma-limit-channels
- refs/heads/vksnk/dma-min-max
- refs/heads/vksnk/expr-match-shuffle
- refs/heads/vksnk/extract-from-scalar
- refs/heads/vksnk/f16-load
- refs/heads/vksnk/fix-packvr
- refs/heads/vksnk/fix_halide_xtensa_narrow_with_rounding_shift_i16
- refs/heads/vksnk/fused-compute-with
- refs/heads/vksnk/hoist-storage-bug
- refs/heads/vksnk/lerp-intrinsics
- refs/heads/vksnk/lower-signed-shifts
- refs/heads/vksnk/missing-exception
- refs/heads/vksnk/non-widening-halves
- refs/heads/vksnk/optimize-shuffles
- refs/heads/vksnk/replace-all
- refs/heads/vksnk/restrict
- refs/heads/vksnk/roll-buffer
- refs/heads/vksnk/roundeven-arm
- refs/heads/vksnk/rvar-bounds
- refs/heads/vksnk/simplify-slice
- refs/heads/vksnk/skip-semaphores
- refs/heads/vksnk/storage-folding
- refs/heads/vksnk/strided-load-of-4_2
- refs/heads/vksnk/typed-scope
- refs/heads/vksnk/update-simd-driver
- refs/heads/vksnk/vectorize-bug
- refs/heads/vksnk/vectorize-scalarize
- refs/heads/vksnk/widening_absd
- refs/heads/vksnk/xtensa-codegen-fp16
- refs/heads/vksnk/xtensa-dma-improvements
- refs/heads/vksnk/xtensa-regroup-pass
- refs/heads/vksnk/xtensa/lift-allocs
- refs/heads/vulkan
- refs/heads/vulkan-diagnose-alloc-failures
- refs/heads/vulkan-phase0-adts
- refs/heads/vulkan-phase1-spirv
- refs/heads/vulkan-phase2-runtime
- refs/heads/vulkan2
- refs/heads/vulkan_fix_gpu_dynamic_shared_test
- refs/heads/vulkan_fix_subregion_memory_offsets
- refs/heads/webassembly-old
- refs/heads/winograd
- refs/heads/wording_fix
- refs/heads/xtensa-codegen
- refs/heads/xtensa-codegen-parallel
- refs/heads/xuanda/fix-serialize-bad-partition-always
- refs/remotes/origin/rootjalex/add_autosched_caching
- refs/tags/release_2018_02_15
- refs/tags/release_2019_08_27
- refs/tags/release_8.0.0
- refs/tags/v10.0.0
- refs/tags/v10.0.1
- refs/tags/v11.0.0
- refs/tags/v11.0.1
- refs/tags/v12.0.0
- refs/tags/v12.0.1
- refs/tags/v13.0.0
- refs/tags/v13.0.1
- refs/tags/v13.0.2
- refs/tags/v13.0.3
- refs/tags/v13.0.4
- refs/tags/v14.0.0
- refs/tags/v15.0.0
- refs/tags/v15.0.1
- refs/tags/v16.0.0
- refs/tags/v17.0.0
- refs/tags/v17.0.1
- refs/tags/v8.0.0
Permalinks
To reference or cite the objects present in the Software Heritage archive, permalinks based on SoftWare Hash IDentifiers (SWHIDs) must be used.
Select below a type of object currently browsed in order to display its associated SWHID and permalink.
Revision | Author | Date | Message | Commit Date |
---|---|---|---|---|
e50ed85 | Andrew Adams | 24 March 2023, 17:40:18 UTC | Remove HALIDE_NO_USER_CODE_INLINE | 24 March 2023, 17:40:18 UTC |
e1e638c | Andrew Adams | 24 March 2023, 17:00:50 UTC | Pacify clang-tidy | 24 March 2023, 17:00:50 UTC |
6942fbf | Andrew Adams | 24 March 2023, 16:21:56 UTC | Promote fixed-point intrinsics out of the Internal namespace and add deprecated wrappers for them in the Internal namespace so that we don't break any existing code | 24 March 2023, 16:21:56 UTC |
4fa913e | Volodymyr Kysenko | 23 March 2023, 16:46:58 UTC | Add missing #include <exception> (#7445) | 23 March 2023, 16:46:58 UTC |
4cb6dba | Andrew Adams | 19 March 2023, 20:55:40 UTC | Redo CPU schedule for bilateral grid (#7436) | 19 March 2023, 20:55:40 UTC |
badf486 | Steven Johnson | 17 March 2023, 17:41:48 UTC | Disable performance_boundary_conditions under WebGPU pending #7420 (#7435) | 17 March 2023, 17:41:48 UTC |
643b2f1 | Steven Johnson | 16 March 2023, 21:42:34 UTC | Modify runtime calls to always return a valid halide_error_code_t value (#7404) * Modify runtime calls to always return a valid halide_error_code_t value Currently, the return values from our runtime code are a mishmash -- there's lots of code that returns any random nonzero value to indicate an error. This isn't wrong per se, but it's not clean, and it's desirable that the return values are predictable. This PR doesn't change the call signature of any (public) Halide Runtime functions, but modifies the internal logic so that all return values are valid values of `enum halide_error_code_t`. Generally, there should be minimal change to the code otherwise, although I did leave in a few drive-by changes that I couldn't resist (e.g., better error-checking when dynamically loading symbols). My long-term goal here is to eventually propose changing the signature of runtime functions that return errors to actually return `enum halide_error_code_t`; as you might imagine, making that transition might be controversial for a number of reasons. This PR is intended to be a way to make such a future transition easier to reason about, while arguably improving the code quality of the runtime slightly. * tidy * Update opencl.cpp * trigger buildbots * trigger buildbots * Fix merge mistake * Update gpu_context_common.h * Update cuda.cpp * Werror * status_ * if-with-initializer format * Update cuda.cpp * Update opencl.cpp * Update cuda.cpp * Update device_interface.h * Update hexagon_cache_allocator.cpp * Update printer.h * remove prefixes * "device field is already non-zero" * Update opencl.cpp * Update cuda.cpp * Fix error spacing * trigger buildbots | 16 March 2023, 21:42:34 UTC |
88b3ef8 | Steven Johnson | 16 March 2023, 21:35:52 UTC | Split WebGPU runtime into two variants (#7248 workaround) (#7419) * Split WebGPU runtime into two variants (#7248 workaround) Halide promises that you can crosscompile to *any* supported target from a 'stock' build of libHalide. Unfortunately, the initial landing of WebGPU support breaks that promise: we compile the webgpu runtime support (webgpu.cpp) with code that is predicated on `WITH_DAWN_NATIVE` (for Dawn vs Emscripten, respectively). This means that if you build Halide with `WITH_DAWN_NATIVE` defined, you can *only* target Dawn with that build of Halide; similarly, if you build with `WITH_DAWN_NATIVE` not-defined, you can only target Emscripten. (Trying to use the 'wrong' version will produce link-time errors.) For people who build everything from source, this isn't a big deal, but for people who just pull binary builds, this is a big problem. This PR proposes a temporary workaround until the API discrepancies are resolved: - Compile the existing webgpu.cpp runtime *both* ways - in LLVM_Runtime_Linker.cpp, select the correct variant based on whether the Target is targeting wasm or not - Profit! This is a rather ugly hack, but it should hopefully be (relatively) temporary. * A few more fixes * Update HalideGeneratorHelpers.cmake * Update interpreter.cpp * Update interpreter.cpp | 16 March 2023, 21:35:52 UTC |
cc74ee8 | Steven Johnson | 16 March 2023, 18:22:23 UTC | Move long boilerplate C/C++ code into template files (#7426) * Move long boilerplate C/C++ code into template files Codegen_C has a couple of long strings with boilerplate code that is conditionally emitted; at least one of these is too long for a single string literal under MSVC. Let's try moving these into standalone files instead; this may make it easier to use conventional tooling on the C++ code, and make Codegen_C easier to read and think about. (Note that the not-yet-landed Xtensa branch should also use this approach, if we decide this approach is good, since it has even more such code.) TODO: probably would be good to augment `binary2cpp` to allow option comments in the source file that are stripped in the output file (e.g. "This file is used in CodeGen_C.cpp for blah blah blah, look out for blah") * Remove detritus | 16 March 2023, 18:22:23 UTC |
50f8c85 | Steven Johnson | 15 March 2023, 17:16:01 UTC | Disable correctness_atomics on Windows with Cuda, alas (#7423) (#7424) | 15 March 2023, 17:16:01 UTC |
ae59b91 | Steven Johnson | 14 March 2023, 20:56:33 UTC | Ignore assertions inside WebGPU kernels (#7418) This is the approach that (e.g.) the OpenCL backend takes to assertions inside kernel code. | 14 March 2023, 20:56:33 UTC |
e966163 | Steven Johnson | 14 March 2023, 20:33:34 UTC | Log name of failing function for !function_takes_user_context (#7417) | 14 March 2023, 20:33:34 UTC |
05fa61a | Steven Johnson | 14 March 2023, 18:40:40 UTC | Fix for top-of-tree LLVM (#7416) | 14 March 2023, 18:40:40 UTC |
9199849 | Steven Johnson | 14 March 2023, 18:01:46 UTC | A few minor cleanups in WebGPU backend (#7413) * A few minor cleanups in WebGPU backend Mostly just using using halide_error_code_t values everywhere. * trigger buildbots * Update webgpu.cpp | 14 March 2023, 18:01:46 UTC |
b63139e | James Price | 14 March 2023, 17:01:25 UTC | Update mini_webgpu.h with latest changes from Dawn (#7415) | 14 March 2023, 17:01:25 UTC |
d383cb9 | James Price | 14 March 2023, 16:58:12 UTC | Fix null device crash during WebGPU initialization (#7414) If RequestDevice fails, make sure we exit initialization early instead of trying to create a staging buffer with a nullptr device. | 14 March 2023, 16:58:12 UTC |
6d19cec | Steven Johnson | 13 March 2023, 22:36:38 UTC | printer.h should include HalideRuntime.h (#7412) printer.h uses uint64_t, so it needs to include something that ensures that type is defined. HalideRuntime.h is probably the right choice (since it always transitively includes runtime_internal.h when compiling runtime. | 13 March 2023, 22:36:38 UTC |
078465c | James Price | 13 March 2023, 20:30:38 UTC | Add initial support for WebGPU (#6492) * [WebGPU] Add runtime stubs and codegen skeleton All runtime functions are currently unimplemented, and the shader codegen just emits an empty compute shader. * [WebGPU] Implement lazy device initialization Sychronize device access using a WebGpuContext object, as is done in the other GPU runtimes. Device initialization is asynchronous, so we rely on Emscripten's Asyncify mechanism to wait for it to complete. * [WebGPU] Implement device release Release the device and adapter. * [WebGPU] Add scoped error handling mechanism WebGPU uses asynchronous callbacks for error handling, so we need to spin-lock while waiting for them to fire in order to determine success/failure status for any WebGPU APIs that we call. * [WebGPU] Implement device malloc/free Create a WGPUBuffer with Storage|CopyDst|CopySrc usages. We'll need a staging buffer to perform host<->device transfers, as a buffer used as a storage buffer cannot be mapped. * [WebGPU] Implement basic host<->device copies Use a staging buffer to copy data from the device, since we cannot map a buffer that is used a storage buffer. This logic will need an overhaul in order to support buffers that represent sub-regions of larger buffers. * [WebGPU] Implement halide_webgpu_device_sync Just wait for all submitted commands to complete. * [WebGPU] Implement shader compilation * [WebGPU] Implement core of WGSL codegen This implements enough of the WGSL codegen required to generate code for a 32-bit version of the apps/blur example. Buffer arguments are emitted as read_write storage buffers. Non-buffer arguments are collected into a structure and generated as a single uniform buffer. Workgroup sizes are currently required to be constant. This can relaxed in the future by using pipeline-overridable constants. * [WebGPU] Implement Cast node WGSL uses different syntax for type casts. * [WebGPU] Implement the float_from_bits() intrinsic * [WebGPU] Implement run function Create a compute pipeline, create a uniform buffer for non-buffer arguments, and dispatch a compute command to the queue. Does not handle workgroup storage yet. * [WebGPU] Move queue into WgpuContext class The queue handle received from wgpuDeviceGetQueue needs to be released, so it makes more sense to automatically get/release this handle in the WgpuContext constructor/destructor. * [WebGPU] Add support for JIT This requires a native implementation of WebGPU (e.g. Dawn or wgpu). * [WebGPU] Enable the gpu_only AOT generator test Link generator tests against the native WebGPU library if specified. * [WebGPU] Add support for targeting dawn-native We have to make this a compile-time switch for now, as the supported APIs currently differ between Dawn and Emscripten. We should be able to remove all of these conditionals when the native API stabilizes. * [WebGPU] Add support for AOT tests when using WASM Use Dawn's node bindings to run these tests, by pre-pending some initialization code to the script that nodejs will invoke. * [WebGPU] Print explicit types for let declarations This is makes the generated code a little more human-readable. * [WebGPU] Address first round of review comments * [WebGPU] Add copy of webgpu.h This is a verbatim copy of the Emscripten version, just without the include directives. * [WebGPU] Add comment about ASYNCIFY requirement * [WebGPU] Remove -Wno-atomic-alignment Tweak atomic primitive usage to avoid the need for this. * pacify clang-tidy * Fix more clang-tidy errors * Only use ASYNCIFY for tests when targeting WebGPU * Fix even more clang-tidy errors * [WebGPU] Add basic support to Makefile * [WebGPU] Don't wrap buffers inside structures This requirement has been removed from the WGSL specification, and the corresponding implementation changes have now landed in Dawn. * [WebGPU] Fix debug message tag * [WebGPU] Update WGPUErrorFilter enum in header * [WebGPU] Update WGSL attribute syntax The WGSL specification recently changed attribute syntax from [[attribute]] to @attribute. * [WebGPU] Add README_webgpu.md Explains how to configure Halide to target WebGPU for both Emscripten and Dawn native. Also lists several known limitations. * [WebGPU] Move native WebGPU library CMake logic This is the correct place, otherwise the link line order is wrong. * [WebGPU] Implement WGSL codegen for serial loops * [WebGPU] Implement WGSL codegen for Allocate Use array types for stack allocations. Leave GPU shared memory unimplemented for now. * [WebGPU] Implement WGSL codegen for Select Use the WGSL select builtin function, which supports bool vector conditions for component-wise selection too. * [WebGPU] Mark 64-bit types as unsupported * [WebGPU] Implement device_and_host_{malloc,free} Just use the default implementations. * [WebGPU] Fixed WGSL codegen for boolean vectors * [WebGPU] Implement f32 math intrinsics * [WebGPU] Implement inverse and inverse sqrt * [WebGPU] Fixup some errors in WGSL codegen * [WebGPU] Implement logical and/or for bool vectors * [WebGPU] Implement WGSL codegen for Broadcast node * [WebGPU] Implement WGSL codegen for Ramp node * [WebGPU] Emulate 8- and 16-bit integers Use atomics to emulate storage, and widen the values to 32-bits when operating on them. * [WebGPU] Avoid buffer name collisions Buffers are declared as global variables in WGSL, so prefix them with the kernel name to avoid collisions. * [WebGPU] Fix divide-by-power-of-two WGSL requires the RHS of a shift to be unsigned. * [WebGPU] Implement codegen for gpu_thread_barrier * [WebGPU] Implement WGSL codegen for Evaluate This fixes an issue with the halide_unused() intrinsic. * [WebGPU] Add support for shared memory This currently only supports statically-sized shared memory allocations; dynamically-sized allocations require using pipeline-overridable constants which are not yet implemented in Dawn. * [WebGPU] Fix 8/16-bit load/store emulation Loads and stores that implicitly casted to/from 32-bits were casting when they should not have been, and not casting when they should. * Use generic 64-bit support query in gpu_mixed_shared_mem_types This is more robust than checking for specific APIs. * [WebGPU] Fix object cleanup during teardown * [WebGPU] Do not re-allocate device buffers * [WebGPU] Mark maximum vector width as 4 in tests * [WebGPU] Add functions to object lifetime tracker * [WebGPU] Override FloatImm handling We need to explicitly suffix the integer literal with a `u` to make it unsigned. * [WebGPU] Scalarize predicated loads * [WebGPU] Implement if_then_else intrinsic This is generated when scalarizing predicated loads. * [WebGPU] Enable gpu_free_sync test * [WebGPU] Implement print_reinterpret Use WGSL's bitcast<> operator. * [WebGPU] Implement print_extern_call This is just a regular function call in WGSL. * Add missing include and namespace for isnan/isinf * [WebGPU] Avoid short-circuiting operators These can cause issues with WGSL's uniformity analysis. * [WebGPU] Use commas for struct member seperators * [WebGPU] Update API headers and usages Two API functions were renamed. * [WebGPU] Use CodeGen_GPU_C base class for codegen Introduce a new enum value for WGSL's vector syntax. * [WebGPU] Add warning for emulated narrow integers * [WebGPU] Update README with latest status * [WebGPU] Add support for non-contiguous copies Also adds support for buffer cropping/slicing. * [WebGPU] Fix clang-tidy error * [WebGPU] Use atomicCmpXchg for 8/16-bit emulation Halide assumes that write-write data-races are benign when both threads are writing the same value, but this is not true when those writes are implemented using atomicXor. We need to use atomicCompareExchangeWeak in a loop to perform this emulation instead. Unfortunately this makes things even slower :-( * [WebGPU] Support non-32-bit parameter types Expand them to 32-bits in the shader and the runtime. * [WebGPU] Fix mixed types in buffers The atomic emulation of narrow types shouldn't use the allocation type, since the same buffer may be re-used for multiple types. This means we also sometimes need to perform 32-bit accesses using atomics as well. Instead of using the buffer allocation type, we now pre-traverse the IR to check for accesses that will require emulation, and mark the corresponding buffer as such. * [WebGPU] Show validation errors for failed maps * [WebGPU] Round up buffer offsets and sizes The WebGPU API requires that these are multiples of 4. * [WebGPU] Update implementation status in README * [WebGPU] Replace @stage(compute) with @compute * [WebGPU] Polyfill the pow_f32 intrinsic The pow builtin in WGSL only has the correct semantics for x>0, so we need to emulate the behavior for the other cases. * [WebGPU] Skip the gpu_allocation_cache test Allocation cache is not yet implemented for the WebGPU backend, and so this test takes forever. * [WebGPU] Use builtins for inverse hyperbolics * [WebGPU] Map rint() to round() WGSL's round() builtin has round-to-nearest-even semantics. * [WebGPU] Set device lost callback This provides more information about events that cause the device to become inaccessible. * [WebGPU] Use i32 for bool parameters The bool type cannot be used in the uniform storage class in WGSL. * [WebGPU] Raise limits for buffer size and workgroup storage * [WebGPU] Update mini_webgpu.h Disable clang-format for this file. * [WebGPU] Add support for dynamic workgroups This is implemented using pipeline-overridable constants. * [WebGPU] Avoid using 'new' as an identifier This is a reserved keyword in WGSL. * [WebGPU] Do not merge workgroup allocations Since we promote 8- and 16-bit integers to 32-bit in workgroup memory, merging multiple different GPUShared allocations into a single one can cause different regions to overlap and data races ensue. * [WebGPU] Fix Emscripten support * [WebGPU] Use const for integer immediates This avoids some MSL compiler ICEs with the code generated by Dawn, and also makes it easier for the WGSL compiler to constant-fold these values. * [WebGPU] Squelch clang-tidy error * [WebGPU] Note Dawn's dependency on go in README * [WebGPU] Add links to Emscripten vs Dawn issue * [WebGPU] Show error finding WebGPU library fails * [WebGPU] Add link to issue about Windows support * [WebGPU] Rename roundUpToMultipleOf4 ...and use it in one place where we weren't. * [WebGPU] Add links to wrap_native issue * [WebGPU] Use debug_assert for some runtime errors * [WebGPU] Stop using designated initializers Add -Wc++20-designator flag to runtime build to prevent this from happening again. * [WebGPU] Update mini_webgpu.h * [WebGPU] Fix a validation issue with ToT Dawn A bitcast that produces a NaN is invalid in WGSL, so implement the `nan()` intrinsic via `float_from_bits()`. * Update README_webgpu.md * Add is_finite_f32, is_inf_f32, is_nan_f32 * Update isinf, isnan, isfinite; add inf_f32, neg_inf_f32 * correctness_isnan should be skipped for WebGPU * Update isnan.cpp * Update atomics_gpu_8_bit.cpp * Fix python_tutorial_lesson_10_aot_compilation_run * Partial fix for generator_aot_acquire_release() This adds the necessary (horrible hackery) to bring the WebGPU case in line with the other backends... but the test still fails, apparently due to the same copy-to-host bug we suspect for generator_aot_gpu_only.Pushing this anyway because it was annoying to write :-) * [WebGPU] Fix AOT test build for non-WASM * [WebGPU] Move README to root * [WebGPU] Address review comments * wip * Update CMakeLists.txt * Work-in-progress for generator_aot_gpu_multi_context_threaded * [WebGPU] Use a per-context staging buffer This fixes the generator_aot_gpu_multi_context_threaded tests. * [WebGPU] Fix clang-format issue * [WebGPU] Move staging buffer creation We shouldn't be doing this inside a callback as we use async error checking routines. * Revert "Merge branch 'webgpu' of https://github.com/jrprice/Halide into pr/6492" This reverts commit 9d79ac60153e12ed51ea9f1a91a544f3fc7ccc78, reversing changes made to 0673e6754e53e62cf6029efdb94a57447db2f03a. * Skip correctness_multi_way_select on x86 Macs (for now) * clang-format * [WebGPU] Recommit 742db3feb888394ff6529c357da3f32ae9286ea4 | 13 March 2023, 20:30:38 UTC |
78097a7 | Steven Johnson | 13 March 2023, 17:01:35 UTC | Dont use negative values for exit() (#7405) * Dont use negative values for exit() A program that terminates via `exit(-1)` leaves `$? = 255` (per the exit(3) manpage), as the value is chopped with `& 0377`. This makes it hard to use `git bisect` to track down bugs, as considers any exit code > 127 (or = 15, oddly enough) equivalent to an `abort()` and terminates the bisect. IMHO this is perverse behavior on the part of `git bisect`, but it is what it is, so let's revamp our tests to avoid calling `exit(-1);`. * trigger buildbots | 13 March 2023, 17:01:35 UTC |
822f5da | Steven Johnson | 13 March 2023, 17:00:50 UTC | Don't return negative values from main() (#7406) * Don't return negative values from main() A program that returns `-1` from `main()` leaves $? = 255 (per the exit(3) manpage), as the value is chopped with & 0377. This makes it hard to use git bisect to track down bugs, as considers any exit code > 127 (or = 15, oddly enough) equivalent to an abort() and terminates the bisect. IMHO this is perverse behavior on the part of git bisect, but it is what it is, so let's revamp our tests to avoid returning `-1` from main to indicate failure, and return `1` instead. * trigger buildbots * validate_gpu_object_lifetime() handling | 13 March 2023, 17:00:50 UTC |
f76bcc8 | Steven Johnson | 12 March 2023, 17:49:46 UTC | Cleanup of src/runtime/internal and test/runtime (#7399) * Cleanup of src/runtime/internal and test/runtime - Don't include .cpp files. - Don't use header-only "libraries" that rely on include order or being included only once (to wit: test/runtime/common.h -> common.cpp) - All files should explicitly #include what they need, even if they think it's already included (ie, order of include files should not matter) - In src/runtime/internal, change all `halide_abort_if_false` -> `halide_debug_assert` - in test/runtime, add HALIDE_CHECK to common.h and use it for tests instead of `halide_abort_if_false` * 32-bit fixes | 12 March 2023, 17:49:46 UTC |
0beb081 | Ttayu | 11 March 2023, 20:45:17 UTC | docs: Use halide.imageio instead of imageio (#7409) * docs: Use halide.imageio instead of imageio * trigger buildbots --------- Co-authored-by: Steven Johnson <srj@google.com> | 11 March 2023, 20:45:17 UTC |
4bf660e | Andrew Adams | 11 March 2023, 00:53:44 UTC | Fix overflow in x86 absd lowering (#7407) * Fix overflow in x86 absd lowering * Fix default lowering too | 11 March 2023, 00:53:44 UTC |
ead83e6 | Ttayu | 10 March 2023, 18:23:21 UTC | Docs/update readme python (#7402) * docs: Update simple usage in README_python.md * docs: Add imports | 10 March 2023, 18:23:21 UTC |
b02de4b | Steven Johnson | 09 March 2023, 17:46:43 UTC | Update all use of actions/checkout to v3 (#7400) | 09 March 2023, 17:46:43 UTC |
6a6c842 | Steven Johnson | 07 March 2023, 23:56:03 UTC | Backport changes to CodegenC from the xtensa branch (#7396) Backported with some trivial style changes (virtual method instead of flag, since this should never changes for a given instance) | 07 March 2023, 23:56:03 UTC |
a2df5bd | Steven Johnson | 07 March 2023, 19:29:08 UTC | Attempt to give proper labels to pip dev builds (v2) (#7395) A much-simpler alternative to #7394 | 07 March 2023, 19:29:08 UTC |
4a80251 | Andrew Adams | 03 March 2023, 20:57:47 UTC | Destringify CanonicalizeGPUVars (#7386) * Destringify CanonicalizeGPUVars This new implementation takes the high-water marks of each type of GPU loop, instead of filtering using the prefix of the loop name. * Better comments * fix typo | 03 March 2023, 20:57:47 UTC |
aa8fcad | Terry Heo | 03 March 2023, 20:56:52 UTC | hannk: Provide weak symbol functions to use op profiling (#7388) * hannk: Provide weak symbol functions to use op profiling You can add your own profiler with strong symbols. * hannk: Guard profiler feature with HANNK_PROFILER | 03 March 2023, 20:56:52 UTC |
387a19c | Xuanda Yang | 03 March 2023, 01:10:50 UTC | Make README consistent on supported LLVM versions (#7390) Make README consistent with https://github.com/halide/Halide/pull/7093 | 03 March 2023, 01:10:50 UTC |
122b5b6 | Terry Heo | 01 March 2023, 22:22:20 UTC | hannk: Add device_sync method hannk::Tensor class (#7387) hannk: Add devic_sync method hannk::Tensor class This method is useful to debug performance with synchronous execution. | 01 March 2023, 22:22:20 UTC |
303a90c | Steven Johnson | 01 March 2023, 17:21:48 UTC | Minor cleanup of GPUCompilationCache (#7376) * Minor cleanup of GPUCompilationCache While tracking down an apparently-unrelated threading bug in the webgpu backend, I made some tweaks to this code that I think are worth keeping. The main one of importance is that `release_hold()` and `release_context()` really should acquire the mutex -- they weren't before -- so now all public methods are properly mutexed. The other changes are mostly cosmetic: - Moved helper methods to be private rather than public - Changed the id value size to be `uintptr_t` rather than `uint32_t`; the space allocated for them is sizeof(void*). (Not sure this moves the needle but it felt right.) - Removed unused ctor for CachedCompilation Also, drive-by change in printer.h to capture some logging improvements. * Fix deadlock release_all() and release_context() were contending for the mutex | 01 March 2023, 17:21:48 UTC |
fba892a | Volodymyr Kysenko | 01 March 2023, 00:29:42 UTC | Specify a full type for llvm::IRBuilder (#7384) Specify full type for llvm::IRBuilder | 01 March 2023, 00:29:42 UTC |
5c02ae2 | Adrian Lebioda | 28 February 2023, 02:14:41 UTC | Bounds visitor for div was missing single_point mutated case (#7379) * Bounds visitor for div was missing single_point mutated case Signed-off-by: Adrian Lebioda <adrian.lebioda@hexagon.com> * Add test Signed-off-by: Adrian Lebioda <adrian.lebioda@hexagon.com> --------- Signed-off-by: Adrian Lebioda <adrian.lebioda@hexagon.com> Co-authored-by: Adrian Lebioda <adrian.lebioda@hexagon.com> | 28 February 2023, 02:14:41 UTC |
bdba694 | Steven Johnson | 28 February 2023, 02:14:24 UTC | Add Callable default ctor + `defined()` method (#7380) * Add Callable default ctor + `defined()` method This allows it to behave like * Add user_assert + test | 28 February 2023, 02:14:24 UTC |
c42a5b2 | Andrew Adams | 27 February 2023, 16:46:41 UTC | Remove a gross hack from gpu_only_aottest (#7378) * Remove a gross hack from gpu_only_aottest Also add metal support * Add missing include | 27 February 2023, 16:46:41 UTC |
09400f6 | Andrew Adams | 25 February 2023, 17:16:32 UTC | Bounds visitors for min/max were missing single_point mutated case (#7377) * Bounds visitors for min/max were missing single_point mutated case Partially fixes #7374 * Add test | 25 February 2023, 17:16:32 UTC |
b6a18b8 | Steven Johnson | 23 February 2023, 19:36:32 UTC | Update WABT to 1.0.32; Increase stack size for WASM AOT apps (#7373) | 23 February 2023, 19:36:32 UTC |
144c1a4 | Steven Johnson | 23 February 2023, 17:59:39 UTC | correctness_round should use Target::supports_type() (#7372) This gives it proper support for new GPU backends | 23 February 2023, 17:59:39 UTC |
b17806d | Steven Johnson | 23 February 2023, 17:59:13 UTC | Use HalideFreeHelper for the register_destructor (#7371) Slightly cleaner code. Also, drive-by change of NULL -> nullptr | 23 February 2023, 17:59:13 UTC |
629da52 | Steven Johnson | 23 February 2023, 17:58:37 UTC | Use single-char form of `unique_name` for semaphores (#7370) The multi-char form of `unique_name` will append a `$` to the identifier, e.g. `sema$4`. This isn't really legal for a C/C++ identifier. | 23 February 2023, 17:58:37 UTC |
386a2d1 | Steven Johnson | 22 February 2023, 23:39:25 UTC | Clean Up HalideFreeHelper code (main) (#7369) * Clean Up HalideFreeHelper code (main) - Revise HalideFreeHelper to be a templated struct, to save the unnecessary stack storage for the function - Add emit_halide_free_helper() method to consolidate usage * Update CodeGen_C.cpp | 22 February 2023, 23:39:25 UTC |
3246844 | Andrew Adams | 21 February 2023, 21:18:30 UTC | Use a std::unique_ptr for the IR Builder (#7356) * Use a std::unique_ptr for the IR Builder instead of a raw owning pointer * Use make_unique | 21 February 2023, 21:18:30 UTC |
e19a036 | Andrew Adams | 21 February 2023, 17:28:43 UTC | Overflow on casts is fine for ints < 32 bits (#7366) | 21 February 2023, 17:28:43 UTC |
ccc085a | Steven Johnson | 16 February 2023, 19:54:27 UTC | Update CMakePresets.json to use VS2022 instead of VS2019 (main) (#7359) | 16 February 2023, 19:54:27 UTC |
b65ea62 | Steven Johnson | 15 February 2023, 18:25:04 UTC | Remove unused code in VectorizeLoops (#7354) | 15 February 2023, 18:25:04 UTC |
18eb7d8 | Andrew Adams | 15 February 2023, 17:18:56 UTC | Permit vectorization of non-recursive atomic operations (#7346) * Vectorization of non-recursive atomic operations * Remove dead Vars | 15 February 2023, 17:18:56 UTC |
e5ed226 | Steven Johnson | 15 February 2023, 17:03:41 UTC | Fix Python error handling (#7352) * Fix Python error handling Error handling in the Python bindings wasn't quite right for JIT: We previously replaced halide_error() to throw a C++ exception. Sounds good, but unfortunately, doesn't work reliably: if called from jitted code (which doesn't know about C++ exceptions), the throw statement may be unable to find the enclosing try block (which is outside jitted code), meaning it will call std::terminate. Now, instead, we just leave the JIT error handler unset, and call with an explicit JITUserContext with a custom print handler; in theory, this meant that the code in JITFuncCallContext::finalize() would check for an error after the call into jitted code, and call `halide_runtime_error` if so (which would then trigger an all-in-C++-exception). Unfortunately... (2) JITFuncCallContext is broken by design; it mutates the input JITUserContext, so that trying to use the same JITUserContext for two calls in a row leaves you with a JITUserContext with (at least) the error_handler set. Since at least one of the realize() calls does this twice (once for bounds query, once for execution), this means that an error in the second call would never be seen, since finalize() only reported errors if there wasn't a custom error handler on input. Per @abadams suggestion, we work around this by treating 'JITErrorBuffer::handler' as 'no custom error handler', which is mostly true. (But really, JITFuncCallContext and JITUserContext are a hard-to-reason-about mess and arguably need to rethought entirely.) (3) Removed entirely-unnecessary overrides of runtime print and error handlers from PyStubImpl; despite the comments, this code is unnecessary. * format | 15 February 2023, 17:03:41 UTC |
7963cd4 | Steven Johnson | 14 February 2023, 17:14:05 UTC | Change early-bound default args in Python bindings to late-bound (#7347) In PyBind11, if you specify a default argument for a method, it is evaluated when the Python module is initialized, *not* when the method is called (as you might expect in C++). For defaults that are just constants/literals, this is no big deal, but when calling get_*_target_from_environment, this means it is called at module init time -- also normally not a big deal (since the values ~never change at runtime anyway), with one big exception (no pun intended): if the function throws an exception (e.g. via calling user_assert() or similar), that exception is thrown at Module-initialization time, which is a much more inscrutable crash, and one that is very hard to recover from. This may seem unlikely, but can happen pretty easily if you set (say) HL_JIT_TARGET=host-cuda (or other gpu) and the given GPU runtime isn't present on the given system; the current behavior is basically "make if impossible for the libHalidePython bindings to run", whereas what we want is "runtime exception thrown when you call the method". This changes the relevant methods to use `Target()` as the default value, and inside the method wrapper, if the value passed equals `Target()`, it replaces the value with the righ `get_*_target_from_environment()` call. (This turned up while doing some testing of https://github.com/halide/Halide/pull/6924 on a system without Vulkan available) | 14 February 2023, 17:14:05 UTC |
8bd07fb | Andrew Adams | 14 February 2023, 01:52:53 UTC | Fix tuple output bounds checks (#7345) Fix #7343 Tuple outputs weren't getting appropriate bounds checks due to overzealous culling of uninteresting code in the add_image_checks pass. | 14 February 2023, 01:52:53 UTC |
22aed20 | Steven Johnson | 11 February 2023, 02:42:05 UTC | Devirtualize the protected compile() methods in Codegen_C (#7341) With the addition of `preprocess_function_body()`, neither of these need to be virtual, and devirtualizing them avoid `hidden overloaded virtual function` warnings in subclasses that don't override them | 11 February 2023, 02:42:05 UTC |
6c5ca8e | Steven Johnson | 11 February 2023, 00:34:41 UTC | Tiny improvements in codegen in C backend (#7337) * Tiny improvements in codegen in C backend (1) Emit `true` or `false` instead of `(bool)(0ull)` etc for bool literals (2) Avoid redundant temporaries in print_cast_expr(), which occur in a small but nonzero number of cases Basically this means that code currently like ``` bool _523 = (bool)(0ull); bool _524 = (bool)(_523); ... foo(_524); ``` becomes ``` foo(false); ``` ...I'm sure this has no output on final object code, but it makes the generated C code less weird to read. * Also avoid extra intermediates for typed nullptr * Also use std::isnan() and std::isinf() * Update CodeGen_C.cpp | 11 February 2023, 00:34:41 UTC |
a6c5be7 | Steven Johnson | 10 February 2023, 21:28:51 UTC | Add a hook to Codegen_C::compile() (#7335) At least one subclass of Codegen_C currently has to replicate ~all of the compile(LoweredFunc) method, with the result that it has often gone stale (and still is stale) wrt changes in the base; this adds an optional method to allow some modifications to the function body just before it is printed, to avoid redundant code. | 10 February 2023, 21:28:51 UTC |
88d40c2 | Steve Suzuki | 10 February 2023, 18:07:56 UTC | Fix issue in find_package in cross-compilation for no OS (#7282) When using toolchain where Threads libs are not available, which is the case in baremetal target cross-compilation, we were not able to load even HalideHelpers pacakge. Co-authored-by: Alex Reinking <alex.reinking@gmail.com> | 10 February 2023, 18:07:56 UTC |
35322c3 | Steven Johnson | 10 February 2023, 00:22:10 UTC | Fix a subtle uninitialized-memory-read in Buffer::for_each_value() (#7330) * Fix a subtle uninitialized-memory-read in Buffer::for_each_value() When we flattened dimensions in for_each_value_prep(), we would copy from one past the end, meaning the last element contained uninitialized garbage. (This wasn't noticed as an out-of-bounds read because we overallocated in structure in for_each_value_impl()). This garbage stride was later used to advance ptrs in for_each_value_helper()... but only on the final iteration, so even if the ptr was wrong, it didn't matter, as the ptr was never used again. Under certain MSAN configurations, though, the read would be (correctly) flagged as uninitialized. This fixes the MSAN bug, and also (slightly) improves the efficiency by returning the post-flattened number of dimensions, potentially reducing the number of iterations f for_each_value_helper() needed. * Oopsie * Update HalideBuffer.h * Update HalideBuffer.h | 10 February 2023, 00:22:10 UTC |
ae3f401 | Steven Johnson | 09 February 2023, 23:08:37 UTC | Explicitly remove -D_GLIBCXX_ASSERTIONS from LLVM definitions (#7332) Explicitly remove -D_GLIBCXX_ASSERTIONS from LLVM definitions as a workaround for https://reviews.llvm.org/D142279 | 09 February 2023, 23:08:37 UTC |
4156c5a | Steven Johnson | 09 February 2023, 21:51:22 UTC | Allow _Float16 as alias for float16_t in halide_type_of<>() (#7325) (#7326) | 09 February 2023, 21:51:22 UTC |
734e34a | Steven Johnson | 09 February 2023, 17:59:12 UTC | Remove deprecated `HVX_shared_object` feature (#7331) This has been marked 'deprecated' for quite a while, and has no affect on codegen or, well, anything else. Let's remove it. | 09 February 2023, 17:59:12 UTC |
0f6003e | Terry Heo | 08 February 2023, 20:26:31 UTC | Float16: Remove unused header dependency (#7324) IRMutator.h is not needed for the Float16.h. | 08 February 2023, 20:26:31 UTC |
c3f3318 | Steven Johnson | 08 February 2023, 20:25:37 UTC | Fixes for top-of-tree LLVM (#7329) * Fixes for top-of-tree LLVM * fix * times ten * Update LLVM_Output.cpp | 08 February 2023, 20:25:37 UTC |
ddb515a | Steve Suzuki | 07 February 2023, 18:41:04 UTC | Improve support for Arm baremetal compilation and runtime (#7286) * Improve support for Arm baremetal compilation and runtime - Add Target feature "semihosting" mode for baremetal runtime - Fix error of aligned_alloc() when compiled by Arm GNU toolchain * Modify comments for Target feature semihosting * Add an example app to guide cross-compilation for baremetal target * Update build steps in HelloBaremetal * Fix line-ending * Set CMake variable BAREMETAL in toolchain file | 07 February 2023, 18:41:04 UTC |
34d256f | Steve Suzuki | 07 February 2023, 18:40:22 UTC | Make auto scheduler libs available in HalideHelpers package (#7285) * Make auto scheduler libs available in HalideHelpers package find_package(HalideHelpers) allows us to use add_halide_library(). But auto scheduler libs are not available unless they are in Halide-Interfaces.cmake. Note: Those libraries are not actually linked to the target application, but need to be available for add_custom_command call. | 07 February 2023, 18:40:22 UTC |
0c7722f | Terry Heo | 07 February 2023, 17:37:09 UTC | Add buffer sync methods hannk::Tensor class (#7323) Add few methods for GPU memory interaction. | 07 February 2023, 17:37:09 UTC |
0b7379f | Steve Suzuki | 07 February 2023, 17:08:32 UTC | Warn emulated float16 equivalent is generated (#7307) * Warn emulated float16 equivalent is generated | 07 February 2023, 17:08:32 UTC |
a55a09a | Dmitry Kurtaev | 07 February 2023, 14:17:58 UTC | Fix Halide cross-compilation (#7073) Use CMAKE_CROSSCOMPILING_EMULATOR for llvm-as and clang imported targets | 07 February 2023, 14:17:58 UTC |
1ad328a | Alex Reinking | 07 February 2023, 01:18:21 UTC | Fix LLVM 17+ build integration on 32-bit systems (#7322) * Fix LLVM 17+ build integration on 32-bit systems Fixes #7319 * add detail and precision to comment | 07 February 2023, 01:18:21 UTC |
91f3ac0 | Steve Suzuki | 06 February 2023, 22:23:47 UTC | Fix segfault by nonconstant bound in Adams2019 (#7321) Fix segmentation fault in Adams2019 in case the estimate or bound of Func is set to nonconstant Expr. | 06 February 2023, 22:23:47 UTC |
01f9e2d | Andrew Adams | 06 February 2023, 19:05:01 UTC | Replace some push_backs with emplace_back (#7317) | 06 February 2023, 19:05:01 UTC |
e9aecee | Steven Johnson | 03 February 2023, 21:17:47 UTC | Make visit_leaf() public in hannk/ops.h (#7318) * Make visit_leaf() public in hannk/ops.h This makes it easier for downstream code to experiment with adding ops * Update ops.h | 03 February 2023, 21:17:47 UTC |
0782d80 | Tom Westerhout | 01 February 2023, 18:24:01 UTC | Make Callable::call_argv_fast public (#7315) * Make Callable::call_argv_fast public * Add rough specification of the calling convention * Fix a typo | 01 February 2023, 18:24:01 UTC |
beba53a | Steven Johnson | 31 January 2023, 21:08:15 UTC | halide_popcount<uint64_t> is broken (#7313) Would not compile for Win32 or any other compiler without __builtin_popcountll available. (How did this get checked in without being tested on MSVC?) | 31 January 2023, 21:08:15 UTC |
fe76ab2 | Steven Johnson | 30 January 2023, 22:01:42 UTC | Minimal updates to allow Halide building with LLVM17 (#7309) * Minimal updates to allow Halide building with LLVM17 (Opening as draft initially until Buildbots build the new LLVM versions) * trigger buildbots | 30 January 2023, 22:01:42 UTC |
dd973f4 | Mikhail Usvyatsov | 25 January 2023, 21:40:54 UTC | Improved halide_popcount (#7225) * Improved halide_popcount * reused popcount64 from Utils.cpp in CodeGen_C * Fixed comment for popcount | 25 January 2023, 21:40:54 UTC |
810bd0b | Andrew Adams | 21 January 2023, 22:08:30 UTC | Hoist vector slices using rewrite rules (#7243) * Hoist slices using rewrite rules This lets us add associative variants more easily, which are helpful in the work on staging strided loads. * Don't hoist extract_element shuffles The Shuffle visitor wants to sink them * Add some static asserts * Add explanatory comment on shuffle hoisting * Fix comment * add lanes predicate to slice hoisting * add vector slice hoisting test cases Co-authored-by: Steven Johnson <srj@google.com> Co-authored-by: Alexander <ajroot@stanford.edu> | 21 January 2023, 22:08:30 UTC |
bafd60f | Alexander Root | 20 January 2023, 18:03:25 UTC | [x86 & wasm] Split up double saturating-narrows from i32 (#7280) * better x86 double sat-cast + add test * fix wasm too + test Co-authored-by: Steven Johnson <srj@google.com> | 20 January 2023, 18:03:25 UTC |
c601e4e | Steven Johnson | 20 January 2023, 17:43:56 UTC | Add workaround for the const-or-not user_context issue (#635) (#7291) Add a workaround for the const-or-not user_context issue (https://github.com/halide/Halide/issues/635) | 20 January 2023, 17:43:56 UTC |
2cc0468 | Steve Suzuki | 20 January 2023, 17:39:41 UTC | Fix issue in add_halide_runtime in cross-compilation (#7284) * Fix issue in add_halide_runtime in cross-compilation add_halide_runtime() tries to build generator executable, but it fails if we are working with cross-compiler toolchain. By using existing generator set as "FROM", we can work around this. | 20 January 2023, 17:39:41 UTC |
d44e99d | Steve Suzuki | 20 January 2023, 13:12:30 UTC | Fix error of add_halide_generator in cross-compilation (#7283) In case the project name is CamelCase, add_halide_generator() was not able to find the generator package, because CMake searches <name>Config.cmake or <lower-case-name>-config.cmake | 20 January 2023, 13:12:30 UTC |
147ff48 | Alex Reinking | 20 January 2023, 12:54:34 UTC | Remove dependency on platform threads library (#7297) * Refactor internal ThreadPool.h into halide_thread_pool.h tool * Drop dependency of libHalide on threads library * Remove other redundant uses of Threads::Threads * Update CMake documentation. | 20 January 2023, 12:54:34 UTC |
314b2fd | Alexander Root | 20 January 2023, 00:35:14 UTC | [HVX] Fix EliminateInterleaves (#7279) * fix EliminateInterleaves Co-authored-by: Steven Johnson <srj@google.com> | 20 January 2023, 00:35:14 UTC |
c9f3602 | Steven Johnson | 19 January 2023, 23:48:26 UTC | Remove the watchdog timer from generator_main(). It was intended to k… (#7295) Remove the watchdog timer from generator_main(). It was intended to kill pathologically slow builds, but in the environment it was added for (Google build servers), it ended up being redundant to existing mechanisms, and removing it allows us to remove a dependency on threading libraries in libHalide. | 19 January 2023, 23:48:26 UTC |
51a4f6c | Steven Johnson | 19 January 2023, 23:36:47 UTC | Emit prototypes for destructor functions in C Backend (#7296) We gathered up the destructors, but only emitted the prototypes if there was at least one non-C++ function declaration needed -- so if you built with cpp_name_mangling enabled, you might omit the right prototype. Fixed and added the right flag to a Generator test to tickle this behavior. | 19 January 2023, 23:36:47 UTC |
e8e1481 | Steven Johnson | 18 January 2023, 21:56:14 UTC | Drop support for MIPS (#7287) (#7289) * Drop support for MIPS (#7287) * Update Target.cpp | 18 January 2023, 21:56:14 UTC |
888c41c | Steven Johnson | 18 January 2023, 00:47:52 UTC | Add CMake support for C++ backend in test/generator (#7274) * Add support for C++ backend in test/generator When the CMake rules were rewritten a while back, the support for building/testing generators with the C++ backend (instead of the standard LLVM, etc) got lost. This adds it back in. Also made some drive-by fixes to the Makefile to enable some tests there that work correctly now. Also made a drive-by fix in in Codegen_C to fix allocation nodes that were just wrappers around buffer_get_host -- this prevented the cleanup_on_error test from building with the C++ backend. | 18 January 2023, 00:47:52 UTC |
0d43318 | Steven Johnson | 10 January 2023, 19:23:58 UTC | Optimize Module::compile() for some edge cases (#7269) * Optimize Module::compile() for some edge cases Avoid redundant `compile_to_buffer()` calls for output requests that can't possibly ever need them. * Avoid mutation | 10 January 2023, 19:23:58 UTC |
a8d88bb | Steven Johnson | 10 January 2023, 17:51:38 UTC | Use ::aligned_alloc() instead of std::aligned_alloc() in HalideBuffer.h (#7268) | 10 January 2023, 17:51:38 UTC |
eea7696 | Steven Johnson | 09 January 2023, 17:58:04 UTC | Update README_python.md (#7266) | 09 January 2023, 17:58:04 UTC |
c070bb8 | Alina Sbirlea | 06 January 2023, 00:00:34 UTC | Update change following LLVM WASM change f841ad30d77eeb4c51663e68efefdb734c7a3d07 (#7264) * Update change following LLVM WASM change https://github.com/llvm/llvm-project/commit/f841ad30d77eeb4c51663e68efefdb734c7a3d07 * Update checks conditional on LLVM version. | 06 January 2023, 00:00:34 UTC |
4b74049 | Andrew Adams | 05 January 2023, 21:12:50 UTC | Inline into extern function args during bounds inference (#7261) * Inline into extern function args during bounds inference Fixes #7260 * Run CSE once at the end * Actually recursively inline * clang-tidy * trigger buildbots * Make test invariant to the number of times the warning is printed as long as it's at least once Co-authored-by: Steven Johnson <srj@google.com> | 05 January 2023, 21:12:50 UTC |
04bb986 | Steven Johnson | 28 December 2022, 17:31:01 UTC | Conditional allocations shouldn't fail for size=0 in C++ backend (#7255) (#7256) * Conditional allocations shouldn't fail for size=0 in C++ backend (#7255) Allocations can be conditional; if the condition evaluates to false, we end up calling `halide_malloc(0)` (or `halide_tcm_malloc(0)` in the xtensa branch). Since it's legal via spec for `malloc(0)` to return nullptr, we need to be cautious here: if we are compiling with assertions enabled, *and* have a malloc() (etc) implementation that returns nullptr for alloc(0), we need to skip the assertion check, since we know the result won't be used. Note: a similar check will be inserted in the xtensa branch separately. Note 2: LLVM backend already has this check via Codegen_Posix.cpp * Update CodeGen_C.cpp | 28 December 2022, 17:31:01 UTC |
ade8b56 | Steven Johnson | 20 December 2022, 20:05:19 UTC | Remove deprecated halide_target_feature_disable_llvm_loop_opt (#7247) * Remove deprecated halide_target_feature_disable_llvm_loop_opt Was deprecated in Halide 15; let's remove in Halide 16 * trigger buildbots * trigger buildbots * Update CodeGen_LLVM.cpp | 20 December 2022, 20:05:19 UTC |
10345d4 | Andrew Adams | 16 December 2022, 17:56:08 UTC | Explicitly stage strided loads (#7230) * Add a pass to do explicit densification of strided loads * densify more types of strided load * Reorder downsample in local laplacian for slightly better performance * Move allocation padding into the IR. Still WIP. * Simplify concat_bits handling * Use evidence from parent scopes to densify * Disallow padding allocations with custom new expressions * Add test for parent scopes * Remove debugging prints. Avoid nested ramps. * Avoid parent scope loops * Update cmakefiles * Fix for large_buffers * Pad stack allocations too * Restore vld2/3/4 generation on non-Apple ARM chips * Appease clang-format and clang-tidy * Silence clang-tidy * Better comments * Comment improvements * Nuke code that reads out of bounds * Fix stage_strided_loads test * Change strategy for loads from external buffers Some backends don't like non-power-of-two vectors. Do two overlapping half-sized loads and shuffle instead of one funny-sized load. * Add explanatory comment to ARM backend * Fix cpp backend shuffling * Fix missing msan annotations * Magnify heap cost effect in stack_vs_heap performance test * Address review comments * clang-tidy * Fix for when same load node occurs in two different allocate nodes | 16 December 2022, 17:56:08 UTC |
382f813 | Steven Johnson | 16 December 2022, 17:54:21 UTC | Fix "may be used uninitialized" warnings in Codegen_C::print_scalarized_expr() (#7244) | 16 December 2022, 17:54:21 UTC |
da6746e | Steven Johnson | 14 December 2022, 05:29:45 UTC | correctness/exception.cpp needs to check HALIDE_WITH_EXCEPTIONS (fixes #7240) (#7241) correctness/exception.cpp needs to check HALIDE_WITH_EXCEPTIONS | 14 December 2022, 05:29:45 UTC |
1a4a469 | Andrew Adams | 13 December 2022, 16:11:54 UTC | Fix some sources of signed integer overflow in the compiler (#7231) * Fix some sources of signed integer overflow in the compiler Also, use compiler intrinsics when possible to handle overflow, as it generates faster code. * Fix msvc macro * Must use result * Actually perform the requested operation | 13 December 2022, 16:11:54 UTC |
533e6e5 | Steven Johnson | 12 December 2022, 20:54:46 UTC | Remove rogue string suffix in simd_op_check_arm.cpp (#7227) * Remove rogue string suffix in simd_op_check_arm.cpp Interestingly, it compiles here, but in some compilers it will fail with "unexpected token". * Update simd_op_check_arm.cpp | 12 December 2022, 20:54:46 UTC |
6ecdcbd | Steven Johnson | 11 December 2022, 18:05:55 UTC | Tighten alignment promises for halide_malloc() (#7222) This makes a couple of changes to the behavior/implementation of `halide_malloc()`: * Currently, halide_malloc must return a pointer aligned to the maximum meaningful alignment for the platform for the purpose of vector loads and stores. This PR also adds the requirement that the memory returned must be legal to access in an integral multple of alignment >= the requested size (in other words: you should be able to do vector load/stores "off the end" without causing any faults). * Currently, the `halide_malloc_alignment()` function is used to determine the default alignment; this cannot be overridden by user code (well, it can be, but the override will have no useful effect). It is intended to be "internal only" but is used in at least one place outside the runtime (apps/hannk). This change removes the call entirely, in favor of a call that is harder to access from outside the runtime and much less likely for end users to attempt to call. (It also changes apps/hannk to stop using it.) | 11 December 2022, 18:05:55 UTC |
16421a7 | Steven Johnson | 09 December 2022, 17:21:30 UTC | Revise simd_op_check tests to ignore HL_TARGET (#7207) (#7216) * Revise simd_op_check tests to ignore HL_TARGET (#7207) The simd_op_check tests have historically only run using the value of HL_TARGET, which mean that the coverage they had was low (since HL_TARGET is only set to values that are runnable on at least one buildbot). This change completely disconnects these tests from HL_TARGET; instead, each test now tests for a range of targets appropriate to the architecture being tested. On all platforms, they still compile to assembly and verify that the correct instructions are generated; additionally, if the host platform can JIT for the given target, it verifies that the results are as expected. * Update simd_op_check_riscv.cpp * Update simd_op_check_x86.cpp * Update simd_op_check_x86.cpp * Update simd_op_check_arm.cpp * Add more features that must match; re-enable the bfloat instructions * Update simd_op_check_x86.cpp * Update simd_op_check_riscv.cpp * trigger buildbots * Fix simd_op_check_wasm | 09 December 2022, 17:21:30 UTC |
ba31688 | Steven Johnson | 09 December 2022, 01:22:54 UTC | Increase __clang_major__ check in Float16.h to 16 (#7224) | 09 December 2022, 01:22:54 UTC |
066559b | Steven Johnson | 08 December 2022, 23:35:01 UTC | Remove check_jit_user_context() from V8 bindings (#7220) Obsolete code from early V8 work, it can trigger inappropriately in some corner-case scenarios. Remove it entirely to avoid false errors. | 08 December 2022, 23:35:01 UTC |