What's Changed
- Merged PR 2391: Update quickstart example, updated docs structure per
feedback. [Lisa Ong]
* Teasers for transformations in the Quickstart sample (to differentiate Accera from others), with benchmarking
* Removed the Miscellaneous section, redistributed various docs to various related locations
* Renamed the cross compilation tutorial so that it is ordered last
- Merged PR 2392: Populate Target.Models based on known devices. [Kern
Handa]
Populate Target.Models based on known devices
- Merged PR 2390: Merge multiple HAT files during project building.
[Kern Handa]
Merge multiple HAT files during project building
Related work items: 3559
- Merged PR 2386: Add support for various targets. [Kern Handa]
Add support for various targets
Related work items: 3631
- Merged PR 2389: [nfc] Doc typos and consistency fixes. [Lisa Ong]
- Merged PR 2388: Update quickstart example, add binder quickstart.
[Lisa Ong]
* Update quickstart example to perform a matmul + ReLU (unoptimized)
* Add Launch in Binder button to run everything in the browser
- Merged PR 2387: Placeholder GPU GridUnit definitions, add library
creation from multiple object files. [Lisa Ong]
Dependent HAT PR: https://github.com/microsoft/hat/pull/21
* GridUnit definitions are static until we have real GPU targets. These are updated just to be consistent with the Manual
* When not cross compiling, combine multiple .obj/.o into .lib/.a
Related work items: 3576
- Merged PR 2384: Update target docs, split Intel generation 8 and 9 for
consistency. [Lisa Ong]
* Update target docs to list the name of the target in the table
* Define separate models for Intel generation 8 and 9 for consistency
Related work items: 3631
- Merged PR 2383: Support dynamic libs from `Package.build` [Lisa Ong]
* Add static and dynamic variants to the HAT and MLIR formats
* MLIR format is also split because we'd want to support MLIR inspection of the cross-compilation scenario without forcing users to switch between dynamic and static
* Updated README sample
Left for future work:
* Combining multiple object files into a static lib or dynamic lib. We'd need to think about how HAT packages can be merged together (for example, how to reconcile the metadata in the HAT file, such as description, author - do we merge all metadata or just pick the first HAT file encountered as the "master", etc)
Related PR: https://github.com/microsoft/hat/pull/18
Related work items: 3576
- Merged PR 2382: [nfc] Move Case Studies out of the Accera repo. [Lisa
Ong]
Case Studies will live in other repositories, and be cross linked from the Accera repo's Case Studies README.md (to be added in the future).
Related work items: 3632
- Merged PR 2379: Specify dynamic lib dependencies from the HAT Package.
[Lisa Ong]
This is the final missing piece before we transition to building static / dynamic libs using `hatlib`.
* `Plan` infers additional dynamic dependencies when the target is GPU or when parallelization is requested.
* `Package.add` collects the dependency info the various `Plan` instances.
* `Package.build`, the platform parameter is used to resolve to the appropriate library (either a path or a -l directive).
* For library paths that cannot be fully determined in advance, we default to the current working directory, so perhaps the user can put the lib in the same path as the binaries. (this needs to be fleshed out more)
* Removed dead code
Dependent hatlib PR: https://github.com/microsoft/hat/pull/16/files
Related work items: 3576
- Merged PR 2380: Add Raspberry Pi 4 (B) support. [Kern Handa]
Related work items: 3631
- Merged PR 2368: Update and optimize acc-translate. [Abdul Dakkak]
- propagate constants while generating C++ code
- inline mlir within the C++ code to ease debugging
- increase support for vector ops
- silence a lot of warnings that were being emitted in the acc-translate codebase
The following
// CONFIG: {"K":2048,"M":2048,"N":2048,"block":{"x":16,"y":16,"z":1},"grid":{"x":128,"y":128,"z":1}}
module gemm_naive_14479263422999410716_module attributes {gpu.binary = "HSACO"} {
func gemm_naive_14479263422999410716(%arg0: memref<2048x2048xf32> loc(unknown), %arg1: memref<2048x2048xf32> loc(unknown), %arg2: memref<2048x2048xf32> loc(unknown)) {
%c16 = constant 16 : index loc(unknown)
%c0 = constant 0 : index loc(unknown)
%c2048 = constant 2048 : index loc(unknown)
%c1 = constant 1 : index loc(unknown)
%cst = constant 0.000000e+00 : f32 loc(unknown)
%0 = "gpu.thread_id"() {dimension = "x"} : () -> index loc(unknown)
%1 = "gpu.thread_id"() {dimension = "y"} : () -> index loc(unknown)
%2 = "gpu.block_id"() {dimension = "x"} : () -> index loc(unknown)
%3 = "gpu.block_id"() {dimension = "y"} : () -> index loc(unknown)
%4 = scf.for %arg3 = %c0 to %c2048 step %c1 iter_args(%arg4 = %cst) -> (f32) {
%11 = muli %3, %c16 : index loc(unknown)
%12 = addi %1, %11 : index loc(unknown)
%13 = memref.load %arg0[%12, %arg3] : memref<2048x2048xf32> loc(unknown)
%14 = muli %2, %c16 : index loc(unknown)
%15 = addi %0, %14 : index loc(unknown)
%16 = memref.load %arg1[%arg3, %15] : memref<2048x2048xf32> loc(unknown)
%17 = mulf %13, %16 {RelaxedPrecision} : f32 loc(unknown)
%18 = addf %arg4, %17 {RelaxedPrecision} : f32 loc(unknown)
scf.yield %18 : f32 loc(unknown)
} loc(unknown)
%5 = muli %3, %c16 : index loc(unknown)
%6 = addi %1, %5 : index loc(unknown)
%7 = muli %2, %c16 : index loc(unknown)
%8 = addi %0, %7 : index loc(unknown)
%9 = memref.load %arg2[%6, %8] : memref<2048x2048xf32> loc(unknown)
%10 = addf %9, %4 {RelaxedPrecision} : f32 loc(unknown)
memref.store %10, %arg2[%6, %8] : memref<2048x2048xf32> loc(unknown)
return loc(unknown)
} loc(unknown)
} loc(unknown)
generates the following cpp file
if defined(__HIP_PLATFORM_AMD__)
include <hip/hip_runtime.h>
using vfloatx2_t = float __attribute__((ext_vector_type(2)));
using vfloatx4_t = float __attribute__((ext_vector_type(4)));
using vfloatx16_t = float __attribute__((ext_vector_type(16)));
else
include "cuda_fp16.h"
endif // !defined(__HIP_PLATFORM_AMD__)
include <math.h>
include <stdint.h>
__global__ void gemm_naive_14479263422999410716(float (*arg0)[2048], float (*arg1)[2048], float (*arg2)[2048])
{
/*%0 = "gpu.thread_id"() {dimension = "x"} : () -> index*/
const uint threadIdx_x_0 = threadIdx.x;
/*%1 = "gpu.thread_id"() {dimension = "y"} : () -> index*/
const uint threadIdx_y_1 = threadIdx.y;
/*%2 = "gpu.block_id"() {dimension = "x"} : () -> index*/
const uint blockIdx_x_2 = blockIdx.x;
/*%3 = ...
- Merged PR 2376: [build] Install acc-lsp-server as an internal tool.
[Lisa Ong]
Removes acc-lsp-server from accera-compilers
Minor CMake macro renames to (hopefully) improve usability
- Merged PR 2378: [doc] Update doc links after DSL changes, fix missing
file warnings. [Lisa Ong]
Verified by:
cd <accera_root>
pip install mkdocs-material mkdocs-git-revision-date-plugin
mkdocs serve
- Merged PR 2377: Retire Benchmark.py, use hatlib for benchmarking and
shared library creation. [Lisa Ong]
This cleanup work precedes the actual work to produce static or dynamic libraries by migrating existing HAT Python scripts to consume hatlib. Next PRs will consume hatlib to produce those libraries.
hatlib defines a HAT package as .hat files and a library.
* Remove accera.tuning.AutoBenchmark and replace usages with hat.run_benchmark in case studies
* Removed accera.tuning.CorrectnessCheck. Baked correctness checking into accera.test.verifiers
* Disabled some tests in preparation for coming work (next PRs)
* parallelization tests: need to specify lomp as a link target dependency in the HAT file, and update hatlib to honor this flag
* emit_unpacked_buffer_tests: to resolve multi-MLIR-module scenario where we have a globals module in addition to the package module
Depends on this PR: https://github.com/microsoft/hat/pull/15
Related work items: 3556
- Merged PR 2374: Retain and honor the order of functions added to the
package. [Kern Handa]
Retain and honor the order of functions added to the package
Related work items: 3629
- Merged PR 2371: add lsp server for accera. [Abdul Dakkak]
this adds an lsp server to be used with the mlir vscode extension https://marketplace.visualstudio.com/items?itemName=llvm-vs-code-extensions.vscode-mlir . You will have to specify the lsp server in your settings.json . On my system this means to add the following setting
"mlir.server_path": "${workspaceFolder}/build/accera/acc-lsp-server/acc-lsp-server",
It's not super robust though
- Merged PR 2372: reduce install size. For example, on linux the install
size goes from 873M to 742M on Linux. [Abdul Dakkak]
reduce install size. For example, on linux the install dir goes from 873M to 742M. More can be done along those lines
- Merged PR 2369: run clang-format on acc_translate. [Abdul Dakkak]
run clang-format on acc_translate. There are no modifications to the code
- Merged PR 2367: Selectively emit GPU utilities. [Kern Handa]
Selectively emit GPU utilities
Related work items: 3559
- Merged PR 2366: [build] Fix manylinux package build. [Lisa Ong]
Apply updated requirements.txt without rebuilding docker image
- Merged PR 2365: Unify Package.add_function and Package.add_functions
into Package.add. [Kern Handa]
Related work items: 3549
- Merged PR 2363: Initial quickstart example in the main README. [Lisa
Ong]
* The quickstart example demonstrates how to do everything (including calling the function) from Python
* hatlib is now a runtime dependency as a result.
* We should consider updating at least the HelloMatMul Tutorials to also cover how to call functions from Python for quick testing. Calling from C++ is still the mainline scenario for performance
Dependent PR: https://github.com/microsoft/hat/pull/11
Related work items: 3630
- Merged PR 2364: Rename action plan references to plan. [Kern Handa]
Related work items: 3563
- Merged PR 2362: [hygiene] Move manylinux pipeline triggers from
classic to YAML. [Lisa Ong]
For maintainability, so that the triggers for that pipeline are in one place
- Merged PR 2350: LLVM update to 13.0.0. [Lisa Ong]
Updated LLVM to the "llvmorg-13.0.0" tag
Related work items: 3618
- Merged PR 2361: Build release versions of binaries for packaging
purposes, workaround auditwheel compression bug. [Lisa Ong]
* We currently build RelWithDebInfo instead of Release. This can result in packages that are too big to be uploaded to PyPI. A quick fix is to enable Release builds when invoked by the CI pipelines.
* Add triggering by tags for all pipelines that produce packages intended for PyPI (Windows, ManyLinux, macOS)
* Add pipeline to automate creating an LLVM build environment for the ManyLinux pipeline
* Revert to a last known good version of auditwheel (5.0.0) due to a compression bug (https://github.com/pypa/auditwheel/issues/366)
- [doc] tweaking public links. [Lisa Ong]
- Merged PR 2352: Reference github URLs for links in README.md. [Lisa
Ong]
README.md is referenced in PyPI, so these need to be fully-qualified URLs.
(The links will not work until the repo is published)
Related work items: 3619
- Merged PR 2360: Fix divide-by-0 crash when the active block exceeds
the vectorizable. [Mason Remy]
Fix divide-by-0 crash when the active block exceeds the vectorizable
size in the innermost dimension
- Add smoke test for this case. [Mason Remy]
- Fix divide-by-0 crash when the active block exceeds the vectorizable
size in the innermost dimension. [Mason Remy]
- Squashed commit of the following: [Lisa Ong]
commit add8396adc6e0f4e3cf0ae89796d08ac416c00a4
- Tweaked dark mode for better contrast, added favicon, improved
navigation. [Lisa Ong]
- Merged PR 2359: [docs] Fix rendering issues with code blocks and
bullet points. [Lisa Ong]
Also added sticky nav and tabs
- Minor typos in docs (4) [Lisa Ong]
* Update 00 Introduction.md
* Update mkdocs.yml
* Update Installing_accera_on_MacOS.md
* Update Installing_accera_on_Ubuntu.md
* Update Installing_on_MacOS.md
* Update Installing_on_Ubuntu.md
* Update Optimized_MatMul.md
* Update Hello_MatMul.md
* Update Cross_Compilation_PI3.md
- Add copyright. [Lisa Ong]
- Use mkdocs for documentation (3) [Lisa Ong, Lisa Ong]
* mkdocs integration
* add publishing workflow
* doc the doc
- Backport doc fixes from gh-pages to main (https://github.com/microsoft
/Accera/commit/ff491e3401691124b2aa6c3ee1d317bf264bdc11) [Lisa Ong]
- Merged PR 2345: Infer number of threads from the parallelization
indices. [Lisa Ong]
The number of threads was previously set to Target.num_threads.
This change treats Target.num_threads as a capacity setting, and infers the number of threads from an aggregate of:
* the number of unsplit indices
* the number of split blocks for each outermost index
This gives the user control over how many threads to request.
Examples:
* indices = i, j, k : 3 threads, one per index. Reason is that it doesn't make sense to just use 1 thread. For the future, we may want to add an explicit parameter to control the number of threads for this case
* indices = i, where ii = i.split(N//4): N//4 threads. We could have used ceiling(N/4), but due to loop unswitching, we don't directly apply the extra thread to the boundary loop. (future work?)
* indices = i, j, where ii = i.split(N//4): N//4 + 1 threads.
Implementation detail: if workshare loop collapsing happens because the indices are contiguous, the number of threads assigned is unaffected.
Related work items: 3554