instruction improvements

IntelPython · napetrov · Jan 24, 2026 · Jan 24, 2026 · Jan 24, 2026 · Feb 17, 2026
commit a4d95a026080e7d490291c7e9efba9a4e89b7f53
@@ -7,19 +7,16 @@ applyTo:
 
 # Cython Instructions
 
-See [dpctl/AGENTS.md](/dpctl/AGENTS.md) for conventions and patterns.
+See `dpctl/AGENTS.md` for full conventions.
 
-## Quick Reference
-
-### Required Directives (after license header)
+## Required Directives (after license)
 ```cython
 # distutils: language = c++
 # cython: language_level=3
 # cython: linetrace=True
 ```
 
-### Key Rules
-- `cimport` for C-level declarations, `import` for Python
-- Store C refs as `_*_ref`, clean up in `__dealloc__`
+## Key Rules
+- `cimport` for C-level, `import` for Python-level
+- Store C refs as `_*_ref`, clean up in `__dealloc__` with NULL check
 - Use `with nogil:` for blocking C operations
-- Check NULL before using C API returns
@@ -10,17 +10,17 @@ applyTo:
 
 # DPCTL General Instructions
 
-See [/AGENTS.md](/AGENTS.md) for project overview and architecture.
+See `AGENTS.md` at repository root for project overview and architecture.
+Each major directory has its own `AGENTS.md` with specific conventions.
 
 ## Key References
 
-- **Code style:** See `.pre-commit-config.yaml` for tool versions, `.clang-format` for C++ style
-- **License:** Apache 2.0 with Intel copyright - see existing files for header format
-- **Directory guides:** Each major directory has its own `AGENTS.md`
+- **Code style:** `.pre-commit-config.yaml`, `.clang-format`, `.flake8`
+- **License:** Apache 2.0 with Intel copyright - match existing file headers
 
 ## Critical Rules
 
-1. **Device compatibility:** Not all devices support fp64/fp16 - check capabilities
-2. **Queue consistency:** All arrays in an operation must share compatible queues
-3. **Resource cleanup:** Always clean up C resources in `__dealloc__`
+1. **Device compatibility:** Not all devices support fp64/fp16 - never assume availability
+2. **Queue consistency:** Arrays in same operation must share compatible queues
+3. **Resource cleanup:** Clean up C resources in `__dealloc__` with NULL check
 4. **NULL checks:** Always check C API returns before use
diff --git a/.github/instructions/elementwise.instructions.md b/.github/instructions/elementwise.instructions.md
@@ -1,6 +1,7 @@
 ---
 applyTo:
-  - "dpctl/tensor/libtensor/**/elementwise_functions/**"
+  - "dpctl/tensor/libtensor/include/kernels/elementwise_functions/**"
+  - "dpctl/tensor/libtensor/source/elementwise_functions/**"
   - "dpctl/tensor/_elementwise_*.py"
   - "dpctl/tests/elementwise/**"
 ---
@@ -10,16 +11,14 @@ applyTo:
 Full stack: C++ kernel → pybind11 → Python wrapper → tests
 
 ## References
+- C++ kernels: `dpctl/tensor/libtensor/AGENTS.md`
+- Python wrappers: `dpctl/tensor/AGENTS.md`
+- Tests: `dpctl/tests/AGENTS.md`
 
-- **C++ kernels:** [dpctl/tensor/libtensor/AGENTS.md](/dpctl/tensor/libtensor/AGENTS.md)
-- **Python wrappers:** [dpctl/tensor/AGENTS.md](/dpctl/tensor/AGENTS.md)
-- **Tests:** [dpctl/tests/AGENTS.md](/dpctl/tests/AGENTS.md)
-
-## Adding New Operation Checklist
-
+## Adding New Operation
 1. `libtensor/include/kernels/elementwise_functions/op.hpp` - functor
 2. `libtensor/source/elementwise_functions/op.cpp` - dispatch tables
 3. Register in `tensor_elementwise.cpp`
 4. `_elementwise_funcs.py` - Python wrapper
 5. Export in `__init__.py`
-6. `tests/elementwise/test_op.py` - full coverage
+6. `tests/elementwise/test_op.py` - full dtype/usm coverage
@@ -7,19 +7,17 @@ applyTo:
 
 # C API Instructions
 
