Skip to content

Commit d95b2ae

Browse files
author
Steve Yalovitser
committed
initial commit
Please enter the commit message for your changes. Lines starting with '' will be ignored, and an empty message aborts the commit. On branch main Your branch is up to date with 'origin/main'. Changes to be committed: deleted: CONTRIBUTING.md new file: Config new file: GETTING_STARTED_GUIDE.md modified: LICENSE new file: LICENSE.txt modified: NOTICE modified: README.md new file: agents/neuron-nki-agent.agent-spec.json new file: agents/neuron-nki-agent.md new file: agents/neuron-nki-debugger-agent.agent-spec.json new file: agents/neuron-nki-debugger-agent.md new file: agents/neuron-nki-profile-analysis-agent.agent-spec.json new file: agents/neuron-nki-profile-analysis-agent.md new file: agents/neuron-nki-writer-agent.agent-spec.json new file: agents/neuron-nki-writer-agent.md new file: build new file: hooks/hooks.json new file: hooks/scripts/check-env.sh new file: pyproject.toml new file: setup.cfg new file: setup.py new file: skills/neuron-nki-debugging/SKILL.md new file: skills/neuron-nki-debugging/references/compiler-artifacts.md new file: skills/neuron-nki-debugging/references/compiler-error-codes.md new file: skills/neuron-nki-debugging/references/compiler-flags.md new file: skills/neuron-nki-debugging/references/ncc-memory-resource-errors.md new file: skills/neuron-nki-debugging/references/ncc-type-operation-errors.md new file: skills/neuron-nki-debugging/references/ncc-verification-errors.md new file: skills/neuron-nki-debugging/references/neuron-core-isolation.md new file: skills/neuron-nki-docs/SKILL.md new file: skills/neuron-nki-docs/references/architecture/nki_arch_guides.md new file: skills/neuron-nki-docs/references/architecture/trainium2_arch.md new file: skills/neuron-nki-docs/references/architecture/trainium3_arch.md new file: skills/neuron-nki-docs/references/architecture/trainium_inferentia2_arch.md new file: skills/neuron-nki-docs/references/debugging/error-codes-index.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EARG001.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EBVF030.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EHCA005.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EOOM001.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EOOM002.md new file: skills/neuron-nki-docs/references/debugging/error-codes/ESFH002.md new file: skills/neuron-nki-docs/references/debugging/error-codes/ESPP004.md new file: skills/neuron-nki-docs/references/debugging/error-codes/ESPP047.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EUOC002.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF001.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF004.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF005.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF006.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF007.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF009.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF010.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF011.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF013.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF015.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF016.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF017.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF018.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF019.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF022.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF024.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EVRF031.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EXSP001.md new file: skills/neuron-nki-docs/references/debugging/error-codes/EXTP004.md new file: skills/neuron-nki-docs/references/downloads/average_pool2d_jax.py new file: skills/neuron-nki-docs/references/downloads/average_pool2d_nki_kernels.py new file: skills/neuron-nki-docs/references/downloads/average_pool2d_torch.py new file: skills/neuron-nki-docs/references/downloads/mamba_nki_kernels.py new file: skills/neuron-nki-docs/references/downloads/mamba_torch.py new file: skills/neuron-nki-docs/references/downloads/matrix_multiplication_nki_kernels.py new file: skills/neuron-nki-docs/references/downloads/matrix_multiplication_torch.py new file: skills/neuron-nki-docs/references/downloads/spmd_multiple_nc_tensor_addition_jax.py new file: skills/neuron-nki-docs/references/downloads/spmd_multiple_nc_tensor_addition_nki_kernels.py new file: skills/neuron-nki-docs/references/downloads/spmd_multiple_nc_tensor_addition_torch.py new file: skills/neuron-nki-docs/references/downloads/spmd_tensor_addition_jax.py new file: skills/neuron-nki-docs/references/downloads/spmd_tensor_addition_nki_kernels.py new file: skills/neuron-nki-docs/references/downloads/spmd_tensor_addition_torch.py new file: skills/neuron-nki-docs/references/downloads/test_nki_isa_local_gather.py new file: skills/neuron-nki-docs/references/downloads/transpose2d_jax.py new file: skills/neuron-nki-docs/references/downloads/transpose2d_nki_kernels.py new file: skills/neuron-nki-docs/references/downloads/transpose2d_torch.py new file: skills/neuron-nki-docs/references/indices/hierarchical-toc.md new file: skills/neuron-nki-docs/references/indices/symbol-lookup.md new file: skills/neuron-nki-docs/references/indices/task-routing.md new file: skills/neuron-nki-docs/references/optimization/deep-dives-overview.md new file: skills/neuron-nki-docs/references/optimization/mxfp-matmul.md new file: skills/neuron-nki-docs/references/optimization/nki-beta-versions.md new file: skills/neuron-nki-docs/references/optimization/nki_perf_guide.md new file: skills/neuron-nki-docs/references/optimization/use-neuron-profile.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-isa-memory.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-isa-misc.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-isa-scalar.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-isa-tensor.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-isa-utility.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-isa-vector.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-language-creation.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-language-dims.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-language-memory.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-language-misc.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-language-types.md new file: skills/neuron-nki-docs/references/programming/api/api-nki-tools.md new file: skills/neuron-nki-docs/references/programming/api/api-overview.md new file: skills/neuron-nki-docs/references/programming/api/index.md new file: skills/neuron-nki-docs/references/programming/api/nki.api.shared.md new file: skills/neuron-nki-docs/references/programming/api/nki.isa.md new file: skills/neuron-nki-docs/references/programming/api/nki.language.md new file: skills/neuron-nki-docs/references/programming/api/nki.md new file: skills/neuron-nki-docs/references/programming/api/view-old-api-ref-pages.md new file: skills/neuron-nki-docs/references/programming/data-representation-overview.md new file: skills/neuron-nki-docs/references/programming/framework_custom_op.md new file: skills/neuron-nki-docs/references/programming/indexing-overview.md new file: skills/neuron-nki-docs/references/programming/lnc.md new file: skills/neuron-nki-docs/references/programming/memory-hierarchy-overview.md new file: skills/neuron-nki-docs/references/programming/nki-aps.md new file: skills/neuron-nki-docs/references/programming/nki-compiler.md new file: skills/neuron-nki-docs/references/programming/nki-dma-overview.md new file: skills/neuron-nki-docs/references/programming/nki-introduction.md new file: skills/neuron-nki-docs/references/programming/nki-language-guide.md new file: skills/neuron-nki-docs/references/programming/quickstart-implement-run-kernel.md new file: skills/neuron-nki-docs/references/programming/setup-env.md new file: skills/neuron-nki-docs/references/programming/tiling-overview.md new file: skills/neuron-nki-docs/references/programming/tutorial-use-a-prebuilt-kernel.md new file: skills/neuron-nki-docs/references/programming/tutorials/average_pool2d.md new file: skills/neuron-nki-docs/references/programming/tutorials/fused_mamba.md new file: skills/neuron-nki-docs/references/programming/tutorials/kernel-optimization.md new file: skills/neuron-nki-docs/references/programming/tutorials/matrix_multiplication.md new file: skills/neuron-nki-docs/references/programming/tutorials/spmd_multiple_nc_tensor_addition.md new file: skills/neuron-nki-docs/references/programming/tutorials/spmd_tensor_addition.md new file: skills/neuron-nki-docs/references/programming/tutorials/transpose2d.md new file: skills/neuron-nki-docs/references/programming/tutorials/tutorials.md new file: skills/neuron-nki-docs/references/reference/library/attention-cte.md new file: skills/neuron-nki-docs/references/reference/library/attention-tkg.md new file: skills/neuron-nki-docs/references/reference/library/design-rmsnorm-quant.md new file: skills/neuron-nki-docs/references/reference/library/mlp.md new file: skills/neuron-nki-docs/references/reference/library/output-projection-cte.md new file: skills/neuron-nki-docs/references/reference/library/output-projection-tkg.md new file: skills/neuron-nki-docs/references/reference/library/qkv.md new file: skills/neuron-nki-docs/references/reference/library/rmsnorm-quant.md new file: skills/neuron-nki-docs/references/reference/migration/nki-migration-guide.md new file: skills/neuron-nki-docs/references/reference/migration/nki_block_dimension_migration_guide.md new file: skills/neuron-nki-docs/references/reference/nki_faq.md new file: skills/neuron-nki-docs/references/reference/nki_rn.md new file: skills/neuron-nki-profile-querying/SKILL.md new file: skills/neuron-nki-profile-querying/references/example-bounds-analysis.md new file: skills/neuron-nki-profile-querying/references/getting-started.md new file: skills/neuron-nki-profile-querying/references/investigations/dma_efficiency.md new file: skills/neuron-nki-profile-querying/references/investigations/redundant_dma_transfers.md new file: skills/neuron-nki-profile-querying/references/investigations/redundant_te_transposes.md new file: skills/neuron-nki-profile-querying/references/investigations/te_inefficiency.md new file: skills/neuron-nki-profile-querying/references/performance-bounds.md new file: skills/neuron-nki-profile-querying/references/schema/ActiveTime.yaml new file: skills/neuron-nki-profile-querying/references/schema/DmaPacket.yaml new file: skills/neuron-nki-profile-querying/references/schema/DmaPacketAggregated.yaml new file: skills/neuron-nki-profile-querying/references/schema/DmaQueuesInfo.yaml new file: skills/neuron-nki-profile-querying/references/schema/DmaUsage.yaml new file: skills/neuron-nki-profile-querying/references/schema/Error.yaml new file: skills/neuron-nki-profile-querying/references/schema/Flow.yaml new file: skills/neuron-nki-profile-querying/references/schema/HbmUsage.yaml new file: skills/neuron-nki-profile-querying/references/schema/HbmUsageSummaryByType.yaml new file: skills/neuron-nki-profile-querying/references/schema/HostMemUsage.yaml new file: skills/neuron-nki-profile-querying/references/schema/Instruction.yaml new file: skills/neuron-nki-profile-querying/references/schema/KernelInstructions.yaml new file: skills/neuron-nki-profile-querying/references/schema/KernelIterationVariables.yaml new file: skills/neuron-nki-profile-querying/references/schema/KernelStackFrames.yaml new file: skills/neuron-nki-profile-querying/references/schema/Metadata.yaml new file: skills/neuron-nki-profile-querying/references/schema/NeffHeader.yaml new file: skills/neuron-nki-profile-querying/references/schema/PendingDma.yaml new file: skills/neuron-nki-profile-querying/references/schema/PsumUsage.yaml new file: skills/neuron-nki-profile-querying/references/schema/SbufAllocation.yaml new file: skills/neuron-nki-profile-querying/references/schema/SbufUsage.yaml new file: skills/neuron-nki-profile-querying/references/schema/SchemaFields.yaml new file: skills/neuron-nki-profile-querying/references/schema/SemaphoreUpdate.yaml new file: skills/neuron-nki-profile-querying/references/schema/Summary.yaml new file: skills/neuron-nki-profile-querying/references/schema/TensorInfo.yaml new file: skills/neuron-nki-profile-querying/references/schema/Throttle.yaml new file: skills/neuron-nki-profile-querying/references/schema/ThrottleSummary.yaml new file: skills/neuron-nki-profile-querying/references/schema/Warning.yaml new file: skills/neuron-nki-profiling/SKILL.md new file: skills/neuron-nki-profiling/examples/basic-profiling-workflow.py new file: skills/neuron-nki-profiling/references/metrics-reference.md new file: skills/neuron-nki-profiling/scripts/identify-neffs.py new file: skills/neuron-nki-profiling/scripts/setup-perfetto.sh new file: skills/neuron-nki-profiling/scripts/trace_processor new file: skills/neuron-nki-writing/SKILL.md new file: skills/neuron-nki-writing/examples/associative_scan.py new file: skills/neuron-nki-writing/examples/elementwise_exp.py new file: skills/neuron-nki-writing/examples/simple_matmul.py new file: skills/neuron-nki-writing/references/api-translation.md new file: skills/neuron-nki-writing/references/common-patterns.md new file: skills/neuron-nki-writing/references/indexing-patterns.md new file: skills/neuron-nki-writing/references/kernel-template.md new file: skills/neuron-nki-writing/references/memory-patterns.md new file: skills/neuron-nki-writing/references/nki-language-constraint.md new file: skills/neuron-nki-writing/references/nkilib/core/allocator.md new file: skills/neuron-nki-writing/references/nkilib/core/kernel-helpers.md new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/__init__.py new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/find_nonzero_indices.py new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/indexed_flatten.py new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/indexed_flatten_torch.py new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/layernorm_tkg.py new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/layernorm_torch.py new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/norm_tkg_utils.py new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/rmsnorm_mx_quantize_tkg.py new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/rmsnorm_tkg.py new file: skills/neuron-nki-writing/references/nkilib/core/subkernels/rmsnorm_torch.py new file: skills/neuron-nki-writing/references/nkilib/core/tensor-view.md new file: skills/neuron-nki-writing/references/nkilib/core/tile-info.md new file: skills/neuron-nki-writing/references/nkilib/core/tiled-range.md new file: skills/neuron-nki-writing/references/nkilib/core/utils/__init__.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/allocator.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/common_types.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/interleave_copy.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/kernel_assert.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/kernel_helpers.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/lnc_subscriptable.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/logging.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/modular_allocator.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/stream_shuffle_broadcast.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/tensor_view.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/tile_info.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/tiled_range.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/tp_broadcast.py new file: skills/neuron-nki-writing/references/nkilib/core/utils/tree_logger.py new file: skills/neuron-nki-writing/references/nkilib/ops/stream-shuffle-broadcast.md new file: skills/neuron-nki-writing/references/nkilib/ops/tp-broadcast.md new file: skills/neuron-nki-writing/references/nkilib/patterns/layout-conversion.md new file: skills/neuron-nki-writing/references/nkilib/patterns/moe-patterns.md new file: skills/neuron-nki-writing/references/nkilib/patterns/normalization-patterns.md new file: skills/neuron-nki-writing/references/nkilib/patterns/quantization-helpers.md new file: skills/neuron-nki-writing/references/nkilib/types/common-types.md new file: skills/neuron-nki-writing/references/nkilib/types/logging.md new file: skills/neuron-nki-writing/references/performance-basics.md new file: skills/neuron-nki-writing/references/transpose-and-layout.md new file: src/neuron_agentic_development.egg-info/PKG-INFO new file: src/neuron_agentic_development.egg-info/SOURCES.txt new file: src/neuron_agentic_development.egg-info/dependency_links.txt new file: src/neuron_agentic_development.egg-info/entry_points.txt new file: src/neuron_agentic_development.egg-info/top_level.txt new file: src/neuron_agentic_development/__init__.py new file: src/neuron_agentic_development/artifacts/agents/neuron-nki-agent.agent-spec.json new file: src/neuron_agentic_development/artifacts/agents/neuron-nki-agent.md new file: src/neuron_agentic_development/artifacts/agents/neuron-nki-debugger-agent.agent-spec.json new file: src/neuron_agentic_development/artifacts/agents/neuron-nki-debugger-agent.md new file: src/neuron_agentic_development/artifacts/agents/neuron-nki-profile-analysis-agent.agent-spec.json new file: src/neuron_agentic_development/artifacts/agents/neuron-nki-profile-analysis-agent.md new file: src/neuron_agentic_development/artifacts/agents/neuron-nki-writer-agent.agent-spec.json new file: src/neuron_agentic_development/artifacts/agents/neuron-nki-writer-agent.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/README.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/benchmarks/EVAL_REPORT.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/benchmarks/evals.json new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/conv2d_scaling_min.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/iterations/conv2d_scale_min_baseline.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/iterations/conv2d_scale_min_v1.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/iterations/conv2d_scale_min_v2.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/iterations/conv2d_scale_min_v3.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/iterations/conv2d_scale_min_v4.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/iterations/conv2d_scale_min_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/output/conv2d_scale_min_v4.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/optimization/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/output/conv2d_scale_min_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/output/test_conv2d_scale_min.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/output/test_full_scale.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/conv2d_scaling_min/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/cumsum.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/iterations/cumsum_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/iterations/cumsum_v1.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/iterations/cumsum_v2.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/iterations/cumsum_v3.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/iterations/cumsum_v4.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/iterations/cumsum_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/output/cumsum_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/optimization/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/output/cumsum_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/output/test_cumsum.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/cumsum/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/gelu.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/iterations/gelu_nki_kernels.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/iterations/gelu_v1.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/iterations/gelu_v2.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/iterations/gelu_v3.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/iterations/gelu_v4.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/iterations/gelu_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/output/gelu_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/optimization/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/output/gelu_nki_kernels.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/output/test_gelu.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/gelu/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/matmul_with_BT.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/iterations/matmul_bt_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/iterations/matmul_bt_v1.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/iterations/matmul_bt_v2.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/iterations/matmul_bt_v3.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/iterations/matmul_bt_v4.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/iterations/matmul_bt_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/output/matmul_bt_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/optimization/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/output/matmul_bt_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/output/test_matmul_bt.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/matmul_subtract_multiply_relu.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/iterations/matmul_subtract_multiply_relu_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/iterations/matmul_subtract_multiply_relu_v1.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/iterations/matmul_subtract_multiply_relu_v2.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/iterations/matmul_subtract_multiply_relu_v3.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/iterations/matmul_subtract_multiply_relu_v4.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/iterations/matmul_subtract_multiply_relu_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/output/matmul_subtract_multiply_relu_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/optimization/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/output/matmul_subtract_multiply_relu_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/output/test_matmul_subtract_multiply_relu.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/matmul_subtract_multiply_relu/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/mlp.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/iterations/swiglu_mlp_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/iterations/swiglu_mlp_v1.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/iterations/swiglu_mlp_v2.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/iterations/swiglu_mlp_v3.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/iterations/swiglu_mlp_v4.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/iterations/swiglu_mlp_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/output/swiglu_mlp_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/optimization/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/output/swiglu_mlp_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/output/test_full_mlp.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/output/test_large_mlp.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/output/test_swiglu_mlp.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/mlp/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/iterations/rms_fused_qkv_baseline.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/iterations/rms_fused_qkv_v1.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/iterations/rms_fused_qkv_v2.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/iterations/rms_fused_qkv_v3.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/iterations/rms_fused_qkv_v4.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/iterations/rms_fused_qkv_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/output/rms_fused_qkv_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/optimization/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/output/rms_fused_qkv_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/output/test_rms_fused_qkv.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/rms_fused_qkv/rms_fused_qkv.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/iterations/softmax_nki_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/iterations/softmax_v1.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/iterations/softmax_v2.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/iterations/softmax_v3.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/iterations/softmax_v4.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/iterations/softmax_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/output/softmax_v5.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/optimization/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/output/FINDINGS.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/output/output.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/output/softmax_nki_kernel.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/output/test_softmax.py new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/prompt.md new file: src/neuron_agentic_development/artifacts/evals/neuron-nki-agent/examples/softmax/softmax.py new file: src/neuron_agentic_development/artifacts/hooks/hooks.json new file: src/neuron_agentic_development/artifacts/hooks/scripts/check-env.sh new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/SKILL.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/STAGE0.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/STAGE1.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/STAGE2.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/STAGE3.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/STAGE4.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/STAGE5.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/STAGE6.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/STAGE7.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/cpu-component-debugging.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/cpu-e2e-debugging.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/debugging-case-study-gptoss.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/device-component-debugging.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/device-e2e-debugging.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/dump-tensors.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/equiv-concept.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/negative_samples/M2304_K256_N4096_seed42_d2a166b27b5e5073e755ff9ead80243e_debug_cpu32_gpu_trn_diff_qqplot_hist.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/negative_samples/M256_K128_N4096_seed42_916d411da0ea543e2d2afe79671115bb_debug_cpu32_gpu_trn_diff_qqplot_hist.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/negative_samples/Qwen3-8B-feedforward-module-mismatch-2.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/negative_samples/Qwen3-8B-feedforward-module-mismatch.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/negative_samples/bf16_3tensor_tp1_diff_qqplot_hist.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/negative_samples/rmsnorm.buggy_diff_qqplot_hist.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/positive_samples/M2304_K1024_N4096_seed42_481f8ec28e6c2b4f0050e4dee2b5b550_debug_cpu32_cpu_trn_diff_qqplot_hist.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/positive_samples/M2304_K1024_N4096_seed42_481f8ec28e6c2b4f0050e4dee2b5b550_debug_cpu32_gpu_trn_diff_qqplot_hist.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/positive_samples/M2304_K256_N4096_seed42_d2a166b27b5e5073e755ff9ead80243e_debug_cpu32_cpu_trn_rel_diff_qqplot_hist.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/positive_samples/M256_K128_N4096_seed42_916d411da0ea543e2d2afe79671115bb_debug_cpu32_cpu_trn_diff_qqplot_hist.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/example_plots/positive_samples/rmsnorm.fixed_diff_qqplot_hist.png new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/expected_structural_diffs.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/mapping_example.json new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/references/neuronxcc-debugging.md new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/detect_class_divergence.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/get_model_tree.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/run_calibration.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/run_stage0.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/run_stage1.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/run_stage2.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/run_stage3.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/run_stage5.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/run_stage6.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/run_stage7.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/run_teacher_forced_comparison.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/stage0_scaffolding.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/scripts/tensor_compare.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/templates/conftest_template.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/templates/diagnostic_forward_template.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/templates/equiv_config_template.json new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/templates/fast_tp_equiv_test_template.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/templates/neuron_device_validator_template.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/templates/run_inference_cpu_template.py new file: src/neuron_agentic_development/artifacts/skills/neuron-framework-equivalence/templates/test_template.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-debugging/SKILL.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-debugging/references/compiler-artifacts.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-debugging/references/compiler-error-codes.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-debugging/references/compiler-flags.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-debugging/references/ncc-memory-resource-errors.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-debugging/references/ncc-type-operation-errors.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-debugging/references/ncc-verification-errors.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-debugging/references/neuron-core-isolation.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/SKILL.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/architecture/nki_arch_guides.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/architecture/trainium2_arch.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/architecture/trainium3_arch.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/architecture/trainium_inferentia2_arch.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes-index.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EARG001.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EBVF030.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EHCA005.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EOOM001.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EOOM002.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/ESFH002.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/ESPP004.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/ESPP047.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EUOC002.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF001.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF004.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF005.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF006.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF007.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF009.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF010.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF011.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF013.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF015.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF016.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF017.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF018.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF019.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF022.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF024.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EVRF031.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EXSP001.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/debugging/error-codes/EXTP004.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/average_pool2d_jax.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/average_pool2d_nki_kernels.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/average_pool2d_torch.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/mamba_nki_kernels.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/mamba_torch.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/matrix_multiplication_nki_kernels.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/matrix_multiplication_torch.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/spmd_multiple_nc_tensor_addition_jax.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/spmd_multiple_nc_tensor_addition_nki_kernels.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/spmd_multiple_nc_tensor_addition_torch.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/spmd_tensor_addition_jax.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/spmd_tensor_addition_nki_kernels.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/spmd_tensor_addition_torch.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/test_nki_isa_local_gather.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/transpose2d_jax.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/transpose2d_nki_kernels.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/downloads/transpose2d_torch.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/indices/hierarchical-toc.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/indices/symbol-lookup.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/indices/task-routing.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/optimization/deep-dives-overview.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/optimization/mxfp-matmul.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/optimization/nki-beta-versions.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/optimization/nki_perf_guide.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/optimization/use-neuron-profile.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-isa-memory.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-isa-misc.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-isa-scalar.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-isa-tensor.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-isa-utility.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-isa-vector.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-language-creation.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-language-dims.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-language-memory.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-language-misc.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-language-types.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-nki-tools.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/api-overview.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/index.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/nki.api.shared.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/nki.isa.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/nki.language.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/nki.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/api/view-old-api-ref-pages.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/data-representation-overview.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/framework_custom_op.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/indexing-overview.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/lnc.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/memory-hierarchy-overview.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/nki-aps.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/nki-compiler.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/nki-dma-overview.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/nki-introduction.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/nki-language-guide.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/quickstart-implement-run-kernel.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/setup-env.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tiling-overview.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tutorial-use-a-prebuilt-kernel.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tutorials/average_pool2d.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tutorials/fused_mamba.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tutorials/kernel-optimization.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tutorials/matrix_multiplication.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tutorials/spmd_multiple_nc_tensor_addition.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tutorials/spmd_tensor_addition.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tutorials/transpose2d.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/programming/tutorials/tutorials.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/library/attention-cte.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/library/attention-tkg.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/library/design-rmsnorm-quant.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/library/mlp.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/library/output-projection-cte.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/library/output-projection-tkg.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/library/qkv.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/library/rmsnorm-quant.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/migration/nki-migration-guide.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/migration/nki_block_dimension_migration_guide.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/nki_faq.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-docs/references/reference/nki_rn.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/SKILL.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/example-bounds-analysis.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/getting-started.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/investigations/dma_efficiency.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/investigations/redundant_dma_transfers.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/investigations/redundant_te_transposes.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/investigations/te_inefficiency.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/performance-bounds.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/ActiveTime.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/DmaPacket.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/DmaPacketAggregated.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/DmaQueuesInfo.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/DmaUsage.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/Error.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/Flow.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/HbmUsage.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/HbmUsageSummaryByType.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/HostMemUsage.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/Instruction.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/KernelInstructions.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/KernelIterationVariables.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/KernelStackFrames.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/Metadata.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/NeffHeader.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/PendingDma.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/PsumUsage.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/SbufAllocation.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/SbufUsage.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/SchemaFields.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/SemaphoreUpdate.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/Summary.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/TensorInfo.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/Throttle.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/ThrottleSummary.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profile-querying/references/schema/Warning.yaml new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profiling/SKILL.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profiling/examples/basic-profiling-workflow.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profiling/references/metrics-reference.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profiling/scripts/identify-neffs.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profiling/scripts/setup-perfetto.sh new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-profiling/scripts/trace_processor new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/SKILL.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/examples/associative_scan.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/examples/elementwise_exp.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/examples/simple_matmul.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/api-translation.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/common-patterns.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/indexing-patterns.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/kernel-template.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/memory-patterns.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nki-language-constraint.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nkilib/core/allocator.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nkilib/core/kernel-helpers.md new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nkilib/core/subkernels/__init__.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nkilib/core/subkernels/find_nonzero_indices.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nkilib/core/subkernels/indexed_flatten.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nkilib/core/subkernels/indexed_flatten_torch.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nkilib/core/subkernels/layernorm_tkg.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nkilib/core/subkernels/layernorm_torch.py new file: src/neuron_agentic_development/artifacts/skills/neuron-nki-writing/references/nkilib/core/subkernels/norm_tkg_utils.py new file: src/neuron_agentic_developme…
1 parent 59d8d5c commit d95b2ae

