diff --git a/.clang-format b/.clang-format index 1022bb0407e6..d69055f670f8 100644 --- a/.clang-format +++ b/.clang-format @@ -1,48 +1,39 @@ --- AccessModifierOffset: -4 AlignEscapedNewlines: Left -AlignTrailingComments: true +AlignTrailingComments: + Kind: Always + OverEmptyLines: 0 AllowAllParametersOfDeclarationOnNextLine: true AllowShortBlocksOnASingleLine: Empty AllowShortFunctionsOnASingleLine: None AllowShortIfStatementsOnASingleLine: Always AllowShortLoopsOnASingleLine: false AlwaysBreakBeforeMultilineStrings: false -AlwaysBreakTemplateDeclarations: Yes BinPackParameters: true +BreakAfterReturnType: Automatic BreakBeforeBinaryOperators: None BreakBeforeBraces: Attach BreakBeforeTernaryOperators: false -BreakConstructorInitializersBeforeComma: false +BreakConstructorInitializers: BeforeColon +BreakTemplateDeclarations: Yes ColumnLimit: 0 -ConstructorInitializerAllOnOneLineOrOnePerLine: false ConstructorInitializerIndentWidth: 4 ContinuationIndentWidth: 4 -Cpp11BracedListStyle: true -DerivePointerAlignment: false -ExperimentalAutoDetectBinPacking: false IndentCaseLabels: false -IndentWrappedFunctionNames: false IndentWidth: 4 +IndentWrappedFunctionNames: false MaxEmptyLinesToKeep: 1 NamespaceIndentation: None -ObjCSpaceBeforeProtocolList: true -PenaltyBreakBeforeFirstCallParameter: 19 -PenaltyBreakComment: 60 -PenaltyBreakFirstLessLess: 120 -PenaltyBreakString: 1000 -PenaltyExcessCharacter: 1000000 -PenaltyReturnTypeOnItsOwnLine: 60 +PackConstructorInitializers: BinPack PointerAlignment: Right -SpaceBeforeParens: ControlStatements SpaceAfterCStyleCast: false SpaceAfterTemplateKeyword: false SpaceBeforeAssignmentOperators: true -SpaceInEmptyParentheses: false +SpaceBeforeParens: ControlStatements SpacesBeforeTrailingComments: 2 -SpacesInAngles: false -SpacesInCStyleCastParentheses: false -SpacesInParentheses: false +SpacesInAngles: Never +SpacesInParens: Never Standard: c++17 TabWidth: 8 UseTab: Never diff --git a/.github/workflows/presubmit.yml b/.github/workflows/presubmit.yml index 1b6b549214f4..b58103a45c92 100644 --- a/.github/workflows/presubmit.yml +++ b/.github/workflows/presubmit.yml @@ -16,14 +16,15 @@ permissions: jobs: check_clang_format: name: Check clang-format - runs-on: ubuntu-20.04 + runs-on: macos-14 steps: - uses: actions/checkout@v3 - - uses: DoozyX/clang-format-lint-action@v0.17 - with: - source: '.' - extensions: 'h,c,cpp' - clangFormatVersion: 17 + - name: Install clang-format + run: brew install llvm@19 + - name: Check clang-format + run: ./run-clang-format.sh -c + env: + CLANG_FORMAT_LLVM_INSTALL_DIR: /opt/homebrew/opt/llvm@19 check_clang_tidy: name: Check clang-tidy runs-on: macos-14 diff --git a/apps/HelloAndroidCamera2/jni/LockedSurface.cpp b/apps/HelloAndroidCamera2/jni/LockedSurface.cpp index 6b8dac2b7d14..e5a32b01f9a5 100644 --- a/apps/HelloAndroidCamera2/jni/LockedSurface.cpp +++ b/apps/HelloAndroidCamera2/jni/LockedSurface.cpp @@ -5,7 +5,7 @@ // Round x up to a multiple of mask. // E.g., ALIGN(x, 16) means round x up to the nearest multiple of 16. -#define ALIGN(x, mask) (((x) + (mask)-1) & ~((mask)-1)) +#define ALIGN(x, mask) (((x) + (mask) - 1) & ~((mask) - 1)) LockedSurface *LockedSurface::lock(JNIEnv *env, jobject surface) { LockedSurface *output = new LockedSurface; diff --git a/apps/HelloWasm/core.cpp b/apps/HelloWasm/core.cpp index ae952e188bc6..2d7f8b38cdf5 100644 --- a/apps/HelloWasm/core.cpp +++ b/apps/HelloWasm/core.cpp @@ -76,10 +76,7 @@ void mainloop(void *arg) { "Frame rate: %2.0f fps", ctx->smoothed_runtime, ctx->smoothed_blit_time, ctx->smoothed_fps); // Run some javascript inline to update the web-page - EM_ASM({ - document.getElementById(UTF8ToString($0)).innerHTML = UTF8ToString($1); - }, - "runtime", buf); + EM_ASM({ document.getElementById(UTF8ToString($0)).innerHTML = UTF8ToString($1); }, "runtime", buf); // Read the threads slider from the UI int threads = EM_ASM_INT({ diff --git a/apps/blur/test.cpp b/apps/blur/test.cpp index 3ce299541c8a..565275b66185 100644 --- a/apps/blur/test.cpp +++ b/apps/blur/test.cpp @@ -70,61 +70,61 @@ Buffer blur_fast(Buffer in) { } } #elif __ARM_NEON - uint16x4_t one_third = vdup_n_u16(21846); + uint16x4_t one_third = vdup_n_u16(21846); #pragma omp parallel for - for (int yTile = 0; yTile < out.height(); yTile += 32) { - uint16x8_t tmp[(128 / 8) * (32 + 2)]; - for (int xTile = 0; xTile < out.width(); xTile += 128) { - uint16_t *tmpPtr = (uint16_t *)tmp; - for (int y = 0; y < 32 + 2; y++) { - const uint16_t *inPtr = &(in(xTile, yTile + y)); - for (int x = 0; x < 128; x += 8) { - uint16x8_t a = vld1q_u16(inPtr); - uint16x8_t b = vld1q_u16(inPtr + 1); - uint16x8_t c = vld1q_u16(inPtr + 2); - uint16x8_t sum = vaddq_u16(vaddq_u16(a, b), c); - uint16x4_t sumlo = vget_low_u16(sum); - uint16x4_t sumhi = vget_high_u16(sum); - uint16x4_t avglo = vshrn_n_u32(vmull_u16(sumlo, one_third), 16); - uint16x4_t avghi = vshrn_n_u32(vmull_u16(sumhi, one_third), 16); - uint16x8_t avg = vcombine_u16(avglo, avghi); - vst1q_u16(tmpPtr, avg); - tmpPtr += 8; - inPtr += 8; - } + for (int yTile = 0; yTile < out.height(); yTile += 32) { + uint16x8_t tmp[(128 / 8) * (32 + 2)]; + for (int xTile = 0; xTile < out.width(); xTile += 128) { + uint16_t *tmpPtr = (uint16_t *)tmp; + for (int y = 0; y < 32 + 2; y++) { + const uint16_t *inPtr = &(in(xTile, yTile + y)); + for (int x = 0; x < 128; x += 8) { + uint16x8_t a = vld1q_u16(inPtr); + uint16x8_t b = vld1q_u16(inPtr + 1); + uint16x8_t c = vld1q_u16(inPtr + 2); + uint16x8_t sum = vaddq_u16(vaddq_u16(a, b), c); + uint16x4_t sumlo = vget_low_u16(sum); + uint16x4_t sumhi = vget_high_u16(sum); + uint16x4_t avglo = vshrn_n_u32(vmull_u16(sumlo, one_third), 16); + uint16x4_t avghi = vshrn_n_u32(vmull_u16(sumhi, one_third), 16); + uint16x8_t avg = vcombine_u16(avglo, avghi); + vst1q_u16(tmpPtr, avg); + tmpPtr += 8; + inPtr += 8; } - tmpPtr = (uint16_t *)tmp; - for (int y = 0; y < 32; y++) { - uint16_t *outPtr = &(out(xTile, yTile + y)); - for (int x = 0; x < 128; x += 8) { - uint16x8_t a = vld1q_u16(tmpPtr + (2 * 128)); - uint16x8_t b = vld1q_u16(tmpPtr + 128); - uint16x8_t c = vld1q_u16(tmpPtr); - uint16x8_t sum = vaddq_u16(vaddq_u16(a, b), c); - uint16x4_t sumlo = vget_low_u16(sum); - uint16x4_t sumhi = vget_high_u16(sum); - uint16x4_t avglo = vshrn_n_u32(vmull_u16(sumlo, one_third), 16); - uint16x4_t avghi = vshrn_n_u32(vmull_u16(sumhi, one_third), 16); - uint16x8_t avg = vcombine_u16(avglo, avghi); - vst1q_u16(outPtr, avg); - tmpPtr += 8; - outPtr += 8; - } + } + tmpPtr = (uint16_t *)tmp; + for (int y = 0; y < 32; y++) { + uint16_t *outPtr = &(out(xTile, yTile + y)); + for (int x = 0; x < 128; x += 8) { + uint16x8_t a = vld1q_u16(tmpPtr + (2 * 128)); + uint16x8_t b = vld1q_u16(tmpPtr + 128); + uint16x8_t c = vld1q_u16(tmpPtr); + uint16x8_t sum = vaddq_u16(vaddq_u16(a, b), c); + uint16x4_t sumlo = vget_low_u16(sum); + uint16x4_t sumhi = vget_high_u16(sum); + uint16x4_t avglo = vshrn_n_u32(vmull_u16(sumlo, one_third), 16); + uint16x4_t avghi = vshrn_n_u32(vmull_u16(sumhi, one_third), 16); + uint16x8_t avg = vcombine_u16(avglo, avghi); + vst1q_u16(outPtr, avg); + tmpPtr += 8; + outPtr += 8; } } } + } #else - // No intrinsics enabled, do a naive thing. - for (int y = 0; y < out.height(); y++) { - for (int x = 0; x < out.width(); x++) { - int tmp[3] = { - (in(x, y) + in(x + 1, y) + in(x + 2, y)) / 3, - (in(x, y + 1) + in(x + 1, y + 1) + in(x + 2, y + 1)) / 3, - (in(x, y + 2) + in(x + 1, y + 2) + in(x + 2, y + 2)) / 3, - }; - out(x, y) = (tmp[0] + tmp[1] + tmp[2]) / 3; - } + // No intrinsics enabled, do a naive thing. + for (int y = 0; y < out.height(); y++) { + for (int x = 0; x < out.width(); x++) { + int tmp[3] = { + (in(x, y) + in(x + 1, y) + in(x + 2, y)) / 3, + (in(x, y + 1) + in(x + 1, y + 1) + in(x + 2, y + 1)) / 3, + (in(x, y + 2) + in(x + 1, y + 2) + in(x + 2, y + 2)) / 3, + }; + out(x, y) = (tmp[0] + tmp[1] + tmp[2]) / 3; } + } #endif }); diff --git a/apps/hexagon_dma/process_raw_linear_interleaved_basic.cpp b/apps/hexagon_dma/process_raw_linear_interleaved_basic.cpp index 483066b83ad3..e24b0b6f0538 100644 --- a/apps/hexagon_dma/process_raw_linear_interleaved_basic.cpp +++ b/apps/hexagon_dma/process_raw_linear_interleaved_basic.cpp @@ -42,9 +42,9 @@ typedef struct { #define _SCHEDULE_STR(s) #s #define _SCHEDULE_NAME(data, direction, schedule) pipeline_##data##_##direction##_##schedule #define _SCHEDULE_PAIR(data, direction, schedule) \ - { _SCHEDULE_STR(scheduled - pipeline(data, direction, schedule)), _SCHEDULE_NAME(data, direction, schedule) } + {_SCHEDULE_STR(scheduled - pipeline(data, direction, schedule)), _SCHEDULE_NAME(data, direction, schedule)} #define _SCHEDULE_DUMMY_PAIR \ - { NULL, NULL } + {NULL, NULL} #define SCHEDULE_FUNCTION_RW(schedule) _SCHEDULE_PAIR(raw_linear_interleaved, rw, schedule) #ifdef SCHEDULE_ALL diff --git a/apps/hexagon_dma/process_yuv_linear_basic.cpp b/apps/hexagon_dma/process_yuv_linear_basic.cpp index bfaf1e0207a9..fb9cc38ad94c 100644 --- a/apps/hexagon_dma/process_yuv_linear_basic.cpp +++ b/apps/hexagon_dma/process_yuv_linear_basic.cpp @@ -55,9 +55,9 @@ typedef struct { #define _SCHEDULE_STR(s) #s #define _SCHEDULE_NAME(data, direction, schedule) pipeline_##data##_##direction##_##schedule #define _SCHEDULE_PAIR(data, direction, schedule) \ - { _SCHEDULE_STR(scheduled - pipeline(data, direction, schedule)), _SCHEDULE_NAME(data, direction, schedule) } + {_SCHEDULE_STR(scheduled - pipeline(data, direction, schedule)), _SCHEDULE_NAME(data, direction, schedule)} #define _SCHEDULE_DUMMY_PAIR \ - { NULL, NULL } + {NULL, NULL} #define SCHEDULE_FUNCTION_RW(type, schedule) _SCHEDULE_PAIR(type##_linear, rw, schedule) #ifdef SCHEDULE_ALL diff --git a/apps/resnet_50/Resnet50Generator.cpp b/apps/resnet_50/Resnet50Generator.cpp index 7320a80b7e1d..b0040ae2a523 100644 --- a/apps/resnet_50/Resnet50Generator.cpp +++ b/apps/resnet_50/Resnet50Generator.cpp @@ -31,35 +31,35 @@ class Resnet50Generator : public Halide::Generator { Input> input{"input"}; /** parameter values for scaling layers **/ Input> conv1_gamma{"conv1_gamma"}; - Input[4]> br1_gamma { "br1_gamma" }; - Input[16]> br2a_gamma { "br2a_gamma" }; - Input[16]> br2b_gamma { "br2b_gamma" }; - Input[16]> br2c_gamma { "br2c_gamma" }; + Input[4]> br1_gamma{"br1_gamma"}; + Input[16]> br2a_gamma{"br2a_gamma"}; + Input[16]> br2b_gamma{"br2b_gamma"}; + Input[16]> br2c_gamma{"br2c_gamma"}; Input> conv1_beta{"conv1_beta"}; - Input[4]> br1_beta { "br1_beta" }; - Input[16]> br2a_beta { "br2a_beta" }; - Input[16]> br2b_beta { "br2b_beta" }; - Input[16]> br2c_beta { "br2c_beta" }; + Input[4]> br1_beta{"br1_beta"}; + Input[16]> br2a_beta{"br2a_beta"}; + Input[16]> br2b_beta{"br2b_beta"}; + Input[16]> br2c_beta{"br2c_beta"}; Input> conv1_mu{"conv1_mu"}; - Input[4]> br1_mu { "br1_mu" }; - Input[16]> br2a_mu { "br2a_mu" }; - Input[16]> br2b_mu { "br2b_mu" }; - Input[16]> br2c_mu { "br2c_mu" }; + Input[4]> br1_mu{"br1_mu"}; + Input[16]> br2a_mu{"br2a_mu"}; + Input[16]> br2b_mu{"br2b_mu"}; + Input[16]> br2c_mu{"br2c_mu"}; Input> conv1_sig{"conv1_sig"}; - Input[4]> br1_sig { "br1_sig" }; - Input[16]> br2a_sig { "br2a_sig" }; - Input[16]> br2b_sig { "br2b_sig" }; - Input[16]> br2c_sig { "br2c_sig" }; + Input[4]> br1_sig{"br1_sig"}; + Input[16]> br2a_sig{"br2a_sig"}; + Input[16]> br2b_sig{"br2b_sig"}; + Input[16]> br2c_sig{"br2c_sig"}; /** weights and biases for convolutions **/ Input> conv1_weights{"conv1_weights"}; - Input[4]> br1_conv_weights { "br1_conv_weights" }; - Input[16]> br2a_conv_weights { "br2a_conv_weights" }; - Input[16]> br2b_conv_weights { "br2b_conv_weights" }; - Input[16]> br2c_conv_weights { "br2c_conv_weights" }; + Input[4]> br1_conv_weights{"br1_conv_weights"}; + Input[16]> br2a_conv_weights{"br2a_conv_weights"}; + Input[16]> br2b_conv_weights{"br2b_conv_weights"}; + Input[16]> br2c_conv_weights{"br2c_conv_weights"}; Input> fc1000_weights{"fc1000_weights"}; Input> fc1000_bias{"fc1000_bias"}; diff --git a/python_bindings/src/halide/halide_/PyBuffer.cpp b/python_bindings/src/halide/halide_/PyBuffer.cpp index 6d1461a33ec4..5e6ca85d4b39 100644 --- a/python_bindings/src/halide/halide_/PyBuffer.cpp +++ b/python_bindings/src/halide/halide_/PyBuffer.cpp @@ -380,9 +380,9 @@ void define_buffer(py::module &m) { .def_static("make_scalar", (Buffer<>(*)(Type, const std::string &))Buffer<>::make_scalar, py::arg("type"), py::arg("name") = "") .def_static("make_interleaved", (Buffer<>(*)(Type, int, int, int, const std::string &))Buffer<>::make_interleaved, py::arg("type"), py::arg("width"), py::arg("height"), py::arg("channels"), py::arg("name") = "") - .def_static( + .def_static( // "make_with_shape_of", [](const Buffer<> &buffer, const std::string &name) -> Buffer<> { - return Buffer<>::make_with_shape_of(buffer, nullptr, nullptr, name); + return Buffer<>::make_with_shape_of(buffer, nullptr, nullptr, name); // }, py::arg("src"), py::arg("name") = "") @@ -401,22 +401,22 @@ void define_buffer(py::module &m) { .def("bottom", (int(Buffer<>::*)() const) & Buffer<>::bottom) .def("left", (int(Buffer<>::*)() const) & Buffer<>::left) .def("right", (int(Buffer<>::*)() const) & Buffer<>::right) - .def("number_of_elements", (size_t(Buffer<>::*)() const) & Buffer<>::number_of_elements) - .def("size_in_bytes", (size_t(Buffer<>::*)() const) & Buffer<>::size_in_bytes) + .def("number_of_elements", (size_t(Buffer<>::*)() const)&Buffer<>::number_of_elements) + .def("size_in_bytes", (size_t(Buffer<>::*)() const)&Buffer<>::size_in_bytes) .def("has_device_allocation", (bool(Buffer<>::*)() const) & Buffer<>::has_device_allocation) .def("host_dirty", (bool(Buffer<>::*)() const) & Buffer<>::host_dirty) .def("device_dirty", (bool(Buffer<>::*)() const) & Buffer<>::device_dirty) - .def( - "set_host_dirty", [](Buffer<> &b, bool dirty) -> void { - b.set_host_dirty(dirty); - }, - py::arg("dirty") = true) - .def( - "set_device_dirty", [](Buffer<> &b, bool dirty) -> void { - b.set_device_dirty(dirty); - }, - py::arg("dirty") = true) + .def("set_host_dirty", // + [](Buffer<> &b, bool dirty) -> void { + b.set_host_dirty(dirty); // + }, + py::arg("dirty") = true) + .def("set_device_dirty", // + [](Buffer<> &b, bool dirty) -> void { + b.set_device_dirty(dirty); // + }, + py::arg("dirty") = true) .def("copy", &Buffer<>::copy) .def("copy_from", &Buffer<>::copy_from::AnyDims>) @@ -426,27 +426,27 @@ void define_buffer(py::module &m) { for (int i = 0; i < b.dimensions(); i++) { order.push_back(d - i - 1); } - return b.transposed(order); + return b.transposed(order); // }) .def("add_dimension", (void(Buffer<>::*)()) & Buffer<>::add_dimension) .def("allocate", [](Buffer<> &b) -> void { - b.allocate(nullptr, nullptr); + b.allocate(nullptr, nullptr); // }) .def("deallocate", (void(Buffer<>::*)()) & Buffer<>::deallocate) .def("device_deallocate", (void(Buffer<>::*)()) & Buffer<>::device_deallocate) - .def( - "crop", [](Buffer<> &b, int d, int min, int extent) -> void { - b.crop(d, min, extent); - }, - py::arg("dimension"), py::arg("min"), py::arg("extent")) - .def( - "crop", [](Buffer<> &b, const std::vector> &rect) -> void { - b.crop(rect); - }, - py::arg("rect")) + .def("crop", // + [](Buffer<> &b, int d, int min, int extent) -> void { + b.crop(d, min, extent); // + }, + py::arg("dimension"), py::arg("min"), py::arg("extent")) + .def("crop", // + [](Buffer<> &b, const std::vector> &rect) -> void { + b.crop(rect); // + }, + py::arg("rect")) // Present in Runtime::Buffer but not Buffer // .def("cropped", [](Buffer<> &b, int d, int min, int extent) -> Buffer<> { @@ -456,60 +456,60 @@ void define_buffer(py::module &m) { // return b.cropped(rect); // }, py::arg("rect")) - .def( - "embed", [](Buffer<> &b, int d, int pos) -> void { - b.embed(d, pos); - }, - py::arg("dimension"), py::arg("pos")) - .def( - "embedded", [](Buffer<> &b, int d, int pos) -> Buffer<> { - return b.embedded(d, pos); - }, - py::arg("dimension"), py::arg("pos")) - - .def( - "embed", [](Buffer<> &b, int d) -> void { - b.embed(d); - }, - py::arg("dimension")) - .def( - "embedded", [](Buffer<> &b, int d) -> Buffer<> { - return b.embedded(d); - }, - py::arg("dimension")) - - .def( - "slice", [](Buffer<> &b, int d, int pos) -> void { - b.slice(d, pos); - }, - py::arg("dimension"), py::arg("pos")) - .def( - "sliced", [](Buffer<> &b, int d, int pos) -> Buffer<> { - return b.sliced(d, pos); - }, - py::arg("dimension"), py::arg("pos")) - - .def( - "slice", [](Buffer<> &b, int d) -> void { - b.slice(d); - }, - py::arg("dimension")) - .def( - "sliced", [](Buffer<> &b, int d) -> Buffer<> { - return b.sliced(d); - }, - py::arg("dimension")) - - .def( - "translate", [](Buffer<> &b, int d, int dx) -> void { - b.translate(d, dx); - }, - py::arg("dimension"), py::arg("dx")) - .def( - "translate", [](Buffer<> &b, const std::vector &delta) -> void { - b.translate(delta); - }, - py::arg("delta")) + .def("embed", // + [](Buffer<> &b, int d, int pos) -> void { + b.embed(d, pos); // + }, + py::arg("dimension"), py::arg("pos")) + .def("embedded", // + [](Buffer<> &b, int d, int pos) -> Buffer<> { + return b.embedded(d, pos); // + }, + py::arg("dimension"), py::arg("pos")) + + .def("embed", // + [](Buffer<> &b, int d) -> void { + b.embed(d); // + }, + py::arg("dimension")) + .def("embedded", // + [](Buffer<> &b, int d) -> Buffer<> { + return b.embedded(d); // + }, + py::arg("dimension")) + + .def("slice", // + [](Buffer<> &b, int d, int pos) -> void { + b.slice(d, pos); // + }, + py::arg("dimension"), py::arg("pos")) + .def("sliced", // + [](Buffer<> &b, int d, int pos) -> Buffer<> { + return b.sliced(d, pos); // + }, + py::arg("dimension"), py::arg("pos")) + + .def("slice", // + [](Buffer<> &b, int d) -> void { + b.slice(d); // + }, + py::arg("dimension")) + .def("sliced", // + [](Buffer<> &b, int d) -> Buffer<> { + return b.sliced(d); // + }, + py::arg("dimension")) + + .def("translate", // + [](Buffer<> &b, int d, int dx) -> void { + b.translate(d, dx); // + }, + py::arg("dimension"), py::arg("dx")) + .def("translate", // + [](Buffer<> &b, const std::vector &delta) -> void { + b.translate(delta); // + }, + py::arg("delta")) // Present in Runtime::Buffer but not Buffer // .def("translated", [](Buffer<> &b, int d, int dx) -> Buffer<> { @@ -519,49 +519,49 @@ void define_buffer(py::module &m) { // return b.translated(delta); // }, py::arg("delta")) - .def( - "transpose", [](Buffer<> &b, int d1, int d2) -> void { - b.transpose(d1, d2); - }, - py::arg("d1"), py::arg("d2")) - - .def( - "transposed", [](Buffer<> &b, int d1, int d2) -> Buffer<> { - return b.transposed(d1, d2); - }, - py::arg("d1"), py::arg("d2")) - - .def( - "transpose", [](Buffer<> &b, const std::vector &order) -> void { - b.transpose(order); - }, - py::arg("order")) - - .def( - "transposed", [](Buffer<> &b, const std::vector &order) -> Buffer<> { - return b.transposed(order); - }, - py::arg("order")) - - .def( - "dim", [](Buffer<> &b, int dimension) -> BufferDimension { - return b.dim(dimension); - }, - py::arg("dimension"), py::keep_alive<0, 1>() // Keep Buffer<> alive while Dimension exists - ) - - .def( - "for_each_element", [](Buffer<> &b, py::function f) -> void { - const int d = b.dimensions(); - std::vector pos_v(d, 0); - b.for_each_element([&f, &pos_v](const int *pos) -> void { - for (size_t i = 0; i < pos_v.size(); ++i) { - pos_v[i] = pos[i]; - } - f(pos_v); - }); - }, - py::arg("f")) + .def("transpose", // + [](Buffer<> &b, int d1, int d2) -> void { + b.transpose(d1, d2); // + }, + py::arg("d1"), py::arg("d2")) + + .def("transposed", // + [](Buffer<> &b, int d1, int d2) -> Buffer<> { + return b.transposed(d1, d2); // + }, + py::arg("d1"), py::arg("d2")) + + .def("transpose", // + [](Buffer<> &b, const std::vector &order) -> void { + b.transpose(order); // + }, + py::arg("order")) + + .def("transposed", // + [](Buffer<> &b, const std::vector &order) -> Buffer<> { + return b.transposed(order); // + }, + py::arg("order")) + + .def("dim", // + [](Buffer<> &b, int dimension) -> BufferDimension { + return b.dim(dimension); // + }, + py::arg("dimension"), py::keep_alive<0, 1>() // Keep Buffer<> alive while Dimension exists + ) + + .def("for_each_element", // + [](Buffer<> &b, py::function f) -> void { + const int d = b.dimensions(); + std::vector pos_v(d, 0); + b.for_each_element([&f, &pos_v](const int *pos) -> void { + for (size_t i = 0; i < pos_v.size(); ++i) { + pos_v[i] = pos[i]; + } + f(pos_v); + }); // + }, + py::arg("f")) .def("fill", &call_fill, py::arg("value")) .def("all_equal", &call_all_equal, py::arg("value")) @@ -573,78 +573,75 @@ void define_buffer(py::module &m) { // .def("for_each_value", [](Buffer<> &b, py::args f, py::args other_buffers) -> void { // }, py::arg("f")) - .def("copy_to_host", [](Buffer<> &b) -> int { + .def("copy_to_host", [](Buffer<> &b) -> int { // return b.copy_to_host(nullptr); }) - .def("device_detach_native", [](Buffer<> &b) -> int { + .def("device_detach_native", [](Buffer<> &b) -> int { // return b.device_detach_native(nullptr); }) - .def("device_free", [](Buffer<> &b) -> int { + .def("device_free", [](Buffer<> &b) -> int { // return b.device_free(nullptr); }) - .def("device_sync", [](Buffer<> &b) -> int { + .def("device_sync", [](Buffer<> &b) -> int { // return b.device_sync(nullptr); }) - .def( - "copy_to_device", [](Buffer<> &b, const Target &target) -> int { - return b.copy_to_device(to_jit_target(target)); - }, - py::arg("target") = Target()) - - .def( - "copy_to_device", [](Buffer<> &b, const DeviceAPI &d, const Target &target) -> int { - return b.copy_to_device(d, to_jit_target(target)); - }, - py::arg("device_api"), py::arg("target") = Target()) - .def( - "device_malloc", [](Buffer<> &b, const Target &target) -> int { - return b.device_malloc(to_jit_target(target)); - }, - py::arg("target") = Target()) - - .def( - "device_malloc", [](Buffer<> &b, const DeviceAPI &d, const Target &target) -> int { - return b.device_malloc(d, to_jit_target(target)); - }, - py::arg("device_api"), py::arg("target") = Target()) - - .def( - "set_min", [](Buffer<> &b, const std::vector &mins) -> void { + .def("copy_to_device", [](Buffer<> &b, const Target &target) -> int { // + return b.copy_to_device(to_jit_target(target)); + }, + py::arg("target") = Target()) + + .def("copy_to_device", [](Buffer<> &b, const DeviceAPI &d, const Target &target) -> int { // + return b.copy_to_device(d, to_jit_target(target)); + }, + py::arg("device_api"), py::arg("target") = Target()) + .def("device_malloc", [](Buffer<> &b, const Target &target) -> int { // + return b.device_malloc(to_jit_target(target)); + }, + py::arg("target") = Target()) + + .def("device_malloc", [](Buffer<> &b, const DeviceAPI &d, const Target &target) -> int { // + return b.device_malloc(d, to_jit_target(target)); + }, + py::arg("device_api"), py::arg("target") = Target()) + + .def( // + "set_min", // + [](Buffer<> &b, const std::vector &mins) -> void { if (mins.size() > (size_t)b.dimensions()) { throw py::value_error("Too many arguments"); } - b.set_min(mins); + b.set_min(mins); // }, py::arg("mins")) - .def( + .def( // "contains", [](Buffer<> &b, const std::vector &coords) -> bool { if (coords.size() > (size_t)b.dimensions()) { throw py::value_error("Too many arguments"); } - return b.contains(coords); + return b.contains(coords); // }, py::arg("coords")) - .def("__getitem__", [](Buffer<> &buf, const int &pos) -> py::object { + .def("__getitem__", [](Buffer<> &buf, const int &pos) -> py::object { // return buffer_getitem_operator(buf, {pos}); }) - .def("__getitem__", [](Buffer<> &buf, const std::vector &pos) -> py::object { + .def("__getitem__", [](Buffer<> &buf, const std::vector &pos) -> py::object { // return buffer_getitem_operator(buf, pos); }) - .def("__getitem__", [](Buffer<> &buf, const Expr &pos) -> Expr { + .def("__getitem__", [](Buffer<> &buf, const Expr &pos) -> Expr { // return buf(std::vector{pos}); }) - .def("__getitem__", [](Buffer<> &buf, const std::vector &pos) -> Expr { + .def("__getitem__", [](Buffer<> &buf, const std::vector &pos) -> Expr { // return buf(pos); }) - .def("__setitem__", [](Buffer<> &buf, const int &pos, const py::object &value) -> py::object { + .def("__setitem__", [](Buffer<> &buf, const int &pos, const py::object &value) -> py::object { // return buffer_setitem_operator(buf, {pos}, value); }) - .def("__setitem__", [](Buffer<> &buf, const std::vector &pos, const py::object &value) -> py::object { + .def("__setitem__", [](Buffer<> &buf, const std::vector &pos, const py::object &value) -> py::object { // return buffer_setitem_operator(buf, pos, value); }) @@ -655,7 +652,7 @@ void define_buffer(py::module &m) { } else { o << ""; } - return o.str(); + return o.str(); // }); } diff --git a/python_bindings/src/halide/halide_/PyDerivative.cpp b/python_bindings/src/halide/halide_/PyDerivative.cpp index 8ffdd5b80632..e6cc8214b706 100644 --- a/python_bindings/src/halide/halide_/PyDerivative.cpp +++ b/python_bindings/src/halide/halide_/PyDerivative.cpp @@ -11,27 +11,28 @@ void define_derivative(py::module &m) { return d(func); }, py::arg("func")) - .def( - "__getitem__", [](const Derivative &d, const Buffer<> &buffer) { - return d(buffer); - }, - py::arg("buffer")) - .def( - "__getitem__", [](const Derivative &d, const Param<> ¶m) { - return d(param); - }, - py::arg("param")) - .def("__getitem__", [](const Derivative &d, const std::tuple &args) { - return d(std::get<0>(args), std::get<1>(args)); - }); + .def("__getitem__", // + [](const Derivative &d, const Buffer<> &buffer) { + return d(buffer); // + }, + py::arg("buffer")) + .def("__getitem__", // + [](const Derivative &d, const Param<> ¶m) { + return d(param); // + }, + py::arg("param")) + .def("__getitem__", // + [](const Derivative &d, const std::tuple &args) { + return d(std::get<0>(args), std::get<1>(args)); // + }); m.def("propagate_adjoints", - (Derivative(*)(const Func &, const Func &, const Region &)) & propagate_adjoints); + (Derivative(*)(const Func &, const Func &, const Region &))&propagate_adjoints); m.def("propagate_adjoints", - (Derivative(*)(const Func &, const Buffer &)) & propagate_adjoints); + (Derivative(*)(const Func &, const Buffer &))&propagate_adjoints); m.def("propagate_adjoints", - (Derivative(*)(const Func &)) & propagate_adjoints); + (Derivative(*)(const Func &))&propagate_adjoints); } } // namespace PythonBindings -} // namespace Halide \ No newline at end of file +} // namespace Halide diff --git a/python_bindings/src/halide/halide_/PyFunc.cpp b/python_bindings/src/halide/halide_/PyFunc.cpp index 495187581aaa..67616ba10251 100644 --- a/python_bindings/src/halide/halide_/PyFunc.cpp +++ b/python_bindings/src/halide/halide_/PyFunc.cpp @@ -227,103 +227,103 @@ void define_func(py::module &m) { .def("store_in", &Func::store_in, py::arg("memory_type")) - .def( - "compile_to", [](Func &f, const std::map &output_files, const std::vector &args, const std::string &fn_name, const Target &target) { - f.compile_to(output_files, args, fn_name, to_aot_target(target)); - }, - py::arg("outputs"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_bitcode", [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - f.compile_to_bitcode(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_bitcode", [](Func &f, const std::string &filename, const std::vector &args, const Target &target) { - f.compile_to_bitcode(filename, args, "", to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) - .def( - "compile_to_llvm_assembly", [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - f.compile_to_llvm_assembly(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_llvm_assembly", [](Func &f, const std::string &filename, const std::vector &args, const Target &target) { - f.compile_to_llvm_assembly(filename, args, "", to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) - .def( - "compile_to_object", [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - f.compile_to_object(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_object", [](Func &f, const std::string &filename, const std::vector &args, const Target &target) { - f.compile_to_object(filename, args, "", to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) - .def( - "compile_to_header", [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - f.compile_to_header(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) - .def( - "compile_to_assembly", [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - f.compile_to_assembly(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_assembly", [](Func &f, const std::string &filename, const std::vector &args, const Target &target) { - f.compile_to_assembly(filename, args, "", to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) - .def( - "compile_to_c", [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - f.compile_to_c(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) - .def( - "compile_to_lowered_stmt", [](Func &f, const std::string &filename, const std::vector &args, StmtOutputFormat fmt, const Target &target) { - f.compile_to_lowered_stmt(filename, args, fmt, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fmt") = Text, py::arg("target") = Target()) - .def( - "compile_to_conceptual_stmt", [](Func &f, const std::string &filename, const std::vector &args, StmtOutputFormat fmt, const Target &target) { - f.compile_to_conceptual_stmt(filename, args, fmt, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fmt") = Text, py::arg("target") = Target()) - .def( - "compile_to_file", [](Func &f, const std::string &filename_prefix, const std::vector &args, const std::string &fn_name, const Target &target) { - f.compile_to_file(filename_prefix, args, fn_name, to_aot_target(target)); - }, - py::arg("filename_prefix"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) - .def( - "compile_to_static_library", [](Func &f, const std::string &filename_prefix, const std::vector &args, const std::string &fn_name, const Target &target) { - f.compile_to_static_library(filename_prefix, args, fn_name, to_aot_target(target)); - }, - py::arg("filename_prefix"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) + .def("compile_to", // + [](Func &f, const std::map &output_files, const std::vector &args, const std::string &fn_name, const Target &target) { + f.compile_to(output_files, args, fn_name, to_aot_target(target)); // + }, + py::arg("outputs"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) + .def("compile_to_bitcode", // + [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + f.compile_to_bitcode(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) + .def("compile_to_bitcode", // + [](Func &f, const std::string &filename, const std::vector &args, const Target &target) { + f.compile_to_bitcode(filename, args, "", to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) + .def("compile_to_llvm_assembly", // + [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + f.compile_to_llvm_assembly(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) + .def("compile_to_llvm_assembly", // + [](Func &f, const std::string &filename, const std::vector &args, const Target &target) { + f.compile_to_llvm_assembly(filename, args, "", to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) + .def("compile_to_object", // + [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + f.compile_to_object(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) + .def("compile_to_object", // + [](Func &f, const std::string &filename, const std::vector &args, const Target &target) { + f.compile_to_object(filename, args, "", to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) + .def("compile_to_header", // + [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + f.compile_to_header(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) + .def("compile_to_assembly", // + [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + f.compile_to_assembly(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) + .def("compile_to_assembly", // + [](Func &f, const std::string &filename, const std::vector &args, const Target &target) { + f.compile_to_assembly(filename, args, "", to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) + .def("compile_to_c", // + [](Func &f, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + f.compile_to_c(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) + .def("compile_to_lowered_stmt", // + [](Func &f, const std::string &filename, const std::vector &args, StmtOutputFormat fmt, const Target &target) { + f.compile_to_lowered_stmt(filename, args, fmt, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fmt") = Text, py::arg("target") = Target()) + .def("compile_to_conceptual_stmt", // + [](Func &f, const std::string &filename, const std::vector &args, StmtOutputFormat fmt, const Target &target) { + f.compile_to_conceptual_stmt(filename, args, fmt, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fmt") = Text, py::arg("target") = Target()) + .def("compile_to_file", // + [](Func &f, const std::string &filename_prefix, const std::vector &args, const std::string &fn_name, const Target &target) { + f.compile_to_file(filename_prefix, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename_prefix"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) + .def("compile_to_static_library", // + [](Func &f, const std::string &filename_prefix, const std::vector &args, const std::string &fn_name, const Target &target) { + f.compile_to_static_library(filename_prefix, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename_prefix"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) .def("compile_to_multitarget_static_library", &Func::compile_to_multitarget_static_library, py::arg("filename_prefix"), py::arg("arguments"), py::arg("targets")) .def("compile_to_multitarget_object_files", &Func::compile_to_multitarget_object_files, py::arg("filename_prefix"), py::arg("arguments"), py::arg("targets"), py::arg("suffixes")) // TODO: useless until Module is defined. - .def( - "compile_to_module", [](Func &f, const std::vector &args, const std::string &fn_name, const Target &target) -> Module { - return f.compile_to_module(args, fn_name, to_aot_target(target)); - }, - py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) - - .def( - "compile_jit", [](Func &f, const Target &target) { - f.compile_jit(to_jit_target(target)); - }, - py::arg("target") = Target()) - - .def( - "compile_to_callable", [](Func &f, const std::vector &args, const Target &target) { - return f.compile_to_callable(args, to_jit_target(target)); - }, - py::arg("arguments"), py::arg("target") = Target()) + .def("compile_to_module", // + [](Func &f, const std::vector &args, const std::string &fn_name, const Target &target) -> Module { + return f.compile_to_module(args, fn_name, to_aot_target(target)); // + }, + py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) + + .def("compile_jit", // + [](Func &f, const Target &target) { + f.compile_jit(to_jit_target(target)); // + }, + py::arg("target") = Target()) + + .def("compile_to_callable", // + [](Func &f, const std::vector &args, const Target &target) { + return f.compile_to_callable(args, to_jit_target(target)); // + }, + py::arg("arguments"), py::arg("target") = Target()) .def("has_update_definition", &Func::has_update_definition) .def("num_update_definitions", &Func::num_update_definitions) @@ -332,11 +332,11 @@ void define_func(py::module &m) { .def("update_args", &Func::update_args, py::arg("idx") = 0) .def("update_value", &Func::update_value, py::arg("idx") = 0) .def("update_value", &Func::update_value, py::arg("idx") = 0) - .def( - "update_values", [](Func &func, int idx) -> py::tuple { - return to_python_tuple(func.update_values(idx)); - }, - py::arg("idx") = 0) + .def("update_values", // + [](Func &func, int idx) -> py::tuple { + return to_python_tuple(func.update_values(idx)); // + }, + py::arg("idx") = 0) .def("rvars", &Func::rvars, py::arg("idx") = 0) .def("trace_loads", &Func::trace_loads) @@ -351,57 +351,65 @@ void define_func(py::module &m) { .def("is_extern", &Func::is_extern) .def("extern_function_name", &Func::extern_function_name) - .def("define_extern", (void(Func::*)(const std::string &, const std::vector &, const std::vector &, const std::vector &, NameMangling, DeviceAPI)) & Func::define_extern, py::arg("function_name"), py::arg("params"), py::arg("types"), py::arg("arguments"), py::arg("mangling") = NameMangling::Default, py::arg("device_api") = DeviceAPI::Host) + .def("define_extern", // + (void(Func::*)(const std::string &, const std::vector &, const std::vector &, const std::vector &, NameMangling, DeviceAPI)) & Func::define_extern, // + py::arg("function_name"), py::arg("params"), py::arg("types"), py::arg("arguments"), py::arg("mangling") = NameMangling::Default, py::arg("device_api") = DeviceAPI::Host) - .def("define_extern", (void(Func::*)(const std::string &, const std::vector &, Type, int, NameMangling, DeviceAPI)) & Func::define_extern, py::arg("function_name"), py::arg("params"), py::arg("type"), py::arg("dimensionality"), py::arg("mangling") = NameMangling::Default, py::arg("device_api") = DeviceAPI::Host) + .def("define_extern", // + (void(Func::*)(const std::string &, const std::vector &, Type, int, NameMangling, DeviceAPI)) & Func::define_extern, // + py::arg("function_name"), py::arg("params"), py::arg("type"), py::arg("dimensionality"), py::arg("mangling") = NameMangling::Default, py::arg("device_api") = DeviceAPI::Host) - .def("define_extern", (void(Func::*)(const std::string &, const std::vector &, const std::vector &, int, NameMangling, DeviceAPI)) & Func::define_extern, py::arg("function_name"), py::arg("params"), py::arg("types"), py::arg("dimensionality"), py::arg("mangling") = NameMangling::Default, py::arg("device_api") = DeviceAPI::Host) + .def("define_extern", // + (void(Func::*)(const std::string &, const std::vector &, const std::vector &, int, NameMangling, DeviceAPI)) & Func::define_extern, // + py::arg("function_name"), py::arg("params"), py::arg("types"), py::arg("dimensionality"), py::arg("mangling") = NameMangling::Default, py::arg("device_api") = DeviceAPI::Host) - .def("define_extern", (void(Func::*)(const std::string &, const std::vector &, Type, const std::vector &, NameMangling, DeviceAPI)) & Func::define_extern, py::arg("function_name"), py::arg("params"), py::arg("type"), py::arg("arguments"), py::arg("mangling") = NameMangling::Default, py::arg("device_api") = DeviceAPI::Host) + .def("define_extern", // + (void(Func::*)(const std::string &, const std::vector &, Type, const std::vector &, NameMangling, DeviceAPI)) & Func::define_extern, // + py::arg("function_name"), py::arg("params"), py::arg("type"), py::arg("arguments"), py::arg("mangling") = NameMangling::Default, py::arg("device_api") = DeviceAPI::Host) .def("output_buffer", &Func::output_buffer) .def("output_buffers", &Func::output_buffers) - .def( - "infer_input_bounds", [](Func &f, const py::object &dst, const Target &target) -> void { - const Target t = to_jit_target(target); - PyJITUserContext juc; - - // dst could be Buffer<>, vector, or vector - try { - Buffer<> b = dst.cast>(); - f.infer_input_bounds(&juc, b, t); - return; - } catch (...) { - // fall thru - } - - try { - std::vector> v = dst.cast>>(); - f.infer_input_bounds(&juc, Realization(std::move(v)), t); - return; - } catch (...) { - // fall thru - } - - try { - std::vector v = dst.cast>(); - f.infer_input_bounds(&juc, v, t); - return; - } catch (...) { - // fall thru - } - - throw py::value_error("Invalid arguments to infer_input_bounds"); - }, - py::arg("dst"), py::arg("target") = Target()) - - .def("in_", (Func(Func::*)(const Func &)) & Func::in, py::arg("f")) - .def("in_", (Func(Func::*)(const std::vector &fs)) & Func::in, py::arg("fs")) - .def("in_", (Func(Func::*)()) & Func::in) - - .def("clone_in", (Func(Func::*)(const Func &)) & Func::clone_in, py::arg("f")) - .def("clone_in", (Func(Func::*)(const std::vector &fs)) & Func::clone_in, py::arg("fs")) + .def("infer_input_bounds", // + [](Func &f, const py::object &dst, const Target &target) -> void { + const Target t = to_jit_target(target); + PyJITUserContext juc; + + // dst could be Buffer<>, vector, or vector + try { + Buffer<> b = dst.cast>(); + f.infer_input_bounds(&juc, b, t); + return; + } catch (...) { + // fall thru + } + + try { + std::vector> v = dst.cast>>(); + f.infer_input_bounds(&juc, Realization(std::move(v)), t); + return; + } catch (...) { + // fall thru + } + + try { + std::vector v = dst.cast>(); + f.infer_input_bounds(&juc, v, t); + return; + } catch (...) { + // fall thru + } + + throw py::value_error("Invalid arguments to infer_input_bounds"); // + }, + py::arg("dst"), py::arg("target") = Target()) + + .def("in_", (Func(Func::*)(const Func &))&Func::in, py::arg("f")) + .def("in_", (Func(Func::*)(const std::vector &fs))&Func::in, py::arg("fs")) + .def("in_", (Func(Func::*)())&Func::in) + + .def("clone_in", (Func(Func::*)(const Func &))&Func::clone_in, py::arg("f")) + .def("clone_in", (Func(Func::*)(const std::vector &fs))&Func::clone_in, py::arg("fs")) .def("copy_to_device", &Func::copy_to_device, py::arg("device_api") = DeviceAPI::Default_GPU) .def("copy_to_host", &Func::copy_to_host) @@ -423,7 +431,7 @@ void define_func(py::module &m) { .def("__repr__", [](const Func &func) -> std::string { std::ostringstream o; o << ""; - return o.str(); + return o.str(); // }); py::implicitly_convertible(); diff --git a/python_bindings/src/halide/halide_/PyIROperator.cpp b/python_bindings/src/halide/halide_/PyIROperator.cpp index 430db4622ec4..05d01e1e5a66 100644 --- a/python_bindings/src/halide/halide_/PyIROperator.cpp +++ b/python_bindings/src/halide/halide_/PyIROperator.cpp @@ -128,7 +128,7 @@ void define_operators(py::module &m) { return py::cast(false_expr_value); }); - m.def("mux", (Expr(*)(const Expr &, const std::vector &)) & mux); + m.def("mux", (Expr(*)(const Expr &, const std::vector &))&mux); m.def("sin", &sin); m.def("asin", &asin); @@ -163,8 +163,8 @@ void define_operators(py::module &m) { m.def("is_nan", &is_nan); m.def("is_inf", &is_inf); m.def("is_finite", &is_finite); - m.def("reinterpret", (Expr(*)(Type, Expr)) & reinterpret); - m.def("cast", (Expr(*)(Type, Expr)) & cast); + m.def("reinterpret", (Expr(*)(Type, Expr))&reinterpret); + m.def("cast", (Expr(*)(Type, Expr))&cast); m.def("print", [](const py::args &args) -> Expr { return print(collect_print_args(args)); }); @@ -186,13 +186,13 @@ void define_operators(py::module &m) { m.def("count_trailing_zeros", &count_trailing_zeros); m.def("div_round_to_zero", &div_round_to_zero); m.def("mod_round_to_zero", &mod_round_to_zero); - m.def("random_float", (Expr(*)()) & random_float); - m.def("random_uint", (Expr(*)()) & random_uint); - m.def("random_int", (Expr(*)()) & random_int); - m.def("random_float", (Expr(*)(Expr)) & random_float, py::arg("seed")); - m.def("random_uint", (Expr(*)(Expr)) & random_uint, py::arg("seed")); - m.def("random_int", (Expr(*)(Expr)) & random_int, py::arg("seed")); - m.def("undef", (Expr(*)(Type)) & undef); + m.def("random_float", (Expr(*)())&random_float); + m.def("random_uint", (Expr(*)())&random_uint); + m.def("random_int", (Expr(*)())&random_int); + m.def("random_float", (Expr(*)(Expr))&random_float, py::arg("seed")); + m.def("random_uint", (Expr(*)(Expr))&random_uint, py::arg("seed")); + m.def("random_int", (Expr(*)(Expr))&random_int, py::arg("seed")); + m.def("undef", (Expr(*)(Type))&undef); m.def( "memoize_tag", [](const Expr &result, const py::args &cache_key_values) -> Expr { return Internal::memoize_tag_helper(result, args_to_vector(cache_key_values)); @@ -200,7 +200,7 @@ void define_operators(py::module &m) { py::arg("result")); m.def("likely", &likely); m.def("likely_if_innermost", &likely_if_innermost); - m.def("saturating_cast", (Expr(*)(Type, Expr)) & saturating_cast); + m.def("saturating_cast", (Expr(*)(Type, Expr))&saturating_cast); m.def("strict_float", &strict_float); m.def("target_arch_is", &target_arch_is); m.def("target_bits", &target_bits); diff --git a/python_bindings/src/halide/halide_/PyImageParam.cpp b/python_bindings/src/halide/halide_/PyImageParam.cpp index 1b989b5443f8..15985070f134 100644 --- a/python_bindings/src/halide/halide_/PyImageParam.cpp +++ b/python_bindings/src/halide/halide_/PyImageParam.cpp @@ -19,7 +19,7 @@ void define_image_param(py::module &m) { .def("set_stride", &Dimension::set_stride, py::arg("stride")) .def("set_bounds", &Dimension::set_bounds, py::arg("min"), py::arg("extent")) .def("set_estimate", &Dimension::set_estimate, py::arg("min"), py::arg("extent")) - .def("dim", (Dimension(Dimension::*)(int)) & Dimension::dim, py::arg("dimension"), py::keep_alive<0, 1>()); + .def("dim", (Dimension(Dimension::*)(int))&Dimension::dim, py::arg("dimension"), py::keep_alive<0, 1>()); auto output_image_param_class = py::class_(m, "OutputImageParam") @@ -27,7 +27,7 @@ void define_image_param(py::module &m) { .def("name", &OutputImageParam::name) .def("type", &OutputImageParam::type) .def("defined", &OutputImageParam::defined) - .def("dim", (Dimension(OutputImageParam::*)(int)) & OutputImageParam::dim, py::arg("dimension"), py::keep_alive<0, 1>()) + .def("dim", (Dimension(OutputImageParam::*)(int))&OutputImageParam::dim, py::arg("dimension"), py::keep_alive<0, 1>()) .def("host_alignment", &OutputImageParam::host_alignment) .def("set_estimates", &OutputImageParam::set_estimates, py::arg("estimates")) .def("set_host_alignment", &OutputImageParam::set_host_alignment) @@ -72,9 +72,9 @@ void define_image_param(py::module &m) { .def("__getitem__", [](ImageParam &im, const std::vector &args) -> Expr { return im(args); }) - .def("in_", (Func(ImageParam::*)(const Func &)) & ImageParam::in) - .def("in_", (Func(ImageParam::*)(const std::vector &)) & ImageParam::in) - .def("in_", (Func(ImageParam::*)()) & ImageParam::in) + .def("in_", (Func(ImageParam::*)(const Func &))&ImageParam::in) + .def("in_", (Func(ImageParam::*)(const std::vector &))&ImageParam::in) + .def("in_", (Func(ImageParam::*)())&ImageParam::in) .def("trace_loads", &ImageParam::trace_loads) .def("__repr__", [](const ImageParam &im) -> std::string { diff --git a/python_bindings/src/halide/halide_/PyInlineReductions.cpp b/python_bindings/src/halide/halide_/PyInlineReductions.cpp index 7597bb022d99..315c87b37d7d 100644 --- a/python_bindings/src/halide/halide_/PyInlineReductions.cpp +++ b/python_bindings/src/halide/halide_/PyInlineReductions.cpp @@ -6,24 +6,24 @@ namespace Halide { namespace PythonBindings { void define_inline_reductions(py::module &m) { - m.def("sum", (Expr(*)(Expr, const std::string &s)) & Halide::sum, + m.def("sum", (Expr(*)(Expr, const std::string &s))&Halide::sum, py::arg("expr"), py::arg("name") = "sum"); - m.def("sum", (Expr(*)(const RDom &, Expr, const std::string &s)) & Halide::sum, + m.def("sum", (Expr(*)(const RDom &, Expr, const std::string &s))&Halide::sum, py::arg("rdom"), py::arg("expr"), py::arg("name") = "sum"); - m.def("product", (Expr(*)(Expr, const std::string &s)) & Halide::product, + m.def("product", (Expr(*)(Expr, const std::string &s))&Halide::product, py::arg("expr"), py::arg("name") = "product"); - m.def("product", (Expr(*)(const RDom &, Expr, const std::string &s)) & Halide::product, + m.def("product", (Expr(*)(const RDom &, Expr, const std::string &s))&Halide::product, py::arg("rdom"), py::arg("expr"), py::arg("name") = "product"); - m.def("maximum", (Expr(*)(Expr, const std::string &s)) & Halide::maximum, + m.def("maximum", (Expr(*)(Expr, const std::string &s))&Halide::maximum, py::arg("expr"), py::arg("name") = "maximum"); - m.def("maximum", (Expr(*)(const RDom &, Expr, const std::string &s)) & Halide::maximum, + m.def("maximum", (Expr(*)(const RDom &, Expr, const std::string &s))&Halide::maximum, py::arg("rdom"), py::arg("expr"), py::arg("name") = "maximum"); - m.def("minimum", (Expr(*)(Expr, const std::string &s)) & Halide::minimum, + m.def("minimum", (Expr(*)(Expr, const std::string &s))&Halide::minimum, py::arg("expr"), py::arg("name") = "minimum"); - m.def("minimum", (Expr(*)(const RDom &, Expr, const std::string &s)) & Halide::minimum, + m.def("minimum", (Expr(*)(const RDom &, Expr, const std::string &s))&Halide::minimum, py::arg("rdom"), py::arg("expr"), py::arg("name") = "minimum"); m.def( diff --git a/python_bindings/src/halide/halide_/PyModule.cpp b/python_bindings/src/halide/halide_/PyModule.cpp index 43d8e0993af9..0b06bd260802 100644 --- a/python_bindings/src/halide/halide_/PyModule.cpp +++ b/python_bindings/src/halide/halide_/PyModule.cpp @@ -60,9 +60,9 @@ void define_module(py::module &m) { }); m.def("link_modules", &link_modules, py::arg("name"), py::arg("modules")); - m.def("compile_standalone_runtime", (void (*)(const std::string &, const Target &)) & compile_standalone_runtime, py::arg("filename"), py::arg("target")); + m.def("compile_standalone_runtime", (void (*)(const std::string &, const Target &))&compile_standalone_runtime, py::arg("filename"), py::arg("target")); using OutputMap = std::map; - m.def("compile_standalone_runtime", (OutputMap(*)(const OutputMap &, const Target &)) & compile_standalone_runtime, py::arg("outputs"), py::arg("target")); + m.def("compile_standalone_runtime", (OutputMap(*)(const OutputMap &, const Target &))&compile_standalone_runtime, py::arg("outputs"), py::arg("target")); // TODO: compile_multitarget() deliberately skipped for now. } diff --git a/python_bindings/src/halide/halide_/PyParam.cpp b/python_bindings/src/halide/halide_/PyParam.cpp index ac5c13d45d35..0294d918cb81 100644 --- a/python_bindings/src/halide/halide_/PyParam.cpp +++ b/python_bindings/src/halide/halide_/PyParam.cpp @@ -28,11 +28,11 @@ void add_param_methods(py::class_> ¶m_class) { param.set(value); }, py::arg("value")) - .def( - "set_estimate", [](Param<> ¶m, TYPE value) -> void { - param.set_estimate(value); - }, - py::arg("value")); + .def("set_estimate", // + [](Param<> ¶m, TYPE value) -> void { + param.set_estimate(value); // + }, + py::arg("value")); } } // namespace diff --git a/python_bindings/src/halide/halide_/PyPipeline.cpp b/python_bindings/src/halide/halide_/PyPipeline.cpp index aeac06fe8092..b636de6b9461 100644 --- a/python_bindings/src/halide/halide_/PyPipeline.cpp +++ b/python_bindings/src/halide/halide_/PyPipeline.cpp @@ -72,7 +72,7 @@ void define_pipeline(py::module &m) { .def("outputs", &Pipeline::outputs) - .def("apply_autoscheduler", (AutoSchedulerResults(Pipeline::*)(const Target &, const AutoschedulerParams &) const) & Pipeline::apply_autoscheduler, + .def("apply_autoscheduler", &Pipeline::apply_autoscheduler, py::arg("target"), py::arg("autoscheduler_params")) .def("get_func", &Pipeline::get_func, py::arg("index")) @@ -84,106 +84,119 @@ void define_pipeline(py::module &m) { }, py::arg("outputs"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_bitcode", [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - p.compile_to_bitcode(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_bitcode", [](Pipeline &p, const std::string &filename, const std::vector &args, const Target &target) { - p.compile_to_bitcode(filename, args, "", to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) - .def( - "compile_to_llvm_assembly", [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - p.compile_to_llvm_assembly(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_llvm_assembly", [](Pipeline &p, const std::string &filename, const std::vector &args, const Target &target) { - p.compile_to_llvm_assembly(filename, args, "", to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) - .def( - "compile_to_object", [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - p.compile_to_object(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_object", [](Pipeline &p, const std::string &filename, const std::vector &args, const Target &target) { - p.compile_to_object(filename, args, "", to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) - .def( - "compile_to_header", [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - p.compile_to_header(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) - .def( - "compile_to_assembly", [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - p.compile_to_assembly(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) - .def( - "compile_to_assembly", [](Pipeline &p, const std::string &filename, const std::vector &args, const Target &target) { - p.compile_to_assembly(filename, args, "", to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) - .def( - "compile_to_c", [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { - p.compile_to_c(filename, args, fn_name, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) - .def( - "compile_to_lowered_stmt", [](Pipeline &p, const std::string &filename, const std::vector &args, StmtOutputFormat fmt, const Target &target) { - p.compile_to_lowered_stmt(filename, args, fmt, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fmt") = Text, py::arg("target") = Target()) - .def( - "compile_to_conceptual_stmt", [](Pipeline &p, const std::string &filename, const std::vector &args, StmtOutputFormat fmt, const Target &target) { - p.compile_to_conceptual_stmt(filename, args, fmt, to_aot_target(target)); - }, - py::arg("filename"), py::arg("arguments"), py::arg("fmt") = Text, py::arg("target") = Target()) - .def( - "compile_to_file", [](Pipeline &p, const std::string &filename_prefix, const std::vector &args, const std::string &fn_name, const Target &target) { - p.compile_to_file(filename_prefix, args, fn_name, to_aot_target(target)); - }, - py::arg("filename_prefix"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) - .def( - "compile_to_static_library", [](Pipeline &p, const std::string &filename_prefix, const std::vector &args, const std::string &fn_name, const Target &target) { - p.compile_to_static_library(filename_prefix, args, fn_name, to_aot_target(target)); - }, - py::arg("filename_prefix"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) + .def("compile_to_bitcode", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + p.compile_to_bitcode(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) + + .def("compile_to_bitcode", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const Target &target) { + p.compile_to_bitcode(filename, args, "", to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) + + .def("compile_to_llvm_assembly", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + p.compile_to_llvm_assembly(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) + + .def("compile_to_llvm_assembly", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const Target &target) { + p.compile_to_llvm_assembly(filename, args, "", to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) + + .def("compile_to_object", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + p.compile_to_object(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) + + .def("compile_to_object", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const Target &target) { + p.compile_to_object(filename, args, "", to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) + + .def("compile_to_header", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + p.compile_to_header(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) + + .def("compile_to_assembly", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + p.compile_to_assembly(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target()) + + .def("compile_to_assembly", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const Target &target) { + p.compile_to_assembly(filename, args, "", to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("target") = Target()) + + .def("compile_to_c", // + [](Pipeline &p, const std::string &filename, const std::vector &args, const std::string &fn_name, const Target &target) { + p.compile_to_c(filename, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) + + .def("compile_to_lowered_stmt", // + [](Pipeline &p, const std::string &filename, const std::vector &args, StmtOutputFormat fmt, const Target &target) { + p.compile_to_lowered_stmt(filename, args, fmt, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fmt") = Text, py::arg("target") = Target()) + + .def("compile_to_conceptual_stmt", // + [](Pipeline &p, const std::string &filename, const std::vector &args, StmtOutputFormat fmt, const Target &target) { + p.compile_to_conceptual_stmt(filename, args, fmt, to_aot_target(target)); // + }, + py::arg("filename"), py::arg("arguments"), py::arg("fmt") = Text, py::arg("target") = Target()) + + .def("compile_to_file", // + [](Pipeline &p, const std::string &filename_prefix, const std::vector &args, const std::string &fn_name, const Target &target) { + p.compile_to_file(filename_prefix, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename_prefix"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) + + .def("compile_to_static_library", // + [](Pipeline &p, const std::string &filename_prefix, const std::vector &args, const std::string &fn_name, const Target &target) { + p.compile_to_static_library(filename_prefix, args, fn_name, to_aot_target(target)); // + }, + py::arg("filename_prefix"), py::arg("arguments"), py::arg("fn_name") = "", py::arg("target") = Target()) .def("compile_to_multitarget_static_library", &Pipeline::compile_to_multitarget_static_library, py::arg("filename_prefix"), py::arg("arguments"), py::arg("targets")) .def("compile_to_multitarget_object_files", &Pipeline::compile_to_multitarget_object_files, py::arg("filename_prefix"), py::arg("arguments"), py::arg("targets"), py::arg("suffixes")) - .def( - "compile_to_module", [](Pipeline &p, const std::vector &args, const std::string &fn_name, const Target &target, LinkageType linkage_type) -> Module { - return p.compile_to_module(args, fn_name, to_aot_target(target), linkage_type); - }, - py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target(), py::arg("linkage") = LinkageType::ExternalPlusMetadata) - - .def( - "compile_jit", [](Pipeline &p, const Target &target) { - p.compile_jit(to_jit_target(target)); - }, - py::arg("target") = Target()) - - .def( - "compile_to_callable", [](Pipeline &p, const std::vector &args, const Target &target) { - return p.compile_to_callable(args, to_jit_target(target)); - }, - py::arg("arguments"), py::arg("target") = Target()) - - .def( - "realize", [](Pipeline &p, Buffer<> buffer, const Target &target) -> void { - py::gil_scoped_release release; - - PyJITUserContext juc; - p.realize(&juc, Realization(std::move(buffer)), target); - }, - py::arg("dst"), py::arg("target") = Target()) + .def("compile_to_module", // + [](Pipeline &p, const std::vector &args, const std::string &fn_name, const Target &target, LinkageType linkage_type) -> Module { + return p.compile_to_module(args, fn_name, to_aot_target(target), linkage_type); // + }, + py::arg("arguments"), py::arg("fn_name"), py::arg("target") = Target(), py::arg("linkage") = LinkageType::ExternalPlusMetadata) + + .def("compile_jit", // + [](Pipeline &p, const Target &target) { + p.compile_jit(to_jit_target(target)); // + }, + py::arg("target") = Target()) + + .def("compile_to_callable", // + [](Pipeline &p, const std::vector &args, const Target &target) { + return p.compile_to_callable(args, to_jit_target(target)); // + }, + py::arg("arguments"), py::arg("target") = Target()) + + .def("realize", // + [](Pipeline &p, Buffer<> buffer, const Target &target) -> void { + py::gil_scoped_release release; + + PyJITUserContext juc; + p.realize(&juc, Realization(std::move(buffer)), target); // + }, + py::arg("dst"), py::arg("target") = Target()) // It's important to have this overload of realize() go first: // passing an empty list [] is ambiguous in Python, and could match to @@ -191,76 +204,76 @@ void define_pipeline(py::module &m) { // (it allows realizing a 0-dimensional/scalar buffer) and the former is // not (it will always assert-fail). Putting this one first allows it to // be the first one chosen by the bindings in this case. - .def( - "realize", [](Pipeline &p, std::vector sizes, const Target &target) -> py::object { - std::optional r; - { - py::gil_scoped_release release; - - PyJITUserContext juc; - r = p.realize(&juc, std::move(sizes), target); - } - return realization_to_object(*r); - }, - py::arg("sizes") = std::vector{}, py::arg("target") = Target()) + .def("realize", // + [](Pipeline &p, std::vector sizes, const Target &target) -> py::object { + std::optional r; + { + py::gil_scoped_release release; + + PyJITUserContext juc; + r = p.realize(&juc, std::move(sizes), target); + } + return realization_to_object(*r); // + }, + py::arg("sizes") = std::vector{}, py::arg("target") = Target()) // This will actually allow a list-of-buffers as well as a tuple-of-buffers, but that's OK. - .def( - "realize", [](Pipeline &p, std::vector> buffers, const Target &t) -> void { - py::gil_scoped_release release; - - PyJITUserContext juc; - p.realize(&juc, Realization(std::move(buffers)), t); - }, - py::arg("dst"), py::arg("target") = Target()) - - .def( - "infer_input_bounds", [](Pipeline &p, const py::object &dst, const Target &target) -> void { - const Target t = to_jit_target(target); - PyJITUserContext juc; - - // dst could be Buffer<>, vector, or vector - try { - Buffer<> b = dst.cast>(); - p.infer_input_bounds(&juc, b, t); - return; - } catch (...) { - // fall thru - } - - try { - std::vector> v = dst.cast>>(); - p.infer_input_bounds(&juc, Realization(std::move(v)), t); - return; - } catch (...) { - // fall thru - } - - try { - std::vector v = dst.cast>(); - p.infer_input_bounds(&juc, v, t); - return; - } catch (...) { - // fall thru - } - - throw py::value_error("Invalid arguments to infer_input_bounds"); - }, - py::arg("dst"), py::arg("target") = Target()) + .def("realize", // + [](Pipeline &p, std::vector> buffers, const Target &t) -> void { + py::gil_scoped_release release; + + PyJITUserContext juc; + p.realize(&juc, Realization(std::move(buffers)), t); // + }, + py::arg("dst"), py::arg("target") = Target()) + + .def("infer_input_bounds", // + [](Pipeline &p, const py::object &dst, const Target &target) -> void { + const Target t = to_jit_target(target); + PyJITUserContext juc; + + // dst could be Buffer<>, vector, or vector + try { + Buffer<> b = dst.cast>(); + p.infer_input_bounds(&juc, b, t); + return; + } catch (...) { + // fall thru + } + + try { + std::vector> v = dst.cast>>(); + p.infer_input_bounds(&juc, Realization(std::move(v)), t); + return; + } catch (...) { + // fall thru + } + + try { + std::vector v = dst.cast>(); + p.infer_input_bounds(&juc, v, t); + return; + } catch (...) { + // fall thru + } + + throw py::value_error("Invalid arguments to infer_input_bounds"); // + }, + py::arg("dst"), py::arg("target") = Target()) .def("infer_arguments", [](Pipeline &p) -> std::vector { - return p.infer_arguments(); + return p.infer_arguments(); // }) .def("defined", &Pipeline::defined) .def("invalidate_cache", &Pipeline::invalidate_cache) - .def( - "add_requirement", [](Pipeline &p, const Expr &condition, const py::args &error_args) -> void { - auto v = collect_print_args(error_args); - p.add_requirement(condition, v); - }, - py::arg("condition")) + .def("add_requirement", // + [](Pipeline &p, const Expr &condition, const py::args &error_args) -> void { + auto v = collect_print_args(error_args); + p.add_requirement(condition, v); // + }, + py::arg("condition")) .def("__repr__", [](const Pipeline &p) -> std::string { std::ostringstream o; @@ -271,7 +284,7 @@ void define_pipeline(py::module &m) { comma = ","; } o << "]>"; - return o.str(); + return o.str(); // }); // TODO: These should really live in PyGenerator.cpp once that lands diff --git a/python_bindings/src/halide/halide_/PyScheduleMethods.h b/python_bindings/src/halide/halide_/PyScheduleMethods.h index 5e11f757da72..75475dc73f9d 100644 --- a/python_bindings/src/halide/halide_/PyScheduleMethods.h +++ b/python_bindings/src/halide/halide_/PyScheduleMethods.h @@ -102,12 +102,16 @@ HALIDE_NEVER_INLINE void add_schedule_methods(PythonClass &class_instance) { .def("hexagon", &T::hexagon, py::arg("x") = Var::outermost()) .def("prefetch", (T & (T::*)(const Func &, const VarOrRVar &, const VarOrRVar &, Expr, PrefetchBoundStrategy)) & T::prefetch, py::arg("func"), py::arg("at"), py::arg("from"), py::arg("offset") = 1, py::arg("strategy") = PrefetchBoundStrategy::GuardWithIf) - .def( - "prefetch", [](T &t, const ImageParam &image, const VarOrRVar &at, const VarOrRVar &from, const Expr &offset, PrefetchBoundStrategy strategy) -> T & { - // Templated function; specializing only on ImageParam for now - return t.template prefetch(image, at, from, offset, strategy); - }, - py::arg("image"), py::arg("at"), py::arg("from"), py::arg("offset") = 1, py::arg("strategy") = PrefetchBoundStrategy::GuardWithIf); + .def("prefetch", // + [](T &t, const ImageParam &image, const VarOrRVar &at, const VarOrRVar &from, const Expr &offset, PrefetchBoundStrategy strategy) -> T & { + // Templated function; specializing only on ImageParam for now + return t.template prefetch(image, at, from, offset, strategy); // + }, + py::arg("image"), // + py::arg("at"), // + py::arg("from"), // + py::arg("offset") = 1, // + py::arg("strategy") = PrefetchBoundStrategy::GuardWithIf); } } // namespace PythonBindings diff --git a/python_bindings/src/halide/halide_/PyStage.cpp b/python_bindings/src/halide/halide_/PyStage.cpp index fac47fa3cf1f..cb29e3cfb66d 100644 --- a/python_bindings/src/halide/halide_/PyStage.cpp +++ b/python_bindings/src/halide/halide_/PyStage.cpp @@ -14,9 +14,9 @@ void define_stage(py::module &m) { .def("dump_argument_list", &Stage::dump_argument_list) .def("name", &Stage::name) - .def("rfactor", (Func(Stage::*)(const std::vector> &)) & Stage::rfactor, + .def("rfactor", (Func(Stage::*)(const std::vector> &))&Stage::rfactor, py::arg("preserved")) - .def("rfactor", (Func(Stage::*)(const RVar &, const Var &)) & Stage::rfactor, + .def("rfactor", (Func(Stage::*)(const RVar &, const Var &))&Stage::rfactor, py::arg("r"), py::arg("v")) .def("unscheduled", &Stage::unscheduled); diff --git a/python_bindings/src/halide/halide_/PyVar.cpp b/python_bindings/src/halide/halide_/PyVar.cpp index 88d08a64e500..8757684c5229 100644 --- a/python_bindings/src/halide/halide_/PyVar.cpp +++ b/python_bindings/src/halide/halide_/PyVar.cpp @@ -25,7 +25,7 @@ void define_var(py::module &m) { .def("is_implicit", (bool(Var::*)() const) & Var::is_implicit) .def("implicit_index", (int(Var::*)() const) & Var::implicit_index) .def("is_placeholder", (bool(Var::*)() const) & Var::is_placeholder) - .def_static("implicit", (Var(*)(int)) & Var::implicit) + .def_static("implicit", (Var(*)(int))&Var::implicit) .def_static("outermost", &Var::outermost) .def("__repr__", &var_repr) .def("__str__", &Var::name); diff --git a/run-clang-format.sh b/run-clang-format.sh index 33a0ac6d7152..4c758e6ca07a 100755 --- a/run-clang-format.sh +++ b/run-clang-format.sh @@ -2,35 +2,74 @@ set -e -ROOT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}" )" >/dev/null 2>&1 && pwd )" +ROOT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd)" -# We are currently standardized on using LLVM/Clang17 for this script. +## + +EXPECTED_VERSION=19 + +## + +usage() { echo -e "Usage: $0 [-c]" 1>&2; exit 1; } + +# Fix the formatting in-place +MODE_FLAGS=(-i --sort-includes) + +while getopts "c" o; do + case "${o}" in + c) + # Only check the files and print formatting errors + MODE_FLAGS=(--dry-run -Werror) + ;; + *) + usage + ;; + esac +done +shift $((OPTIND - 1)) + +if [[ "${MODE_FLAGS[*]}" =~ "-i" ]]; then + if ! git diff-files --quiet --ignore-submodules; then + echo -e "\033[0;31m" # RED + echo "WARNING: There are still uncommited changes in your working tree." + echo " Reverting this formatting action will be difficult." + echo -e "\033[0m" # RESET + git diff-files --ignore-submodules + echo + read -p "Do you wish to continue (Y/N)? " -r + if ! [[ $REPLY =~ ^[Yy]$ ]]; then + exit + fi + fi +fi + +# We are currently standardized on using LLVM/Clang19 for this script. # Note that this is totally independent of the version of LLVM that you -# are using to build Halide itself. If you don't have LLVM17 installed, +# are using to build Halide itself. If you don't have LLVM19 installed, # you can usually install what you need easily via: # -# sudo apt-get install llvm-17 clang-17 libclang-17-dev clang-tidy-17 -# export CLANG_FORMAT_LLVM_INSTALL_DIR=/usr/lib/llvm-17 +# sudo apt-get install llvm-19 clang-19 libclang-19-dev clang-tidy-19 +# export CLANG_FORMAT_LLVM_INSTALL_DIR=/usr/lib/llvm-19 # # On macOS: # -# brew install llvm@17 -# export CLANG_FORMAT_LLVM_INSTALL_DIR=/opt/homebrew/opt/llvm@17 +# brew install llvm@19 +# export CLANG_FORMAT_LLVM_INSTALL_DIR=/opt/homebrew/opt/llvm@19 if [ -z "$CLANG_FORMAT_LLVM_INSTALL_DIR" ]; then - echo "CLANG_FORMAT_LLVM_INSTALL_DIR must point to an LLVM installation dir for this script." - exit 1 + echo "CLANG_FORMAT_LLVM_INSTALL_DIR must point to an LLVM installation dir for this script." + exit 1 fi -echo "CLANG_FORMAT_LLVM_INSTALL_DIR = ${CLANG_FORMAT_LLVM_INSTALL_DIR}" +echo "CLANG_FORMAT_LLVM_INSTALL_DIR=${CLANG_FORMAT_LLVM_INSTALL_DIR}" CLANG_FORMAT="${CLANG_FORMAT_LLVM_INSTALL_DIR}/bin/clang-format" VERSION=$("${CLANG_FORMAT}" --version) -if [[ ${VERSION} =~ .*version\ 17.* ]]; then - echo "clang-format version 17 found." +if [[ ${VERSION} =~ .*version\ $EXPECTED_VERSION.* ]]; then + echo "clang-format version $EXPECTED_VERSION found." else - echo "CLANG_FORMAT_LLVM_INSTALL_DIR must point to an LLVM 17 install!" + echo "CLANG_FORMAT_LLVM_INSTALL_DIR must point to an LLVM $EXPECTED_VERSION install!" exit 1 fi @@ -44,5 +83,6 @@ find "${ROOT_DIR}/apps" \ "${ROOT_DIR}/python_bindings" \ -not -path "${ROOT_DIR}/src/runtime/hexagon_remote/bin/src/*" \ \( -name "*.cpp" -o -name "*.h" -o -name "*.c" \) -and -not -wholename "*/.*" \ - -print0 | \ - xargs -0 "${CLANG_FORMAT}" -i -style=file \ No newline at end of file + -print0 | xargs -0 "${CLANG_FORMAT}" "${MODE_FLAGS[@]}" -style=file + +exit "${PIPESTATUS[1]}" diff --git a/src/IRMatch.h b/src/IRMatch.h index 676bd037bd27..1e24388a462c 100644 --- a/src/IRMatch.h +++ b/src/IRMatch.h @@ -648,14 +648,14 @@ struct BinOp { } const Op &op = (const Op &)e; return (a.template match(*op.a.get(), state) && - b.template match::mask>(*op.b.get(), state)); + b.template match<(bound | bindings::mask)>(*op.b.get(), state)); } template HALIDE_ALWAYS_INLINE bool match(const BinOp &op, MatcherState &state) const noexcept { return (std::is_same::value && a.template match(unwrap(op.a), state) && - b.template match::mask>(unwrap(op.b), state)); + b.template match<(bound | bindings::mask)>(unwrap(op.b), state)); } constexpr static bool foldable = A::foldable && B::foldable; @@ -750,14 +750,14 @@ struct CmpOp { } const Op &op = (const Op &)e; return (a.template match(*op.a.get(), state) && - b.template match::mask>(*op.b.get(), state)); + b.template match<(bound | bindings::mask)>(*op.b.get(), state)); } template HALIDE_ALWAYS_INLINE bool match(const CmpOp &op, MatcherState &state) const noexcept { return (std::is_same::value && a.template match(unwrap(op.a), state) && - b.template match::mask>(unwrap(op.b), state)); + b.template match<(bound | bindings::mask)>(unwrap(op.b), state)); } constexpr static bool foldable = A::foldable && B::foldable; @@ -1357,7 +1357,7 @@ struct Intrin { HALIDE_ALWAYS_INLINE bool match_args(int, const Call &c, MatcherState &state) const noexcept { using T = decltype(std::get(args)); return (std::get(args).template match(*c.args[i].get(), state) && - match_args::mask>(0, c, state)); + match_args::mask)>(0, c, state)); } template @@ -1687,14 +1687,14 @@ struct SelectOp { } const Select &op = (const Select &)e; return (c.template match(*op.condition.get(), state) && - t.template match::mask>(*op.true_value.get(), state) && - f.template match::mask | bindings::mask>(*op.false_value.get(), state)); + t.template match<(bound | bindings::mask)>(*op.true_value.get(), state) && + f.template match<(bound | bindings::mask | bindings::mask)>(*op.false_value.get(), state)); } template HALIDE_ALWAYS_INLINE bool match(const SelectOp &instance, MatcherState &state) const noexcept { return (c.template match(unwrap(instance.c), state) && - t.template match::mask>(unwrap(instance.t), state) && - f.template match::mask | bindings::mask>(unwrap(instance.f), state)); + t.template match<(bound | bindings::mask)>(unwrap(instance.t), state) && + f.template match<(bound | bindings::mask | bindings::mask)>(unwrap(instance.f), state)); } HALIDE_ALWAYS_INLINE @@ -1760,7 +1760,7 @@ struct BroadcastOp { template HALIDE_ALWAYS_INLINE bool match(const BroadcastOp &op, MatcherState &state) const noexcept { return (a.template match(unwrap(op.a), state) && - lanes.template match::mask>(unwrap(op.lanes), state)); + lanes.template match<(bound | bindings::mask)>(unwrap(op.lanes), state)); } HALIDE_ALWAYS_INLINE @@ -1824,8 +1824,8 @@ struct RampOp { } const Ramp &op = (const Ramp &)e; if (a.template match(*op.base.get(), state) && - b.template match::mask>(*op.stride.get(), state) && - lanes.template match::mask | bindings::mask>(op.lanes, state)) { + b.template match<(bound | bindings::mask)>(*op.stride.get(), state) && + lanes.template match<(bound | bindings::mask | bindings::mask)>(op.lanes, state)) { return true; } else { return false; @@ -1835,8 +1835,8 @@ struct RampOp { template HALIDE_ALWAYS_INLINE bool match(const RampOp &op, MatcherState &state) const noexcept { return (a.template match(unwrap(op.a), state) && - b.template match::mask>(unwrap(op.b), state) && - lanes.template match::mask | bindings::mask>(unwrap(op.lanes), state)); + b.template match<(bound | bindings::mask)>(unwrap(op.b), state) && + lanes.template match<(bound | bindings::mask | bindings::mask)>(unwrap(op.lanes), state)); } HALIDE_ALWAYS_INLINE @@ -1887,7 +1887,7 @@ struct VectorReduceOp { const VectorReduce &op = (const VectorReduce &)e; if (op.op == reduce_op && a.template match(*op.value.get(), state) && - lanes.template match::mask>(op.type.lanes(), state)) { + lanes.template match<(bound | bindings::mask)>(op.type.lanes(), state)) { return true; } } @@ -1898,7 +1898,7 @@ struct VectorReduceOp { HALIDE_ALWAYS_INLINE bool match(const VectorReduceOp &op, MatcherState &state) const noexcept { return (reduce_op == reduce_op_2 && a.template match(unwrap(op.a), state) && - lanes.template match::mask>(unwrap(op.lanes), state)); + lanes.template match<(bound | bindings::mask)>(unwrap(op.lanes), state)); } HALIDE_ALWAYS_INLINE @@ -2147,9 +2147,9 @@ struct SliceOp { return v.vectors.size() == 1 && v.is_slice() && vec.template match(*v.vectors[0].get(), state) && - base.template match::mask>(v.slice_begin(), state) && - stride.template match::mask | bindings::mask>(v.slice_stride(), state) && - lanes.template match::mask | bindings::mask | bindings::mask>(v.type.lanes(), state); + base.template match<(bound | bindings::mask)>(v.slice_begin(), state) && + stride.template match<(bound | bindings::mask | bindings::mask)>(v.slice_stride(), state) && + lanes.template match<(bound | bindings::mask | bindings::mask | bindings::mask)>(v.type.lanes(), state); } HALIDE_ALWAYS_INLINE diff --git a/src/Solve.cpp b/src/Solve.cpp index 20d6f5200101..7a206d74f8f3 100644 --- a/src/Solve.cpp +++ b/src/Solve.cpp @@ -1346,8 +1346,8 @@ void solve_test() { continue; } for (int num = 5; num <= 10; num++) { - Expr in[] = {x * den= num, x * den> num, - x / den= num, x / den> num}; + Expr in[] = {x * den < num, x * den <= num, x * den == num, x * den != num, x * den >= num, x * den > num, + x / den < num, x / den <= num, x / den == num, x / den != num, x / den >= num, x / den > num}; for (const auto &e : in) { SolverResult solved = solve_expression(e, "x"); internal_assert(solved.fully_solved) << "Error: failed to solve for x in " << e << "\n"; diff --git a/src/runtime/HalideBuffer.h b/src/runtime/HalideBuffer.h index 9741c0278e3d..8f25bf45e95e 100644 --- a/src/runtime/HalideBuffer.h +++ b/src/runtime/HalideBuffer.h @@ -2061,9 +2061,7 @@ class Buffer { } template - HALIDE_ALWAYS_INLINE - storage_T * - address_of(Args... args) const { + HALIDE_ALWAYS_INLINE storage_T *address_of(Args... args) const { if (T_is_void) { return (storage_T *)(this->buf.host) + offset_of(0, args...) * type().bytes(); } else { @@ -2139,9 +2137,7 @@ class Buffer { template::value>::type> - HALIDE_ALWAYS_INLINE - not_void_T & - operator()(int first, Args... rest) { + HALIDE_ALWAYS_INLINE not_void_T &operator()(int first, Args... rest) { static_assert(!T_is_void, "Cannot use operator() on Buffer types"); constexpr int expected_dims = 1 + (int)(sizeof...(rest)); diff --git a/src/runtime/HalidePyTorchCudaHelpers.h b/src/runtime/HalidePyTorchCudaHelpers.h index f4a865ac47f5..edfea67e419b 100644 --- a/src/runtime/HalidePyTorchCudaHelpers.h +++ b/src/runtime/HalidePyTorchCudaHelpers.h @@ -17,7 +17,7 @@ namespace PyTorch { typedef struct UserContext { UserContext(int id, CUcontext *ctx, cudaStream_t *stream) - : device_id(id), cuda_context(ctx), stream(stream){}; + : device_id(id), cuda_context(ctx), stream(stream) {}; int device_id; CUcontext *cuda_context; diff --git a/src/runtime/hexagon_host.cpp b/src/runtime/hexagon_host.cpp index 80dcca1dc707..6389b73cedbb 100644 --- a/src/runtime/hexagon_host.cpp +++ b/src/runtime/hexagon_host.cpp @@ -105,7 +105,7 @@ ALWAYS_INLINE T *uint64_to_ptr(const uint64_t &u) { template ALWAYS_INLINE uint64_t ptr_to_uint64(T *ptr) { - return (uint64_t) reinterpret_cast(ptr); + return (uint64_t)reinterpret_cast(ptr); } template diff --git a/test/correctness/shared_self_references.cpp b/test/correctness/shared_self_references.cpp index a3e28eb1a450..4f08da7a86d5 100644 --- a/test/correctness/shared_self_references.cpp +++ b/test/correctness/shared_self_references.cpp @@ -14,7 +14,7 @@ int main(int argc, char **argv) { f(0) = e; f(1) = e; } // Destroy e - } // Destroy f + } // Destroy f // f should have been cleaned up. valgrind will complain if it // hasn't been. diff --git a/test/correctness/simd_op_check_sve2.cpp b/test/correctness/simd_op_check_sve2.cpp index d0e614d657a2..cefa4d14aaab 100644 --- a/test/correctness/simd_op_check_sve2.cpp +++ b/test/correctness/simd_op_check_sve2.cpp @@ -1164,7 +1164,7 @@ class SimdOpCheckArmSve : public SimdOpCheckTest { int default_vec_factor, bool is_enabled = true /* false to skip testing */) : parent(p), default_bits(default_bits), default_instr_lanes(default_instr_lanes), - default_vec_factor(default_vec_factor), is_enabled(is_enabled){}; + default_vec_factor(default_vec_factor), is_enabled(is_enabled) {}; AddTestFunctor(SimdOpCheckArmSve &p, int default_bits, @@ -1173,7 +1173,7 @@ class SimdOpCheckArmSve : public SimdOpCheckTest { bool is_enabled = true /* false to skip testing */) : parent(p), default_bits(default_bits), default_instr_lanes(Instruction::get_instr_lanes(default_bits, default_vec_factor, p.target)), - default_vec_factor(default_vec_factor), is_enabled(is_enabled){}; + default_vec_factor(default_vec_factor), is_enabled(is_enabled) {}; // Constructs single Instruction with default parameters void operator()(const string &opcode, Expr e) { diff --git a/test/generator/metadata_tester_generator.cpp b/test/generator/metadata_tester_generator.cpp index 3f200d4989dc..65774bceb703 100644 --- a/test/generator/metadata_tester_generator.cpp +++ b/test/generator/metadata_tester_generator.cpp @@ -39,10 +39,10 @@ class MetadataTester : public Halide::Generator { Input array2_i32{"array2_i32", 32, -32, 127}; Input array_h{"array_h", nullptr}; // must be overridden to size=2 - Input[2]> buffer_array_input1 { "buffer_array_input1" }; - Input[2]> buffer_array_input2 { "buffer_array_input2" }; // buffer_array_input2.dim must be set - Input[2]> buffer_array_input3 { "buffer_array_input3" }; // buffer_array_input2.type must be set - Input[2]> buffer_array_input4 { "buffer_array_input4" }; // dim and type must be set + Input[2]> buffer_array_input1{"buffer_array_input1"}; + Input[2]> buffer_array_input2{"buffer_array_input2"}; // buffer_array_input2.dim must be set + Input[2]> buffer_array_input3{"buffer_array_input3"}; // buffer_array_input2.type must be set + Input[2]> buffer_array_input4{"buffer_array_input4"}; // dim and type must be set // .size must be specified for all of these Input[]> buffer_array_input5{"buffer_array_input5"}; Input[]> buffer_array_input6{"buffer_array_input6"}; // buffer_array_input2.dim must be set @@ -65,9 +65,9 @@ class MetadataTester : public Halide::Generator { Output array_outputs2{"array_outputs2", {Float(32), Float(32)}, 3}; Output array_outputs3{"array_outputs3"}; - Output[2]> array_outputs4 { "array_outputs4" }; - Output[2]> array_outputs5 { "array_outputs5" }; // dimensions will be inferred by usage - Output[2]> array_outputs6 { "array_outputs6" }; // dimensions and type will be inferred by usage + Output[2]> array_outputs4{"array_outputs4"}; + Output[2]> array_outputs5{"array_outputs5"}; // dimensions will be inferred by usage + Output[2]> array_outputs6{"array_outputs6"}; // dimensions and type will be inferred by usage // .size must be specified for all of these Output[]> array_outputs7{"array_outputs7"}; diff --git a/test/generator/stubtest_generator.cpp b/test/generator/stubtest_generator.cpp index 8f5b41640e41..de0a03f1c797 100644 --- a/test/generator/stubtest_generator.cpp +++ b/test/generator/stubtest_generator.cpp @@ -32,7 +32,7 @@ class StubTest : public Halide::Generator { Input> typed_buffer_input{"typed_buffer_input"}; Input> untyped_buffer_input{"untyped_buffer_input"}; - Input[2]> array_buffer_input { "array_buffer_input" }; + Input[2]> array_buffer_input{"array_buffer_input"}; Input simple_input{"simple_input", 3}; // require a 3-dimensional Func but leave Type unspecified Input array_input{"array_input", 3}; // require a 3-dimensional Func but leave Type and ArraySize unspecified // Note that Input does not (yet) support Tuples @@ -46,7 +46,7 @@ class StubTest : public Halide::Generator { Output> untyped_buffer_output{"untyped_buffer_output"}; Output> tupled_output{"tupled_output", {Float(32), Int(32)}}; Output> static_compiled_buffer_output{"static_compiled_buffer_output"}; - Output[2]> array_buffer_output { "array_buffer_output" }; + Output[2]> array_buffer_output{"array_buffer_output"}; Output> float16_output{"float16_output"}; Output> bfloat16_output{"bfloat16_output"}; diff --git a/test/performance/rfactor.cpp b/test/performance/rfactor.cpp index a1121a3f0338..2a1296059dc7 100644 --- a/test/performance/rfactor.cpp +++ b/test/performance/rfactor.cpp @@ -269,11 +269,11 @@ int dot_product() { // Reference implementation Func dot_ref("dot_ref"); dot_ref() = 0.0f; - dot_ref() += (A(r.x)) * B(r.x); + dot_ref() += A(r.x) * B(r.x); Func dot("dot"); dot() = 0.0f; - dot() += (A(r.x)) * B(r.x); + dot() += A(r.x) * B(r.x); RVar rxo, rxi, rxio, rxii; dot.update().split(r.x, rxo, rxi, 4 * 8192);