-See [libsyclinterface/AGENTS.md](/libsyclinterface/AGENTS.md) for conventions.
+See `libsyclinterface/AGENTS.md` for conventions.
 
-## Quick Reference
-
-### Naming
+## Naming
 `DPCTL<ClassName>_<MethodName>` (e.g., `DPCTLDevice_Create`)
 
-### Ownership annotations (from `Support/MemOwnershipAttrs.h`)
+## Ownership annotations (see `include/syclinterface/Support/MemOwnershipAttrs.h`)
 - `__dpctl_give` - caller must free
 - `__dpctl_take` - function takes ownership
 - `__dpctl_keep` - function only observes
 
-### Key Rules
+## Key Rules
 - Annotate all parameters and returns
 - Return NULL on failure
 - Use `DPCTL_API` for exports
diff --git a/.github/instructions/libtensor-cpp.instructions.md b/.github/instructions/libtensor-cpp.instructions.md
@@ -6,15 +6,11 @@ applyTo:
 
 # C++ SYCL Kernel Instructions
 
-See [dpctl/tensor/libtensor/AGENTS.md](/dpctl/tensor/libtensor/AGENTS.md) for patterns and structure.
+See `dpctl/tensor/libtensor/AGENTS.md` for patterns and directory structure.
 
-## Quick Reference
-
-### Type dispatch
-See `include/utils/type_dispatch_building.hpp` for the 14 supported types.
-
-### Key Rules
-- Kernel class names must be unique
+## Key Rules
+- Kernel class names must be globally unique
 - Use `if constexpr` for compile-time type branching
 - Complex types don't support vectorization
 - Return `nullptr` from factory for unsupported types
+- Check `include/kernels/elementwise_functions/common.hpp` for base patterns
@@ -6,16 +6,14 @@ applyTo:
 
 # USM Memory Instructions
 
-See [dpctl/memory/AGENTS.md](/dpctl/memory/AGENTS.md) for details.
+See `dpctl/memory/AGENTS.md` for details.
 
-## Quick Reference
-
-### USM Types
+## USM Types
 - `MemoryUSMDevice` - device-only (fastest)
 - `MemoryUSMShared` - host and device accessible
 - `MemoryUSMHost` - host memory, device accessible
 
-### Lifetime Rules
+## Lifetime Rules
 1. Memory is queue-bound
 2. Keep memory alive until operations complete
 3. Views extend base memory lifetime
diff --git a/.github/instructions/tensor-python.instructions.md b/.github/instructions/tensor-python.instructions.md
@@ -6,16 +6,14 @@ applyTo:
 
 # Tensor Python Instructions
 
-See [dpctl/tensor/AGENTS.md](/dpctl/tensor/AGENTS.md) for patterns.
+See `dpctl/tensor/AGENTS.md` for patterns.
 
-## Quick Reference
-
-### Queue validation (required for all operations)
+## Queue validation (required)
 ```python
 exec_q = dpctl.utils.get_execution_queue([x.sycl_queue, y.sycl_queue])
 if exec_q is None:
     raise ExecutionPlacementError("...")
 ```
 
-### Adding operations
-See checklist in [dpctl/tensor/AGENTS.md](/dpctl/tensor/AGENTS.md).
+## Adding operations
+See checklist in `dpctl/tensor/AGENTS.md`.
@@ -6,20 +6,18 @@ applyTo:
 
 # Testing Instructions
 
-See [dpctl/tests/AGENTS.md](/dpctl/tests/AGENTS.md) for patterns and dtype lists.
+See `dpctl/tests/AGENTS.md` for patterns.
 
-## Quick Reference
-
-### Essential helpers (from `helper/_helper.py`)
+## Essential helpers (from `helper/_helper.py`)
 ```python
 get_queue_or_skip()           # Create queue or skip
 skip_if_dtype_not_supported() # Skip if device lacks dtype
 ```
 
-### Dtype lists (from `elementwise/utils.py`)
-Use `_all_dtypes`, `_usm_types` for parametrization.
+## Dtype/USM lists
+Import from `elementwise/utils.py` - do not hardcode.
 