647 files changed

Lines changed: 120236 additions & 69 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎CONTRIBUTING.md‎

Lines changed: 0 additions & 59 deletions
This file was deleted.

‎Config‎

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
package.NeuronAgenticDevelopment = {
2+
interfaces = (1.0);
3+
build-system = brazilpython;
4+
5+
build-tools = {
6+
1.0 = {
7+
AIMBuild = 1.0;
8+
BrazilPython = 3.0;
9+
BrazilPython-WheelBuilder = 1.0;
10+
NodeJS = 20.x;
11+
};
12+
};
13+
14+
dependencies = {
15+
1.0 = {
16+
AmazonBuilderCoreContext = 1.0;
17+
NeuroborosFoundations = 1.0;
18+
};
19+
};
20+
21+
test-dependencies = {
22+
1.0 = {
23+
BrazilPython-Pytest = any;
24+
Pytest = 8.x;
25+
BrazilPython-Pytest-Guard = any;
26+
};
27+
};
28+
29+
runtime-dependencies = {
30+
1.0 = {
31+
};
32+
};
33+
34+
targets = {
35+
NeuronAgenticDevelopment = {
36+
type = ai-capabilities;
37+
};
38+
python = {
39+
type = python;
40+
};
41+
};
42+
};

‎GETTING_STARTED_GUIDE.md‎

