Repository navigation
[Docs] State the compile boundary in the manifest guide as the entry derives it - #74
Merged
Merged
Conversation
…derives it The manifest guide still described a compile boundary as something an op class declares, and listed SymInt evaluability among the validator's checks. tile-ai/TileOPs#2520 derives the boundary from the manifest entry, and the validator checks only that compile_op_names agrees with the entry; the SymInt requirement follows from the cold fullgraph compile the op must pass.
There was a problem hiding this comment.
Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.
An organization owner can choose the repository organization as the billing source for reviews requested by members with a paid Copilot license. Learn about Copilot code review billing.
…ce_facts, and what a non-CUDA backend's tests and reports cover Follows tile-ai/TileOPs#2595. The backend guide said tuning never reaches a target, that the op layer rejects every unsupported dtype or shape, that the op tests run once with the backend's kernel and that reports are produced as usual; the last two did not hold on a device other than CUDA. The guide now states the two optional members a target's kernel may define, how a backend for another device type reports its device facts, which checks belong to the op layer, build_kernel and the kernel, the conditions for a backend's test claim about an op, and how benchmarks time and drop baselines on another device.
The dev and CI images moved out of the shared tileops-runner package into tileops-dev and tileops-ci-runner, tagged with the CUDA, PyTorch and TileLang versions, and the old package is deleted. The development guide and the backend example's docker command now name the new image.
…-device tuning and benchmark statements Follows tile-ai/TileOPs#2603's review. A reader of register_device_facts now receives the call's torch.device, which a cpu device can be; request_tune() is called again on every request and is idempotent per device; the op layer makes the call's device current only for the process's accelerator type; and a cpu device is not benchmarked, so contract item 7 and the summary after it hold on the devices a target runs on.
…eads, config and timing as the code does Contract items 2 and 7 each held several rules in one table cell. Item 2's three layers of checks now form their own table, and performance reports their own section. A device facts reader is called once per call spec, config is reported when non-empty and the op holds none, and CUDA events time a row only where CUPTI fails and the run allows the fallback.
…ts target serves the device tile-ai/TileOPs#2603 settled what a test needs: in_tree_kernels for the in-tree implementations, cuda_only for CUDA itself, and an architecture mark for a CUDA architecture, which skips elsewhere. The test-claim section named only the first two, and did not say the target has to serve the device.
…e and target tuning on main On TileOPs main a call with no tensor and no device runs on the process's accelerator, not on CUDA, and builds run with the call's device made current through torch.accelerator. A target's kernels receive tuning requests and report their config, and benchmarks run on any accelerator. The manifest, op and dispatch guides still described the CUDA-only behaviour and a tuning request that never reached a target. The backend guide states device facts as a table by device type, says read also runs when a kernel declaring supported_archs is built, gives the report timing and baseline rules as tables, and keys the memo by the call device. Chinese first; the English follows it.
…ements to where main holds them The op-layer contract promised every target an opaque operator and fake; only an op with a compile boundary gets one. A condition on tensor values is the target kernel's to check, and the op layer passes its error on. The timing and benchmark-writing pages, and the execution figure, described CUPTI as the only way a round is timed; another accelerator times with device events. The manifest guide's step 3 is split into one rule per item.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
has_compile_boundary()derives it from the manifest entry — a call-time tensor input and nocomposition— andmaybe_install_signatureinstalls it when the class is defined.validate_manifest.pychecks only thatcompile_op_namesagrees with the entry. The SymInt requirement follows from the coldfullgraph=Truecompile. The list renumbers from nine checks to eight, and thecompile_op_namescheck joins the code-vs-spec ones.calls,conceptsandextensions, English and Chinese.index.mdalready describes it this way, so these three pages contradicted it.This is commit
6290443, pushed todocs/op-base-guidetwenty minutes after #69 merged and stranded there. Verified against TileOPsf5d666e.docs/op-base-guidecan be deleted once this lands.Backend guide and user guides, for tile-ai/TileOPs#2603
Commits
e89df26,6b3d6b1,2478640,46d657a,cd9c557,d7243c8andf647a50match the guide to TileOPs main ate1293df2, which includes tile-ai/TileOPs#2603. Chinese first; the English follows it.backends: a target's kernel may definerequest_tune()andconfig; a kernel withoutrequest_tune()makes the op instance warn once.backends:register_device_factsgets its own section, with a table of where each device type's facts come from.readruns when a call spec first reads facts and when a kernel declaringsupported_archsis constructed.backends: contract item 6 applies to ops with a compile boundary. A condition on tensor values is the target kernel's to check, and it raisesValueError.backends: performance reports give timing and baseline handling as two tables. The test claim states the three markers and the numbers a claim comes with. The memo key is the call device.manifest/callsand its figure, andop/lifecycle: a call with no tensor and no device runs on the process accelerator's current device, and builds run with that device made current throughtorch.accelerator.op/records,op/structure,op/index: tuning requests reach a target's kernels, andkernel_config()reads them.manifest/index,timing,benchmark/writingand its figure: benchmarks run on CUDA or the process's accelerator, with CUPTI on CUDA and device events elsewhere.dispatch/backends: linksregister_kernel_typeto device facts, and a target's kernel to its optional members..pumlsources carry the same text.Dev image, for tile-ai/TileOPs#2602
Commit
f775414follows the image rename: the dev image isghcr.io/tile-ai/tileops-dev:cuda13.2-torch2.13-tilelang0.1.12, and the oldtileops-runnerpackage is deleted.user-guide/developmentnames the new package and its tag form, English and Chinese.backendsruns the example's tests in the new image; the oldtl-afcebed1-devtag no longer exists.timingandshared-memory-accesskeep the old image name: they record the image those numbers were measured on.