-### Coverage requirements
-- All supported dtypes
-- All USM types (device, shared, host)
-- Edge cases: empty arrays, scalars
+## Coverage
+- All dtypes from `_all_dtypes`
+- All USM types: device, shared, host
+- Edge cases: empty, scalar, broadcast
@@ -10,9 +10,10 @@
 ## Architecture
 
 ```
-Python API  →  Cython Bindings  →  C API  →  C++ Kernels  →  SYCL Runtime
-   dpctl/       _sycl_*.pyx      libsycl-    libtensor/
-                                 interface/
+Python API  →  Cython Bindings  →  C API           →  SYCL Runtime
+   dpctl/       _sycl_*.pyx       libsyclinterface/
+
+dpctl.tensor  →  pybind11  →  C++ Kernels (libtensor/)  →  SYCL Runtime
 ```
 
 ## Directory Guide
@@ -23,33 +24,36 @@ Python API  →  Cython Bindings  →  C API  →  C++ Kernels  →  SYCL Runtim
 | `dpctl/tensor/` | [dpctl/tensor/AGENTS.md](dpctl/tensor/AGENTS.md) | Array API tensor operations |
 | `dpctl/tensor/libtensor/` | [dpctl/tensor/libtensor/AGENTS.md](dpctl/tensor/libtensor/AGENTS.md) | C++ SYCL kernels |
 | `dpctl/memory/` | [dpctl/memory/AGENTS.md](dpctl/memory/AGENTS.md) | USM memory management |
+| `dpctl/program/` | [dpctl/program/AGENTS.md](dpctl/program/AGENTS.md) | SYCL kernel compilation |
+| `dpctl/utils/` | [dpctl/utils/AGENTS.md](dpctl/utils/AGENTS.md) | Utility functions |
 | `dpctl/tests/` | [dpctl/tests/AGENTS.md](dpctl/tests/AGENTS.md) | Test suite |
 | `libsyclinterface/` | [libsyclinterface/AGENTS.md](libsyclinterface/AGENTS.md) | C API layer |
 
 ## Code Style
 
-Style tools and configuration are defined in:
-- **Python/Cython:** `.pre-commit-config.yaml` (black, isort, flake8, cython-lint)
+Configuration files (do not hardcode versions - check these files):
+- **Python/Cython:** `.pre-commit-config.yaml`
 - **C/C++:** `.clang-format`
-- **Linting config:** `.flake8`
+- **Linting:** `.flake8`
 
 ## License Header
 
-All source files require Apache 2.0 header with Intel copyright. See existing files for format.
+All source files require Apache 2.0 header with Intel copyright. Reference existing files for exact format.
 
 ## Quick Reference
 
 ```python
 import dpctl
 import dpctl.tensor as dpt
 
-q = dpctl.SyclQueue("gpu")                    # Create queue
-x = dpt.ones((100, 100), dtype="f4", device=q) # Create array
-np_array = dpt.asnumpy(x)                      # Transfer to host
+q = dpctl.SyclQueue("gpu")                          # Create queue
+x = dpt.ones((100, 100), dtype="f4", sycl_queue=q)  # Create array
+np_array = dpt.asnumpy(x)                           # Transfer to host
 ```
 
 ## Key Concepts
 
 - **Queue:** Execution context binding device + context
 - **USM:** Unified Shared Memory (device/shared/host types)
 - **Filter string:** Device selector syntax `"backend:device_type:num"`