Lines changed: 180 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,180 @@
1+
# Getting Started with Neuron Agentic Development
2+
3+
Neuron Agentic Development provides a collection of AI agents and skills for developing custom NKI (Neuron Kernel Interface) kernels on AWS Trainium and Inferentia hardware. These capabilities equip coding agents in Kiro and Claude to author, debug, profile, and optimize NKI kernels — abstracting away the proprietary hardware knowledge required to work with Neuron.
4+
5+
## What are NKI Kernels?
6+
7+
NKI is the programming interface for writing custom compute kernels that run directly on AWS Trainium and Inferentia NeuronCores. NKI gives you fine-grained control over the hardware's tensor engines, vector engines, and DMA subsystem, enabling performance that exceeds what framework-level compilation alone can achieve.
8+
9+
Common use cases include fused attention kernels, custom normalization operations, quantized matrix multiplications, and any operation where the default compiler output leaves performance on the table.
10+
11+
## Prerequisites
12+
13+
| # | Requirement | Details | Needed for |
14+
|---|-------------|---------|------------|
15+
| 1 | Trainium/Inferentia instance | trn1, trn2, inf2 EC2 instance (AL2023 DLAMI recommended) | Compiling, profiling, optimization |
16+
| 2 | Neuron SDK | `aws-neuronx-tools` (pre-installed on DLAMI) | All on-device skills |
17+
| 3 | Python venv with Neuron packages | `neuronx-cc`, `torch-neuronx`, `neuron-explorer` | Compilation, profiling, analysis |
18+
| 4 | Kiro or Claude Code | Installed on the Trainium instance | Running agents and skills |
19+
| 5 | Anthropic API key | For Claude model inference | Agent reasoning |
20+
21+
> **Important:** The agent runs on the same machine as the hardware. There is no laptop-to-remote-box file transfer — everything is co-located. Writing and documentation skills work anywhere (no hardware needed), but profiling and debugging require on-instance execution.
22+
23+
## Step 1: Launch and Verify Your Trainium Instance
24+
25+
Launch a `trn2.3xlarge` instance in `São Paulo (sa-east-1)` or `Melbourne (ap-southeast-4)` using the [Neuron Deep Learning AMI (DLAMI)](https://awsdocs-neuron.readthedocs-hosted.com/en/latest/setup/pytorch/dlami.html), then SSH into the instance.
26+
27+
Verify the Neuron devices are visible:
28+
29+
```bash
30+
# Confirm Neuron devices are visible
31+
neuron-ls
32+
33+
# Confirm neuron-explorer is available
34+
which neuron-explorer && neuron-explorer --version
35+
```
36+
37+
## Step 2: Activate Your Python Environment
38+
39+
The DLAMI comes with a pre-installed virtual environment. Activate it:
40+
41+
```bash
42+
source ~/opt/aws_neuronx_venv_pytorch_2_9/bin/activate
43+
```
44+
45+
## Step 3: Install Neuron Agentic Development
46+
47+
**Option A — Install from Neuron PyPI repository:**
48+
49+
```bash
50+
pip install --upgrade neuron-agentic-development \
51+
--extra-index-url https://pip.repos.neuron.amazonaws.com
52+
```
53+
54+
**Option B — Clone from GitHub (for customization or contribution):**
55+
56+
```bash
57+
git clone https://github.com/aws-neuron/neuron-agentic-development.git
58+
cd neuron-agentic-development
59+
pip install .
60+
```
61+
62+
Then deploy to your preferred tool:
63+
64+
```bash
65+
# For Kiro
66+
deploy-neuron-agentic-development-to-kiro
67+
68+
# For Claude Code
69+
deploy-neuron-agentic-development-to-claude
70+
```
71+
72+
## Step 4: Install Kiro (if not already installed)
73+
74+
```bash
75+
curl -fsSL https://cli.kiro.dev/install | bash
76+
```
77+
78+
## Step 5: Start Using Neuron Agentic Development
79+
80+
```bash
81+
kiro-cli chat --agent neuron-nki-agent
82+
```
83+
84+
The `neuron-nki-agent` is the unified entry point. It automatically selects the right workflow based on your request and orchestrates the appropriate skills.
85+
86+
### Example Prompts
87+
88+
| What you want to do | What to say | Hardware needed? |
89+
|---------------------|-------------|------------------|
90+
| Write a new kernel | "Write a fused softmax kernel for bf16 inputs" | No |
91+
| Debug a compilation error | "Fix this kernel" (with error output) | Yes |
92+
| Profile a kernel | "Profile my kernel and show me the metrics" | Yes |
93+
| Analyze a profile | "What's the bottleneck in this kernel?" | Yes (neuron-explorer) |
94+
95+
## Skills
96+
97+
The package provides five specialized skills that follow the natural kernel development pipeline: **write → debug → profile → analyze**.
98+
99+
| # | Skill | Category | Use when |
100+
|---|-------|----------|----------|
101+
| 1 | `neuron-nki-writing` | Authoring | Writing new kernels or modifying existing ones |
102+
| 2 | `neuron-nki-debugging` | Debugging | Resolving compiler errors or numerical mismatches |
103+
| 3 | `neuron-nki-docs` | Documentation | Looking up API signatures, tutorials, error codes |
104+
| 4 | `neuron-nki-profiling` | Profiling | Capturing execution traces on hardware |
105+
| 5 | `neuron-nki-profile-querying` | Analysis | SQL-based performance bounds and bottleneck analysis |
106+
107+
### Kernel Authoring (`neuron-nki-writing`)
108+
109+
Your starting point for creating NKI kernels. Translates PyTorch, NumPy, or natural language descriptions into correct NKI code. Covers tiling strategies respecting hardware constraints (e.g., 128 partition dimension, 512/4096 PSUM free dimension), memory access patterns, and efficiency guidelines for DMA sizing and SBUF reuse.
110+
111+
### Debugging (`neuron-nki-debugging`)
112+
113+
Systematic workflow for resolving NKI compilation and execution errors. Covers environment setup with correct `--target` flags, compiler error resolution with a categorized index of all 28 NCC error codes, and numerical validation against CPU-computed references.
114+
115+
### Profiling (`neuron-nki-profiling`)
116+
117+
Captures execution profiles on hardware: configures runtime inspection environment variables, runs the kernel, identifies the correct NEFF (Neuron Execution File Format), captures the trace with `neuron-explorer` including DGE notifications for DMA-level detail, and extracts JSON metrics.
118+
119+
### Profile Analysis (`neuron-nki-profile-querying`)
120+
121+
Ingests NEFF/NTFF files and runs SQL queries to compute performance bounds, identify bottleneck engines, and localize inefficiencies to specific NKI source lines. Supports three analysis approaches: the neuron-explorer API server, DuckDB directly on parquet, or Pandas for custom computation.
122+
123+
### Documentation (`neuron-nki-docs`)
124+
125+
Used across all stages of development. Provides API signatures and tutorials during authoring, explains error codes during debugging, and clarifies hardware architecture details during profiling. Ask about any `nisa.*` or `nl.*` API, look up error codes, find tutorials, or browse architecture guides for Trainium 1/2/3.
126+
127+
## Agents
128+
129+
| # | Agent | Focus | What it does |
130+
|---|-------|-------|--------------|
131+
| 1 | `neuron-nki-agent` | Full lifecycle | Top-level entry point. Auto-selects the right workflow based on your request and orchestrates the appropriate skills. |
132+
| 2 | `neuron-nki-writer-agent` | Authoring | Translates PyTorch, NumPy, or natural language descriptions into NKI code. Handles modifications to existing kernels. |
133+
| 3 | `neuron-nki-debugger-agent` | Debugging | Autonomously analyzes compiler errors, searches documentation for fixes, and applies corrections. Tracks iterations (up to 10) and progressively simplifies when stuck. |
134+
| 4 | `neuron-nki-profile-analysis-agent` | Profiling + Analysis | Captures execution profiles on hardware, then runs SQL queries against profile data to compute performance bounds, identify bottleneck engines, and localize inefficiencies. |
135+
136+
## Architecture
137+
138+
```
139+
┌─────────────────────────────────────────────────┐
140+
│ Your Trainium Instance │
141+
│ │
142+
│ ┌──────────────┐ ┌───────────────────────┐ │
143+
│ │ Claude Code │───▶│ Neuron Agentic Dev │ │
144+
│ │ or Kiro │ │ • write kernel │ │
145+
│ └──────────────┘ │ • compile + debug │ │
146+
│ │ │ • profile (NEFF→NTFF) │ │
147+
│ │ │ • analyze (SQL/DuckDB)│ │
148+
│ ▼ │ • optimize + iterate │ │
149+
│ Anthropic API └───────────────────────┘ │
150+
│ (LLM inference) │ │
151+
│ ▼ │
152+
│ ┌────────────────────┐ │
153+
│ │ NeuronCores │ │
154+
│ │ (compile + execute)│ │
155+
│ └────────────────────┘ │
156+
└─────────────────────────────────────────────────┘
157+
```
158+
159+
## Things to Know
160+
161+
- Profiling and debugging skills require execution on actual Trainium/Inferentia hardware. The writing and docs skills work anywhere.
162+
- All skills target the current NKI Beta 3 API.
163+
- Skills support Trainium 1 (gen2), Trainium 2 (gen3), and Trainium 3 (gen4) with appropriate `--target` flags.
164+
- Skills and agents are designed to work together: the top-level agent automatically invokes profiling and debugging skills as needed.
165+
166+
## Sample Workflow: Custom Softmax Kernel
167+
168+
Here's a typical end-to-end workflow:
169+
170+
1. **Write** — Ask the agent: *"Write an NKI kernel that computes scaled softmax: softmax(x * scale) along the last dimension, for input shape [batch, seq_len, hidden_dim] in bfloat16."* The agent produces a complete kernel with proper tiling, hardware-accelerated exp, float32 accumulation, and bfloat16 output cast.
171+
172+
2. **Debug** — Ask the agent to run the kernel and verify numerical parity against a PyTorch reference. If compilation errors occur, the agent diagnoses and fixes them autonomously.
173+
174+
3. **Profile** — Point the agent at your kernel and ask it to profile. It compiles to a NEFF, benchmarks across multiple input shapes, and captures a hardware profile with `neuron-profile`.
175+
176+
4. **Analyze** — The agent runs a full bounds analysis on the captured profile and delivers a diagnosis: which engine is the bottleneck, what the stall rates are, and what optimization opportunities exist (e.g., DMA pipelining, double-buffering).
177+
178+
## Feedback
179+
180+
Give the Neuron Agentic Development capabilities a try and send feedback through [GitHub Issues](https://github.com/aws-neuron/neuron-agentic-development/issues) or your usual AWS Support contacts.

‎LICENSE‎

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -173,3 +173,30 @@
173173
defend, and hold each Contributor harmless for any liability
174174
incurred by, or claims asserted against, such Contributor by reason
175175
of your accepting any such warranty or additional liability.
176+
177+
END OF TERMS AND CONDITIONS
178+
179+
APPENDIX: How to apply the Apache License to your work.
180+
181+
To apply the Apache License to your work, attach the following
182+
boilerplate notice, with the fields enclosed by brackets "[]"
183+
replaced with your own identifying information. (Don't include
184+
the brackets!) The text should be enclosed in the appropriate
185+
comment syntax for the file format. We also recommend that a
186+
file or class name and description of purpose be included on the
187+
same "printed page" as the copyright notice for easier
188+
identification within third-party archives.
189+
190+
Copyright [yyyy] [name of copyright owner]
191+
192+
Licensed under the Apache License, Version 2.0 (the "License");
193+
you may not use this file except in compliance with the License.
194+
You may obtain a copy of the License at
195+
196+
http://www.apache.org/licenses/LICENSE-2.0
197+
198+
Unless required by applicable law or agreed to in writing, software
199+
distributed under the License is distributed on an "AS IS" BASIS,
200+
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
201+
See the License for the specific language governing permissions and
202+
limitations under the License.

0 commit comments

Comments
 (0)