+- **Array API:** Python standard for array operations (https://data-apis.org/array-api/)
@@ -13,12 +13,14 @@ Python/Cython wrappers for SYCL runtime objects: Device, Queue, Context, Event,
 | `_sycl_context.pyx` | `SyclContext` wrapping `sycl::context` |
 | `_sycl_event.pyx` | `SyclEvent` wrapping `sycl::event` |
 | `_sycl_platform.pyx` | `SyclPlatform` wrapping `sycl::platform` |
+| `_sycl_device_factory.pyx` | Device enumeration and selection |
+| `_sycl_queue_manager.pyx` | Queue management utilities |
 | `_backend.pxd` | C API declarations from libsyclinterface |
 | `enum_types.py` | Python enums for SYCL types |
 
 ## Cython Conventions
 
-### Required Directives
+### Required Directives (after license header)
 ```cython
 # distutils: language = c++
 # cython: language_level=3
@@ -40,7 +42,7 @@ cdef class SyclDevice:
 
 ### Key Rules
 - Store C references as `_*_ref` attributes
-- Always clean up in `__dealloc__`
+- Always clean up in `__dealloc__` with NULL check
 - Use `with nogil:` for blocking C calls
 - Check NULL before using C API returns
 

@@ -0,0 +1,39 @@
+# dpctl/program/ - SYCL Kernel Compilation
+
+## Purpose
+
+Compile and manage SYCL kernels from OpenCL C or SPIR-V source.
+
+## Key Files
+
+| File | Purpose |
+|------|---------|
+| `_program.pyx` | `SyclProgram`, `SyclKernel` extension types |
+| `_program.pxd` | Cython declarations |
+| `__init__.py` | Public API exports |
+
+## Classes
+
+- **`SyclProgram`** - Compiled SYCL program containing one or more kernels
+- **`SyclKernel`** - Individual kernel extracted from a program
+
+## Usage Pattern
+
+```python
+from dpctl.program import create_program_from_source
+
+source = """
+__kernel void add(__global float* a, __global float* b, __global float* c) {
+    int i = get_global_id(0);
+    c[i] = a[i] + b[i];
+}
+"""
+
+program = create_program_from_source(queue, source)
+kernel = program.get_sycl_kernel("add")
+```
+
+## Notes
+
+- Programs are context-bound
+- Follows same Cython patterns as core dpctl (see `../AGENTS.md`)
diff --git a/dpctl/tensor/AGENTS.md b/dpctl/tensor/AGENTS.md
@@ -10,11 +10,15 @@ Python Array API-compliant tensor operations using `usm_ndarray` on SYCL devices
 |------|---------|
 | `_usmarray.pyx` | `usm_ndarray` extension type |
 | `_elementwise_funcs.py` | Elementwise operation wrappers |
-| `_elementwise_common.py` | `UnaryElementwiseFunc`, `BinaryElementwiseFunc` |
-| `_reduction.py` | Reduction operations (sum, prod, etc.) |
-| `_manipulation_functions.py` | reshape, concat, stack, etc. |
-| `_ctors.py` | Array constructors (empty, zeros, ones) |
+| `_elementwise_common.py` | `UnaryElementwiseFunc`, `BinaryElementwiseFunc` base classes |
+| `_reduction.py` | Reduction operations (sum, prod, mean, etc.) |
+| `_manipulation_functions.py` | reshape, concat, stack, split, etc. |
+| `_ctors.py` | Array constructors (empty, zeros, ones, arange, etc.) |
 | `_type_utils.py` | Type promotion and validation |
+| `_sorting.py` | sort, argsort |
+| `_searchsorted.py` | searchsorted, digitize |
+| `_dlpack.pyx` | DLPack interoperability |
+| `_copy_utils.py` | Copy and type casting |
 
 See [libtensor/AGENTS.md](libtensor/AGENTS.md) for C++ kernel implementation.
 
@@ -26,15 +30,13 @@ import dpctl.tensor._tensor_impl as ti
 
 abs = UnaryElementwiseFunc(
     "abs",                # Operation name
-    ti._abs_result_type,  # Type inference
-    ti._abs,              # Kernel implementation
+    ti._abs_result_type,  # Type inference function
+    ti._abs,              # Kernel implementation (from pybind11)
     _abs_docstring_
 )
 ```
 
-## Queue Validation
-
-All operations must validate queue compatibility:
+## Queue Validation (required for all operations)
 
 ```python
 exec_q = dpctl.utils.get_execution_queue([x.sycl_queue, y.sycl_queue])
@@ -44,9 +46,9 @@ if exec_q is None:
 
 ## Adding New Operations
 
-1. C++ kernel in `libtensor/include/kernels/`
-2. C++ source in `libtensor/source/`
-3. Register in `libtensor/source/tensor_elementwise.cpp`
-4. Python wrapper in `_elementwise_funcs.py`
+1. C++ kernel header: `libtensor/include/kernels/<category>/op.hpp`
+2. C++ source: `libtensor/source/<category>/op.cpp`
+3. Register in appropriate `tensor_*.cpp` entry point
+4. Python wrapper in appropriate `_*.py` module
 5. Export in `__init__.py`
-6. Tests in `../tests/elementwise/`
+6. Tests in `../tests/` with full dtype/usm coverage