mirror of
https://github.com/pytorch/pytorch.git
synced 2025-11-05 00:14:54 +08:00
Compare commits
44 Commits
ciflow/tru
...
cpp-docs-d
| Author | SHA1 | Date | |
|---|---|---|---|
| de4bd2b3a4 | |||
| 0571524a0e | |||
| 55f9503b47 | |||
| e122994d51 | |||
| 7f9450a68c | |||
| 6e0311b37e | |||
| 20f8edab38 | |||
| 3ef57af18f | |||
| 104b868618 | |||
| 94f2657c4b | |||
| 3f6538febd | |||
| f33abae695 | |||
| 73da7a40b6 | |||
| 335b5c7d4b | |||
| 76bb27e248 | |||
| a2da69385a | |||
| d177900723 | |||
| 61bcc8d75a | |||
| 1656b253c5 | |||
| 5d6230779d | |||
| a4077b568f | |||
| ae038f871b | |||
| defac66e39 | |||
| 061fa73c97 | |||
| 9501405de6 | |||
| e0791fc11d | |||
| e1d011d6eb | |||
| 3f5401020b | |||
| 5a3930abbc | |||
| a5f00077fc | |||
| 69fb3ebb5d | |||
| 1c4ced2eaf | |||
| 392acee68a | |||
| fee1ac927d | |||
| 4a7fefd7c7 | |||
| 3b4315940d | |||
| 3eddf04922 | |||
| 7c203b8420 | |||
| 3ca216ae17 | |||
| 9c22bbb2dc | |||
| 6268883f9c | |||
| 16212f0d6b | |||
| c8adc08b3b | |||
| 23b57a445c |
@ -13,3 +13,4 @@ exclude:
|
||||
- "**/benchmarks/**"
|
||||
- "**/test_*.py"
|
||||
- "**/*_test.py"
|
||||
- "tools/**"
|
||||
|
||||
@ -1,15 +1,11 @@
|
||||
sphinx==5.3.0
|
||||
sphinx==7.2.6
|
||||
#Description: This is used to generate PyTorch docs
|
||||
#Pinned versions: 5.3.0
|
||||
#Pinned versions: 7.2.6
|
||||
|
||||
standard-imghdr==3.13.0; python_version >= "3.13"
|
||||
#Description: This is needed by Sphinx, so it needs to be added here.
|
||||
# The reasons are as follows:
|
||||
# 1) This module has been removed from the Python standard library since Python 3.13(https://peps.python.org/pep-0594/#imghdr);
|
||||
# 2) The current version of Sphinx (5.3.0) is not compatible with Python 3.13.
|
||||
# Once Sphinx is upgraded to a version compatible with Python 3.13 or later, we can remove this dependency.
|
||||
pytorch_sphinx_theme2==0.2.0
|
||||
#Description: This is needed to generate PyTorch docs
|
||||
#Pinned versions: 0.2.0
|
||||
|
||||
-e git+https://github.com/pytorch/pytorch_sphinx_theme.git@71e55749be14ceb56e7f8211a9fb649866b87ad4#egg=pytorch_sphinx_theme2
|
||||
# TODO: sphinxcontrib.katex 0.9.0 adds a local KaTeX server to speed up pre-rendering
|
||||
# but it doesn't seem to work and hangs around idly. The initial thought that it is probably
|
||||
# something related to Docker setup. We can investigate this later.
|
||||
@ -36,17 +32,17 @@ tensorboard==2.18.0 ; python_version >= "3.13"
|
||||
#Description: This is used to generate PyTorch docs
|
||||
#Pinned versions: 2.13.0
|
||||
|
||||
breathe==4.34.0
|
||||
breathe==4.36.0
|
||||
#Description: This is used to generate PyTorch C++ docs
|
||||
#Pinned versions: 4.34.0
|
||||
#Pinned versions: 4.36.0
|
||||
|
||||
exhale==0.2.3
|
||||
exhale==0.3.7
|
||||
#Description: This is used to generate PyTorch C++ docs
|
||||
#Pinned versions: 0.2.3
|
||||
#Pinned versions: 0.3.7
|
||||
|
||||
docutils==0.16
|
||||
docutils==0.20
|
||||
#Description: This is used to generate PyTorch C++ docs
|
||||
#Pinned versions: 0.16
|
||||
#Pinned versions: 0.20
|
||||
|
||||
bs4==0.0.1
|
||||
#Description: This is used to generate PyTorch C++ docs
|
||||
@ -56,13 +52,13 @@ IPython==8.12.0
|
||||
#Description: This is used to generate PyTorch functorch docs
|
||||
#Pinned versions: 8.12.0
|
||||
|
||||
myst-nb==0.17.2
|
||||
myst-nb==1.3.0
|
||||
#Description: This is used to generate PyTorch functorch and torch.compile docs.
|
||||
#Pinned versions: 0.17.2
|
||||
#Pinned versions: 1.3.0
|
||||
|
||||
# The following are required to build torch.distributed.elastic.rendezvous.etcd* docs
|
||||
python-etcd==0.4.5
|
||||
sphinx-copybutton==0.5.0
|
||||
sphinx-design==0.4.0
|
||||
sphinx-design==0.6.1
|
||||
sphinxcontrib-mermaid==1.0.0
|
||||
myst-parser==0.18.1
|
||||
myst-parser==4.0.1
|
||||
|
||||
@ -89,20 +89,23 @@ if [ "$is_main_doc" = true ]; then
|
||||
|
||||
make coverage
|
||||
# Now we have the coverage report, we need to make sure it is empty.
|
||||
# Count the number of lines in the file and turn that number into a variable
|
||||
# $lines. The `cut -f1 ...` is to only parse the number, not the filename
|
||||
# Skip the report header by subtracting 2: the header will be output even if
|
||||
# there are no undocumented items.
|
||||
# Sphinx 7.2.6+ format: python.txt contains a statistics table with a TOTAL row
|
||||
# showing the undocumented count in the third column.
|
||||
# Example: | TOTAL | 99.83% | 2 |
|
||||
#
|
||||
# Also: see docs/source/conf.py for "coverage_ignore*" items, which should
|
||||
# be documented then removed from there.
|
||||
lines=$(wc -l build/coverage/python.txt 2>/dev/null |cut -f1 -d' ')
|
||||
undocumented=$((lines - 2))
|
||||
if [ $undocumented -lt 0 ]; then
|
||||
|
||||
# Extract undocumented count from TOTAL row in Sphinx 7.2.6 statistics table
|
||||
# The table format is: | Module | Coverage | Undocumented |
|
||||
# Extract the third column (undocumented count) from the TOTAL row
|
||||
undocumented=$(grep "| TOTAL" build/coverage/python.txt | awk -F'|' '{print $4}' | tr -d ' ')
|
||||
|
||||
if [ -z "$undocumented" ] || ! [[ "$undocumented" =~ ^[0-9]+$ ]]; then
|
||||
echo coverage output not found
|
||||
exit 1
|
||||
elif [ $undocumented -gt 0 ]; then
|
||||
echo undocumented objects found:
|
||||
elif [ "$undocumented" -gt 0 ]; then
|
||||
echo "undocumented objects found:"
|
||||
cat build/coverage/python.txt
|
||||
echo "Make sure you've updated relevant .rsts in docs/source!"
|
||||
echo "You can reproduce locally by running 'cd docs && make coverage && cat build/coverage/python.txt'"
|
||||
|
||||
@ -60,9 +60,11 @@ performance-*,
|
||||
readability-container-size-empty,
|
||||
readability-delete-null-pointer,
|
||||
readability-duplicate-include,
|
||||
readability-named-parameter,
|
||||
readability-misplaced-array-index,
|
||||
readability-redundant*,
|
||||
readability-simplify-subscript-expr,
|
||||
readability-static-definition-in-anonymous-namespace
|
||||
readability-string-compare,
|
||||
-readability-redundant-access-specifiers,
|
||||
-readability-redundant-control-flow,
|
||||
|
||||
319
.claude/skills/add-uint-support/SKILL.md
Normal file
319
.claude/skills/add-uint-support/SKILL.md
Normal file
@ -0,0 +1,319 @@
|
||||
---
|
||||
name: add-uint-support
|
||||
description: Add unsigned integer (uint) type support to PyTorch operators by updating AT_DISPATCH macros. Use when adding support for uint16, uint32, uint64 types to operators, kernels, or when user mentions enabling unsigned types, barebones unsigned types, or uint support.
|
||||
---
|
||||
|
||||
# Add Unsigned Integer (uint) Support to Operators
|
||||
|
||||
This skill helps add support for unsigned integer types (uint16, uint32, uint64) to PyTorch operators by updating their AT_DISPATCH macros.
|
||||
|
||||
## When to use this skill
|
||||
|
||||
Use this skill when:
|
||||
- Adding uint16, uint32, or uint64 support to an operator
|
||||
- User mentions "unsigned types", "uint support", "barebones unsigned types"
|
||||
- Enabling support for kUInt16, kUInt32, kUInt64 in kernels
|
||||
- Working with operator implementations that need expanded type coverage
|
||||
|
||||
## Quick reference
|
||||
|
||||
**Add unsigned types to existing dispatch:**
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_ALL_TYPES));
|
||||
|
||||
// After (method 1: add unsigned types explicitly)
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES));
|
||||
|
||||
// After (method 2: use V2 integral types if AT_INTEGRAL_TYPES present)
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_INTEGRAL_TYPES_V2), AT_EXPAND(AT_FLOATING_TYPES));
|
||||
```
|
||||
|
||||
## Type group reference
|
||||
|
||||
**Unsigned type groups:**
|
||||
- `AT_BAREBONES_UNSIGNED_TYPES`: kUInt16, kUInt32, kUInt64
|
||||
- `AT_INTEGRAL_TYPES_V2`: AT_INTEGRAL_TYPES + AT_BAREBONES_UNSIGNED_TYPES
|
||||
|
||||
**Relationship:**
|
||||
```cpp
|
||||
AT_INTEGRAL_TYPES // kByte, kChar, kInt, kLong, kShort
|
||||
AT_BAREBONES_UNSIGNED_TYPES // kUInt16, kUInt32, kUInt64
|
||||
AT_INTEGRAL_TYPES_V2 // INTEGRAL_TYPES + BAREBONES_UNSIGNED_TYPES
|
||||
```
|
||||
|
||||
## Instructions
|
||||
|
||||
### Step 1: Determine if conversion to V2 is needed
|
||||
|
||||
Check if the file uses AT_DISPATCH_V2:
|
||||
|
||||
**If using old AT_DISPATCH:**
|
||||
- First convert to AT_DISPATCH_V2 using the at-dispatch-v2 skill
|
||||
- Then proceed with adding uint support
|
||||
|
||||
**If already using AT_DISPATCH_V2:**
|
||||
- Proceed directly to Step 2
|
||||
|
||||
### Step 2: Analyze the current dispatch macro
|
||||
|
||||
Identify what type groups are currently in use:
|
||||
|
||||
```cpp
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
// body
|
||||
}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBFloat16);
|
||||
^^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
Current type coverage
|
||||
```
|
||||
|
||||
Common patterns:
|
||||
- `AT_EXPAND(AT_ALL_TYPES)` → includes AT_INTEGRAL_TYPES + AT_FLOATING_TYPES
|
||||
- `AT_EXPAND(AT_INTEGRAL_TYPES)` → signed integers only
|
||||
- `AT_EXPAND(AT_FLOATING_TYPES)` → floating point types
|
||||
|
||||
### Step 3: Choose the uint addition method
|
||||
|
||||
Two approaches:
|
||||
|
||||
**Method 1: Add AT_BAREBONES_UNSIGNED_TYPES explicitly**
|
||||
- Use when: You want to be explicit about adding uint support
|
||||
- Add `AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES)` to the type list
|
||||
|
||||
**Method 2: Substitute AT_INTEGRAL_TYPES with AT_INTEGRAL_TYPES_V2**
|
||||
- Use when: The dispatch already uses `AT_EXPAND(AT_INTEGRAL_TYPES)`
|
||||
- More concise: replaces one type group with its superset
|
||||
- Only applicable if AT_INTEGRAL_TYPES is present
|
||||
|
||||
### Step 4: Apply the transformation
|
||||
|
||||
**Method 1 example:**
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_V2(
|
||||
dtype,
|
||||
"min_values_cuda",
|
||||
AT_WRAP([&]() {
|
||||
kernel_impl<scalar_t>(iter);
|
||||
}),
|
||||
AT_EXPAND(AT_ALL_TYPES),
|
||||
kBFloat16, kHalf, kBool
|
||||
);
|
||||
|
||||
// After (add unsigned types)
|
||||
AT_DISPATCH_V2(
|
||||
dtype,
|
||||
"min_values_cuda",
|
||||
AT_WRAP([&]() {
|
||||
kernel_impl<scalar_t>(iter);
|
||||
}),
|
||||
AT_EXPAND(AT_ALL_TYPES),
|
||||
AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES),
|
||||
kBFloat16, kHalf, kBool
|
||||
);
|
||||
```
|
||||
|
||||
**Method 2 example:**
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_V2(
|
||||
dtype,
|
||||
"integral_op",
|
||||
AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}),
|
||||
AT_EXPAND(AT_INTEGRAL_TYPES)
|
||||
);
|
||||
|
||||
// After (substitute with V2)
|
||||
AT_DISPATCH_V2(
|
||||
dtype,
|
||||
"integral_op",
|
||||
AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}),
|
||||
AT_EXPAND(AT_INTEGRAL_TYPES_V2)
|
||||
);
|
||||
```
|
||||
|
||||
### Step 5: Handle AT_ALL_TYPES vs individual type groups
|
||||
|
||||
If the dispatch uses `AT_EXPAND(AT_ALL_TYPES)`:
|
||||
- `AT_ALL_TYPES` = `AT_INTEGRAL_TYPES` + `AT_FLOATING_TYPES`
|
||||
- To add uint: add `AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES)` to the list
|
||||
|
||||
If the dispatch separately lists INTEGRAL and FLOATING:
|
||||
```cpp
|
||||
// Before
|
||||
AT_EXPAND(AT_INTEGRAL_TYPES), AT_EXPAND(AT_FLOATING_TYPES)
|
||||
|
||||
// After (Method 2 preferred)
|
||||
AT_EXPAND(AT_INTEGRAL_TYPES_V2), AT_EXPAND(AT_FLOATING_TYPES)
|
||||
```
|
||||
|
||||
### Step 6: Verify all dispatch sites
|
||||
|
||||
Check the file for ALL dispatch macros that need uint support:
|
||||
- Some operators have multiple dispatch sites (CPU, CUDA, different functions)
|
||||
- Apply the transformation consistently across all sites
|
||||
- Ensure each gets the same type coverage updates
|
||||
|
||||
### Step 7: Validate the changes
|
||||
|
||||
Check that:
|
||||
- [ ] AT_DISPATCH_V2 format is used (not old AT_DISPATCH)
|
||||
- [ ] Unsigned types are added via one of the two methods
|
||||
- [ ] All relevant dispatch sites in the file are updated
|
||||
- [ ] Type groups use `AT_EXPAND()`
|
||||
- [ ] Arguments are properly formatted and comma-separated
|
||||
|
||||
## Common patterns
|
||||
|
||||
### Pattern 1: AT_ALL_TYPES + extras
|
||||
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBFloat16);
|
||||
|
||||
// After
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES), kHalf, kBFloat16);
|
||||
```
|
||||
|
||||
### Pattern 2: Separate INTEGRAL + FLOATING
|
||||
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_INTEGRAL_TYPES), AT_EXPAND(AT_FLOATING_TYPES));
|
||||
|
||||
// After
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_INTEGRAL_TYPES_V2), AT_EXPAND(AT_FLOATING_TYPES));
|
||||
```
|
||||
|
||||
### Pattern 3: Old dispatch needs conversion first
|
||||
|
||||
```cpp
|
||||
// Before (needs v2 conversion first)
|
||||
AT_DISPATCH_ALL_TYPES_AND2(kHalf, kBFloat16, dtype, "op", [&]() {
|
||||
kernel<scalar_t>();
|
||||
});
|
||||
|
||||
// After v2 conversion
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBFloat16);
|
||||
|
||||
// After adding uint support
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES), kHalf, kBFloat16);
|
||||
```
|
||||
|
||||
## Multiple dispatch sites example
|
||||
|
||||
For a file with multiple functions:
|
||||
|
||||
```cpp
|
||||
void min_values_kernel_cuda(TensorIterator& iter) {
|
||||
AT_DISPATCH_V2(iter.dtype(), "min_values_cuda", AT_WRAP([&]() {
|
||||
impl<scalar_t>(iter);
|
||||
}), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES), kBFloat16, kHalf);
|
||||
// ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
// Added uint support
|
||||
}
|
||||
|
||||
void min_launch_kernel(TensorIterator &iter) {
|
||||
AT_DISPATCH_V2(iter.input_dtype(), "min_cuda", AT_WRAP([&]() {
|
||||
gpu_reduce_kernel<scalar_t>(iter);
|
||||
}), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES), kBFloat16, kHalf);
|
||||
// ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
// Added uint support here too
|
||||
}
|
||||
```
|
||||
|
||||
## Decision tree
|
||||
|
||||
Use this decision tree to determine the approach:
|
||||
|
||||
```
|
||||
Is the file using AT_DISPATCH_V2?
|
||||
├─ No → Use at-dispatch-v2 skill first, then continue
|
||||
└─ Yes
|
||||
└─ Does it use AT_EXPAND(AT_INTEGRAL_TYPES)?
|
||||
├─ Yes → Replace with AT_EXPAND(AT_INTEGRAL_TYPES_V2)
|
||||
└─ No → Add AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES) to type list
|
||||
```
|
||||
|
||||
## Edge cases
|
||||
|
||||
### Case 1: Dispatch with only floating types
|
||||
|
||||
If the operator only supports floating point types, don't add uint support:
|
||||
|
||||
```cpp
|
||||
// Leave as-is - floating point only operator
|
||||
AT_DISPATCH_V2(dtype, "float_op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_FLOATING_TYPES), kHalf);
|
||||
```
|
||||
|
||||
### Case 2: Complex types present
|
||||
|
||||
Unsigned types work alongside complex types:
|
||||
|
||||
```cpp
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_ALL_TYPES),
|
||||
AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES),
|
||||
AT_EXPAND(AT_COMPLEX_TYPES),
|
||||
kHalf, kBFloat16);
|
||||
```
|
||||
|
||||
### Case 3: Already has uint support
|
||||
|
||||
Check if uint types are already present:
|
||||
- If `AT_INTEGRAL_TYPES_V2` is used → already has uint support
|
||||
- If `AT_BAREBONES_UNSIGNED_TYPES` is already in list → already has uint support
|
||||
- Skip the file if uint support is already present
|
||||
|
||||
## Workflow
|
||||
|
||||
When asked to add uint support:
|
||||
|
||||
1. Read the target file
|
||||
2. Check if using AT_DISPATCH_V2:
|
||||
- If not → use at-dispatch-v2 skill first
|
||||
3. Identify all dispatch macro sites
|
||||
4. For each dispatch:
|
||||
- Analyze current type groups
|
||||
- Choose method (add BAREBONES_UNSIGNED or upgrade to V2)
|
||||
- Apply transformation with Edit tool
|
||||
5. Show the user the changes
|
||||
6. Explain what was modified
|
||||
|
||||
## Important notes
|
||||
|
||||
- Always check if v2 conversion is needed first
|
||||
- Apply changes consistently across all dispatch sites in the file
|
||||
- Method 2 (AT_INTEGRAL_TYPES_V2) is cleaner when applicable
|
||||
- Method 1 (explicit AT_BAREBONES_UNSIGNED_TYPES) is more explicit
|
||||
- Unsigned types are: kUInt16, kUInt32, kUInt64 (not kByte which is uint8)
|
||||
- Some operators may not semantically support unsigned types - use judgment
|
||||
|
||||
## Testing
|
||||
|
||||
After adding uint support, the operator should accept uint16, uint32, and uint64 tensors. The user is responsible for functional testing.
|
||||
305
.claude/skills/at-dispatch-v2/SKILL.md
Normal file
305
.claude/skills/at-dispatch-v2/SKILL.md
Normal file
@ -0,0 +1,305 @@
|
||||
---
|
||||
name: at-dispatch-v2
|
||||
description: Convert PyTorch AT_DISPATCH macros to AT_DISPATCH_V2 format in ATen C++ code. Use when porting AT_DISPATCH_ALL_TYPES_AND*, AT_DISPATCH_FLOATING_TYPES*, or other dispatch macros to the new v2 API. For ATen kernel files, CUDA kernels, and native operator implementations.
|
||||
---
|
||||
|
||||
# AT_DISPATCH to AT_DISPATCH_V2 Converter
|
||||
|
||||
This skill helps convert PyTorch's legacy AT_DISPATCH macros to the new AT_DISPATCH_V2 format, as defined in `aten/src/ATen/Dispatch_v2.h`.
|
||||
|
||||
## When to use this skill
|
||||
|
||||
Use this skill when:
|
||||
- Converting AT_DISPATCH_* macros to AT_DISPATCH_V2
|
||||
- Porting ATen kernels to use the new dispatch API
|
||||
- Working with files in `aten/src/ATen/native/` that use dispatch macros
|
||||
- User mentions "AT_DISPATCH", "dispatch v2", "Dispatch_v2.h", or macro conversion
|
||||
|
||||
## Quick reference
|
||||
|
||||
**Old format:**
|
||||
```cpp
|
||||
AT_DISPATCH_ALL_TYPES_AND3(kBFloat16, kHalf, kBool, dtype, "kernel_name", [&]() {
|
||||
// lambda body
|
||||
});
|
||||
```
|
||||
|
||||
**New format:**
|
||||
```cpp
|
||||
AT_DISPATCH_V2(dtype, "kernel_name", AT_WRAP([&]() {
|
||||
// lambda body
|
||||
}), AT_EXPAND(AT_ALL_TYPES), kBFloat16, kHalf, kBool);
|
||||
```
|
||||
|
||||
## Key transformations
|
||||
|
||||
1. **Reorder arguments**: `scalar_type` and `name` come first, then lambda, then types
|
||||
2. **Wrap the lambda**: Use `AT_WRAP(lambda)` to handle internal commas
|
||||
3. **Expand type groups**: Use `AT_EXPAND(AT_ALL_TYPES)` instead of implicit expansion
|
||||
4. **List individual types**: Add extra types (kHalf, kBFloat16, etc.) after expanded groups
|
||||
5. **Add include**: `#include <ATen/Dispatch_v2.h>` near other Dispatch includes
|
||||
|
||||
## Instructions
|
||||
|
||||
### Step 1: Add the Dispatch_v2.h include
|
||||
|
||||
Add the v2 header near the existing `#include <ATen/Dispatch.h>`:
|
||||
|
||||
```cpp
|
||||
#include <ATen/Dispatch.h>
|
||||
#include <ATen/Dispatch_v2.h>
|
||||
```
|
||||
|
||||
Keep the old Dispatch.h include for now (other code may still need it).
|
||||
|
||||
### Step 2: Identify the old dispatch pattern
|
||||
|
||||
Common patterns to convert:
|
||||
|
||||
- `AT_DISPATCH_ALL_TYPES_AND{2,3,4}(type1, type2, ..., scalar_type, name, lambda)`
|
||||
- `AT_DISPATCH_FLOATING_TYPES_AND{2,3}(type1, type2, ..., scalar_type, name, lambda)`
|
||||
- `AT_DISPATCH_ALL_TYPES_AND_COMPLEX_AND{2,3}(type1, ..., scalar_type, name, lambda)`
|
||||
- `AT_DISPATCH_FLOATING_AND_COMPLEX_TYPES_AND{2,3}(type1, ..., scalar_type, name, lambda)`
|
||||
|
||||
### Step 3: Map the old macro to type groups
|
||||
|
||||
Identify which type group macro corresponds to the base types:
|
||||
|
||||
| Old macro base | AT_DISPATCH_V2 type group |
|
||||
|----------------|---------------------------|
|
||||
| `ALL_TYPES` | `AT_EXPAND(AT_ALL_TYPES)` |
|
||||
| `FLOATING_TYPES` | `AT_EXPAND(AT_FLOATING_TYPES)` |
|
||||
| `INTEGRAL_TYPES` | `AT_EXPAND(AT_INTEGRAL_TYPES)` |
|
||||
| `COMPLEX_TYPES` | `AT_EXPAND(AT_COMPLEX_TYPES)` |
|
||||
| `ALL_TYPES_AND_COMPLEX` | `AT_EXPAND(AT_ALL_TYPES_AND_COMPLEX)` |
|
||||
|
||||
For combined patterns, use multiple `AT_EXPAND()` entries:
|
||||
```cpp
|
||||
// Old: AT_DISPATCH_ALL_TYPES_AND_COMPLEX_AND2(...)
|
||||
// New: AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_COMPLEX_TYPES), type1, type2
|
||||
```
|
||||
|
||||
### Step 4: Extract the individual types
|
||||
|
||||
From `AT_DISPATCH_*_AND2(type1, type2, ...)` or `AT_DISPATCH_*_AND3(type1, type2, type3, ...)`, extract the individual types (type1, type2, etc.).
|
||||
|
||||
These become the trailing arguments after the type group:
|
||||
```cpp
|
||||
AT_DISPATCH_V2(..., AT_EXPAND(AT_ALL_TYPES), kBFloat16, kHalf, kBool)
|
||||
^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
Individual types from AND3
|
||||
```
|
||||
|
||||
### Step 5: Transform to AT_DISPATCH_V2
|
||||
|
||||
Apply the transformation:
|
||||
|
||||
**Pattern:**
|
||||
```cpp
|
||||
AT_DISPATCH_V2(
|
||||
scalar_type, // 1st: The dtype expression
|
||||
"name", // 2nd: The debug string
|
||||
AT_WRAP(lambda), // 3rd: The lambda wrapped in AT_WRAP
|
||||
type_groups, // 4th+: Type groups with AT_EXPAND()
|
||||
individual_types // Last: Individual types
|
||||
)
|
||||
```
|
||||
|
||||
**Example transformation:**
|
||||
```cpp
|
||||
// BEFORE
|
||||
AT_DISPATCH_ALL_TYPES_AND3(
|
||||
kBFloat16, kHalf, kBool,
|
||||
iter.dtype(),
|
||||
"min_values_cuda",
|
||||
[&]() {
|
||||
min_values_kernel_cuda_impl<scalar_t>(iter);
|
||||
}
|
||||
);
|
||||
|
||||
// AFTER
|
||||
AT_DISPATCH_V2(
|
||||
iter.dtype(),
|
||||
"min_values_cuda",
|
||||
AT_WRAP([&]() {
|
||||
min_values_kernel_cuda_impl<scalar_t>(iter);
|
||||
}),
|
||||
AT_EXPAND(AT_ALL_TYPES),
|
||||
kBFloat16, kHalf, kBool
|
||||
);
|
||||
```
|
||||
|
||||
### Step 6: Handle multi-line lambdas
|
||||
|
||||
For lambdas with internal commas or complex expressions, AT_WRAP is essential:
|
||||
|
||||
```cpp
|
||||
AT_DISPATCH_V2(
|
||||
dtype,
|
||||
"complex_kernel",
|
||||
AT_WRAP([&]() {
|
||||
gpu_reduce_kernel<scalar_t, scalar_t>(
|
||||
iter,
|
||||
MinOps<scalar_t>{},
|
||||
thrust::pair<scalar_t, int64_t>(upper_bound(), 0) // Commas inside!
|
||||
);
|
||||
}),
|
||||
AT_EXPAND(AT_ALL_TYPES)
|
||||
);
|
||||
```
|
||||
|
||||
### Step 7: Verify the conversion
|
||||
|
||||
Check that:
|
||||
- [ ] `AT_WRAP()` wraps the entire lambda
|
||||
- [ ] Type groups use `AT_EXPAND()`
|
||||
- [ ] Individual types don't have `AT_EXPAND()` (just `kBFloat16`, not `AT_EXPAND(kBFloat16)`)
|
||||
- [ ] Argument order is: scalar_type, name, lambda, types
|
||||
- [ ] Include added: `#include <ATen/Dispatch_v2.h>`
|
||||
|
||||
## Type group reference
|
||||
|
||||
Available type group macros (use with `AT_EXPAND()`):
|
||||
|
||||
```cpp
|
||||
AT_INTEGRAL_TYPES // kByte, kChar, kInt, kLong, kShort
|
||||
AT_FLOATING_TYPES // kDouble, kFloat
|
||||
AT_COMPLEX_TYPES // kComplexDouble, kComplexFloat
|
||||
AT_QINT_TYPES // kQInt8, kQUInt8, kQInt32
|
||||
AT_ALL_TYPES // INTEGRAL_TYPES + FLOATING_TYPES
|
||||
AT_ALL_TYPES_AND_COMPLEX // ALL_TYPES + COMPLEX_TYPES
|
||||
AT_INTEGRAL_TYPES_V2 // INTEGRAL_TYPES + unsigned types
|
||||
AT_BAREBONES_UNSIGNED_TYPES // kUInt16, kUInt32, kUInt64
|
||||
AT_FLOAT8_TYPES // Float8 variants
|
||||
```
|
||||
|
||||
## Common patterns
|
||||
|
||||
### Pattern: AT_DISPATCH_ALL_TYPES_AND2
|
||||
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_ALL_TYPES_AND2(kHalf, kBFloat16, dtype, "op", [&]() {
|
||||
kernel<scalar_t>(data);
|
||||
});
|
||||
|
||||
// After
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>(data);
|
||||
}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBFloat16);
|
||||
```
|
||||
|
||||
### Pattern: AT_DISPATCH_FLOATING_TYPES_AND3
|
||||
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_FLOATING_TYPES_AND3(kHalf, kBFloat16, kFloat8_e4m3fn,
|
||||
tensor.scalar_type(), "float_op", [&] {
|
||||
process<scalar_t>(tensor);
|
||||
});
|
||||
|
||||
// After
|
||||
AT_DISPATCH_V2(tensor.scalar_type(), "float_op", AT_WRAP([&] {
|
||||
process<scalar_t>(tensor);
|
||||
}), AT_EXPAND(AT_FLOATING_TYPES), kHalf, kBFloat16, kFloat8_e4m3fn);
|
||||
```
|
||||
|
||||
### Pattern: AT_DISPATCH_ALL_TYPES_AND_COMPLEX_AND2
|
||||
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_ALL_TYPES_AND_COMPLEX_AND2(
|
||||
kComplexHalf, kHalf,
|
||||
self.scalar_type(),
|
||||
"complex_op",
|
||||
[&] {
|
||||
result = compute<scalar_t>(self);
|
||||
}
|
||||
);
|
||||
|
||||
// After
|
||||
AT_DISPATCH_V2(
|
||||
self.scalar_type(),
|
||||
"complex_op",
|
||||
AT_WRAP([&] {
|
||||
result = compute<scalar_t>(self);
|
||||
}),
|
||||
AT_EXPAND(AT_ALL_TYPES),
|
||||
AT_EXPAND(AT_COMPLEX_TYPES),
|
||||
kComplexHalf,
|
||||
kHalf
|
||||
);
|
||||
```
|
||||
|
||||
## Edge cases
|
||||
|
||||
### Case 1: No extra types (rare)
|
||||
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_ALL_TYPES(dtype, "op", [&]() { kernel<scalar_t>(); });
|
||||
|
||||
// After
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_ALL_TYPES));
|
||||
```
|
||||
|
||||
### Case 2: Many individual types (AND4, AND5, etc.)
|
||||
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_FLOATING_TYPES_AND4(kHalf, kBFloat16, kFloat8_e4m3fn, kFloat8_e5m2,
|
||||
dtype, "float8_op", [&]() { kernel<scalar_t>(); });
|
||||
|
||||
// After
|
||||
AT_DISPATCH_V2(dtype, "float8_op", AT_WRAP([&]() {
|
||||
kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_FLOATING_TYPES), kHalf, kBFloat16, kFloat8_e4m3fn, kFloat8_e5m2);
|
||||
```
|
||||
|
||||
### Case 3: Lambda with no captures
|
||||
|
||||
```cpp
|
||||
// Before
|
||||
AT_DISPATCH_ALL_TYPES_AND2(kHalf, kBool, dtype, "op", []() {
|
||||
static_kernel<scalar_t>();
|
||||
});
|
||||
|
||||
// After
|
||||
AT_DISPATCH_V2(dtype, "op", AT_WRAP([]() {
|
||||
static_kernel<scalar_t>();
|
||||
}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBool);
|
||||
```
|
||||
|
||||
## Benefits of AT_DISPATCH_V2
|
||||
|
||||
1. **No arity in macro name**: Don't need different macros for AND2, AND3, AND4
|
||||
2. **Composable type sets**: Mix and match type groups with `AT_EXPAND()`
|
||||
3. **Extensible**: Easy to add more types without hitting macro limits
|
||||
4. **Clearer**: Type groups are explicit, not implicit in macro name
|
||||
|
||||
## Important notes
|
||||
|
||||
- Keep `#include <ATen/Dispatch.h>` - other code may need it
|
||||
- The `AT_WRAP()` is mandatory - prevents comma parsing issues in the lambda
|
||||
- Type groups need `AT_EXPAND()`, individual types don't
|
||||
- The v2 API is in `aten/src/ATen/Dispatch_v2.h` - refer to it for full docs
|
||||
- See the header file for the Python script to regenerate the macro implementation
|
||||
|
||||
## Workflow
|
||||
|
||||
When asked to convert AT_DISPATCH macros:
|
||||
|
||||
1. Read the file to identify all AT_DISPATCH uses
|
||||
2. Add `#include <ATen/Dispatch_v2.h>` if not present
|
||||
3. For each dispatch macro:
|
||||
- Identify the pattern and extract components
|
||||
- Map the base type group
|
||||
- Extract individual types
|
||||
- Construct the AT_DISPATCH_V2 call
|
||||
- Apply with Edit tool
|
||||
4. Show the user the complete converted file
|
||||
5. Explain what was changed
|
||||
|
||||
Do NOT compile or test the code - focus on accurate conversion only.
|
||||
2
.github/ci_commit_pins/xla.txt
vendored
2
.github/ci_commit_pins/xla.txt
vendored
@ -1 +1 @@
|
||||
df6798dfb931ce7c7fe5bed2447cd1092a5981af
|
||||
c8b09f5f77d6bf6fb7ed7a9aa83e5d8156b3a5e9
|
||||
|
||||
15
.github/workflows/lint.yml
vendored
15
.github/workflows/lint.yml
vendored
@ -76,11 +76,12 @@ jobs:
|
||||
|
||||
# NOTE: mypy needs its own job because it depends on --all-files, without assessing all files it sometimes
|
||||
# fails to find types when it should
|
||||
lintrunner-mypy:
|
||||
# NOTE: We should be able to disable this and consolidate with Pyrefly
|
||||
lintrunner-pyrefly:
|
||||
uses: pytorch/test-infra/.github/workflows/linux_job_v2.yml@main
|
||||
name: lintrunner-mypy-${{ needs.get-changed-files.outputs.changed-files == '*' && 'all' || 'partial' }}
|
||||
name: lintrunner-pyrefly-${{ needs.get-changed-files.outputs.changed-files == '*' && 'all' || 'partial' }}
|
||||
needs: [get-label-type, get-changed-files]
|
||||
# Only run if there are changed files relevant to mypy
|
||||
# Only run if there are changed files relevant to pyrefly
|
||||
if: |
|
||||
github.repository_owner == 'pytorch' && (
|
||||
needs.get-changed-files.outputs.changed-files == '*' ||
|
||||
@ -98,8 +99,8 @@ jobs:
|
||||
ref: ${{ github.event_name == 'pull_request' && github.event.pull_request.head.sha || github.sha }}
|
||||
script: |
|
||||
CHANGED_FILES="${{ needs.get-changed-files.outputs.changed-files }}"
|
||||
echo "Running mypy"
|
||||
ADDITIONAL_LINTRUNNER_ARGS="--take MYPY,MYPYSTRICT --all-files" .github/scripts/lintrunner.sh
|
||||
echo "Running pyrefly"
|
||||
ADDITIONAL_LINTRUNNER_ARGS="--take PYREFLY --all-files" .github/scripts/lintrunner.sh
|
||||
|
||||
lintrunner-noclang:
|
||||
uses: pytorch/test-infra/.github/workflows/linux_job_v2.yml@main
|
||||
@ -118,9 +119,9 @@ jobs:
|
||||
CHANGED_FILES="${{ needs.get-changed-files.outputs.changed-files }}"
|
||||
echo "Running all other linters"
|
||||
if [ "$CHANGED_FILES" = '*' ]; then
|
||||
ADDITIONAL_LINTRUNNER_ARGS="--skip CLANGTIDY,CLANGFORMAT,MYPY,MYPYSTRICT,PYREFLY --all-files" .github/scripts/lintrunner.sh
|
||||
ADDITIONAL_LINTRUNNER_ARGS="--skip CLANGTIDY,CLANGFORMAT,PYREFLY --all-files" .github/scripts/lintrunner.sh
|
||||
else
|
||||
ADDITIONAL_LINTRUNNER_ARGS="--skip CLANGTIDY,CLANGFORMAT,MYPY,MYPYSTRICT,PYREFLY ${CHANGED_FILES}" .github/scripts/lintrunner.sh
|
||||
ADDITIONAL_LINTRUNNER_ARGS="--skip CLANGTIDY,CLANGFORMAT,PYREFLY ${CHANGED_FILES}" .github/scripts/lintrunner.sh
|
||||
fi
|
||||
|
||||
quick-checks:
|
||||
|
||||
1
.gitignore
vendored
1
.gitignore
vendored
@ -398,3 +398,4 @@ CLAUDE.local.md
|
||||
/test_*.py
|
||||
/debug_*.py
|
||||
CLAUDE_CONTEXT/
|
||||
/.claude/settings.local.json
|
||||
|
||||
@ -121,94 +121,6 @@ command = [
|
||||
]
|
||||
is_formatter = true
|
||||
|
||||
[[linter]]
|
||||
code = 'MYPY'
|
||||
include_patterns = [
|
||||
'setup.py',
|
||||
'functorch/dim/**/*.py',
|
||||
'torch/**/*.py',
|
||||
'torch/**/*.pyi',
|
||||
'caffe2/**/*.py',
|
||||
'caffe2/**/*.pyi',
|
||||
'test/test_bundled_images.py',
|
||||
'test/test_bundled_inputs.py',
|
||||
'test/test_complex.py',
|
||||
'test/test_datapipe.py',
|
||||
'test/test_futures.py',
|
||||
'test/test_numpy_interop.py',
|
||||
'test/test_torch.py',
|
||||
'test/test_type_hints.py',
|
||||
'test/test_type_info.py',
|
||||
'test/test_utils.py',
|
||||
]
|
||||
exclude_patterns = [
|
||||
'**/fb/**',
|
||||
]
|
||||
command = [
|
||||
'python3',
|
||||
'tools/linter/adapters/mypy_linter.py',
|
||||
'--config=mypy.ini',
|
||||
'--',
|
||||
'@{{PATHSFILE}}'
|
||||
]
|
||||
init_command = [
|
||||
'python3',
|
||||
'tools/linter/adapters/pip_init.py',
|
||||
'--dry-run={{DRYRUN}}',
|
||||
'numpy==1.26.4 ; python_version >= "3.10" and python_version <= "3.11"',
|
||||
'numpy==2.1.0 ; python_version >= "3.12"',
|
||||
'expecttest==0.3.0',
|
||||
'mypy==1.16.0',
|
||||
'sympy==1.13.3',
|
||||
'types-requests==2.27.25',
|
||||
'types-pyyaml==6.0.2',
|
||||
'types-tabulate==0.8.8',
|
||||
'types-protobuf==5.29.1.20250403',
|
||||
'types-setuptools==79.0.0.20250422',
|
||||
'types-jinja2==2.11.9',
|
||||
'types-colorama==0.4.6',
|
||||
'filelock==3.18.0',
|
||||
'junitparser==2.1.1',
|
||||
'rich==14.1.0',
|
||||
'pyyaml==6.0.2',
|
||||
'optree==0.13.0',
|
||||
'dataclasses-json==0.6.7',
|
||||
'pandas==2.2.3',
|
||||
]
|
||||
|
||||
[[linter]]
|
||||
code = 'MYPYSTRICT'
|
||||
include_patterns = [
|
||||
'.github/**/*.py',
|
||||
'benchmarks/instruction_counts/**/*.py',
|
||||
'tools/**/*.py',
|
||||
'torchgen/**/*.py',
|
||||
'torch/utils/_pytree.py',
|
||||
'torch/utils/_cxx_pytree.py',
|
||||
'torch/utils/benchmark/utils/common.py',
|
||||
'torch/utils/benchmark/utils/timer.py',
|
||||
'torch/utils/benchmark/utils/valgrind_wrapper/**/*.py',
|
||||
]
|
||||
exclude_patterns = [
|
||||
# (linbinyu) copied from internal repo
|
||||
'**/fb/**',
|
||||
'tools/code_analyzer/gen_operators_yaml.py',
|
||||
'tools/dynamo/verify_dynamo.py',
|
||||
'tools/gen_vulkan_spv.py',
|
||||
'tools/test/gen_operators_yaml_test.py',
|
||||
'tools/test/gen_oplist_test.py',
|
||||
'tools/test/test_selective_build.py',
|
||||
'tools/experimental/torchfuzz/**',
|
||||
]
|
||||
command = [
|
||||
'python3',
|
||||
'tools/linter/adapters/mypy_linter.py',
|
||||
'--config=mypy-strict.ini',
|
||||
'--code=MYPYSTRICT',
|
||||
'--',
|
||||
'@{{PATHSFILE}}'
|
||||
]
|
||||
|
||||
|
||||
[[linter]]
|
||||
code = 'PYREFLY'
|
||||
@ -230,6 +142,7 @@ init_command = [
|
||||
'python3',
|
||||
'tools/linter/adapters/pip_init.py',
|
||||
'--dry-run={{DRYRUN}}',
|
||||
'numpy==1.26.4 ; python_version >= "3.10" and python_version <= "3.11"',
|
||||
'numpy==2.1.0 ; python_version >= "3.12"',
|
||||
'expecttest==0.3.0',
|
||||
'pyrefly==0.36.2',
|
||||
|
||||
@ -11,7 +11,6 @@ aspects of contributing to PyTorch.
|
||||
<!-- toc -->
|
||||
|
||||
- [Developing PyTorch](#developing-pytorch)
|
||||
- [Setup the development environment](#setup-the-development-environment)
|
||||
- [Tips and Debugging](#tips-and-debugging)
|
||||
- [Nightly Checkout & Pull](#nightly-checkout--pull)
|
||||
- [Codebase structure](#codebase-structure)
|
||||
@ -67,23 +66,6 @@ aspects of contributing to PyTorch.
|
||||
|
||||
Follow the instructions for [installing PyTorch from source](https://github.com/pytorch/pytorch#from-source). If you get stuck when developing PyTorch on your machine, check out the [tips and debugging](#tips-and-debugging) section below for common solutions.
|
||||
|
||||
### Setup the development environment
|
||||
|
||||
First, you need to [fork the PyTorch project on GitHub](https://github.com/pytorch/pytorch/fork) and follow the instructions at [Connecting to GitHub with SSH](https://docs.github.com/en/authentication/connecting-to-github-with-ssh) to setup your SSH authentication credentials.
|
||||
|
||||
Then clone the PyTorch project and setup the development environment:
|
||||
|
||||
```bash
|
||||
git clone git@github.com:<USERNAME>/pytorch.git
|
||||
cd pytorch
|
||||
git remote add upstream git@github.com:pytorch/pytorch.git
|
||||
|
||||
make setup-env
|
||||
# Or run `make setup-env-cuda` for pre-built CUDA binaries
|
||||
# Or run `make setup-env-rocm` for pre-built ROCm binaries
|
||||
source venv/bin/activate # or `. .\venv\Scripts\activate` on Windows
|
||||
```
|
||||
|
||||
### Tips and Debugging
|
||||
|
||||
* If you want to have no-op incremental rebuilds (which are fast), see [Make no-op build fast](#make-no-op-build-fast) below.
|
||||
|
||||
@ -181,7 +181,7 @@ c10::intrusive_ptr<c10::TensorImpl> CPUGeneratorImpl::get_state() const {
|
||||
static const size_t size = sizeof(CPUGeneratorImplState);
|
||||
static_assert(std::is_standard_layout_v<CPUGeneratorImplState>, "CPUGeneratorImplState is not a PODType");
|
||||
|
||||
auto state_tensor = at::detail::empty_cpu({(int64_t)size}, ScalarType::Byte, std::nullopt, std::nullopt, std::nullopt, std::nullopt);
|
||||
auto state_tensor = at::detail::empty_cpu({static_cast<int64_t>(size)}, ScalarType::Byte, std::nullopt, std::nullopt, std::nullopt, std::nullopt);
|
||||
auto rng_state = state_tensor.data_ptr();
|
||||
|
||||
// accumulate generator data to be copied into byte tensor
|
||||
|
||||
@ -223,7 +223,7 @@ void Context::setSDPPriorityOrder(const std::vector<int64_t>& order) {
|
||||
"setSDPPriority order expected ", sdp_priority_order.size() - 1, " but got ",
|
||||
at::num_sdp_backends, " unique backends specified in priority order.");
|
||||
for (uint32_t i = 0; i < order.size(); i++) {
|
||||
sdp_priority_order[i] = (at::SDPBackend) order[i];
|
||||
sdp_priority_order[i] = static_cast<at::SDPBackend>(order[i]);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@ -197,6 +197,7 @@ inline at::ScalarType scalar_type(at::ScalarType s) {
|
||||
/* don't use TYPE again in case it is an expensive or side-effect op */ \
|
||||
at::ScalarType _st = ::detail::scalar_type(the_type); \
|
||||
RECORD_KERNEL_FUNCTION_DTYPE(at_dispatch_name, _st); \
|
||||
C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum") \
|
||||
switch (_st) { \
|
||||
__VA_ARGS__ \
|
||||
default: \
|
||||
@ -208,6 +209,7 @@ inline at::ScalarType scalar_type(at::ScalarType s) {
|
||||
toString(_st), \
|
||||
"'"); \
|
||||
} \
|
||||
C10_DIAGNOSTIC_POP() \
|
||||
}()
|
||||
|
||||
#define AT_DISPATCH_CASE_FLOATING_TYPES(...) \
|
||||
|
||||
@ -252,13 +252,13 @@ MapAllocator::MapAllocator(WithFd /*unused*/, std::string_view filename, int fd,
|
||||
if (!(flags_ & ALLOCATOR_MAPPED_FROMFD)) {
|
||||
if (flags_ & ALLOCATOR_MAPPED_SHARED) {
|
||||
// NOLINTNEXTLINE(bugprone-assignment-in-if-condition)
|
||||
if ((fd = open(filename_.c_str(), flags, (mode_t)0600)) == -1) {
|
||||
if ((fd = open(filename_.c_str(), flags, static_cast<mode_t>(0600))) == -1) {
|
||||
TORCH_CHECK(false, "unable to open file <", filename_, "> in read-write mode: ", c10::utils::str_error(errno), " (", errno, ")");
|
||||
}
|
||||
} else if (flags_ & ALLOCATOR_MAPPED_SHAREDMEM) {
|
||||
#ifdef HAVE_SHM_OPEN
|
||||
// NOLINTNEXTLINE(bugprone-assignment-in-if-condition)
|
||||
if((fd = shm_open(filename_.c_str(), flags, (mode_t)0600)) == -1) {
|
||||
if((fd = shm_open(filename_.c_str(), flags, static_cast<mode_t>(0600))) == -1) {
|
||||
TORCH_CHECK(false, "unable to open shared memory object <", filename_, "> in read-write mode: ", c10::utils::str_error(errno), " (", errno, ")");
|
||||
}
|
||||
#else
|
||||
@ -503,7 +503,7 @@ RefcountedMapAllocator::RefcountedMapAllocator(WithFd /*unused*/, const char *fi
|
||||
|
||||
void RefcountedMapAllocator::initializeAlloc() {
|
||||
TORCH_CHECK(base_ptr_, "base_ptr_ is null");
|
||||
MapInfo *map_info = (MapInfo*)base_ptr_;
|
||||
MapInfo *map_info = static_cast<MapInfo*>(base_ptr_);
|
||||
|
||||
#ifdef _WIN32
|
||||
ReleaseContext* r_ctx = new ReleaseContext;
|
||||
@ -539,7 +539,7 @@ void RefcountedMapAllocator::close() {
|
||||
}
|
||||
#else /* _WIN32 */
|
||||
|
||||
MapInfo *info = (MapInfo*)(data);
|
||||
MapInfo *info = static_cast<MapInfo*>(data);
|
||||
if (--info->refcount == 0) {
|
||||
#ifdef HAVE_SHM_UNLINK
|
||||
if (shm_unlink(filename_.c_str()) == -1) {
|
||||
|
||||
@ -862,7 +862,7 @@ void TensorIteratorBase::narrow(int dim, int64_t start, int64_t size) {
|
||||
shape_[dim] = size;
|
||||
view_offsets_[dim] += start;
|
||||
for (auto& op : operands_) {
|
||||
op.data = ((char*)op.data) + op.stride_bytes[dim] * start;
|
||||
op.data = (static_cast<char*>(op.data)) + op.stride_bytes[dim] * start;
|
||||
}
|
||||
if (size == 1 && !is_reduction_) {
|
||||
coalesce_dimensions();
|
||||
@ -873,7 +873,7 @@ void TensorIteratorBase::select_all_keeping_dim(int start_dim, IntArrayRef indic
|
||||
TORCH_INTERNAL_ASSERT(start_dim <= ndim());
|
||||
for (const auto i : c10::irange(start_dim, ndim())) {
|
||||
for (auto& op : operands_) {
|
||||
op.data = ((char*)op.data) + op.stride_bytes[i] * indices[i - start_dim];
|
||||
op.data = (static_cast<char*>(op.data)) + op.stride_bytes[i] * indices[i - start_dim];
|
||||
}
|
||||
shape_[i] = 1;
|
||||
}
|
||||
|
||||
@ -41,7 +41,7 @@ inline void serial_for_each(
|
||||
IntArrayRef strides,
|
||||
char** base_ptrs,
|
||||
size_t ntensors,
|
||||
typename TensorIteratorBase::loop2d_t loop,
|
||||
TensorIteratorBase::loop2d_t loop,
|
||||
Range range) {
|
||||
const auto ndim = shape.size();
|
||||
TORCH_INTERNAL_ASSERT_DEBUG_ONLY(
|
||||
|
||||
@ -190,12 +190,14 @@ class IListRef;
|
||||
* it to a function (e.g. `ImplT::<dispatch-function>(this_)`).
|
||||
*/
|
||||
#define TORCH_ILISTREF_UNWRAP(TAG, BODY) \
|
||||
C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum") \
|
||||
switch (TAG) { \
|
||||
TORCH_ILISTREF_FORALL_TAGS(TORCH_ILISTREF_UNWRAP_CASE, BODY) \
|
||||
break; \
|
||||
default: \
|
||||
TORCH_INTERNAL_ASSERT(false, "invalid IListRef tag."); \
|
||||
}
|
||||
} \
|
||||
C10_DIAGNOSTIC_POP()
|
||||
|
||||
enum class IListRefTag {
|
||||
#define DEFINE_TAG(tag, ...) tag,
|
||||
|
||||
@ -56,7 +56,7 @@ C10_HOST_DEVICE inline T uniform_int_full_range(V val) {
|
||||
* in this overloaded version
|
||||
*/
|
||||
template <typename T, typename V>
|
||||
C10_HOST_DEVICE inline std::enable_if_t<!(std::is_floating_point_v<T>), T>uniform_int(V val) {
|
||||
C10_HOST_DEVICE inline std::enable_if_t<!std::is_floating_point_v<T>, T>uniform_int(V val) {
|
||||
if constexpr (std::is_same_v<T, bool>) {
|
||||
return static_cast<bool>(val & 1);
|
||||
} else if constexpr (std::is_same_v<T, int64_t>) {
|
||||
|
||||
@ -114,25 +114,25 @@ inline typename remove_symint<T>::type unpackSymInt(T x) {
|
||||
}
|
||||
|
||||
template <>
|
||||
inline typename remove_symint<c10::SymInt>::type unpackSymInt(c10::SymInt x) {
|
||||
inline remove_symint<c10::SymInt>::type unpackSymInt(c10::SymInt x) {
|
||||
return x.guard_int(__FILE__, __LINE__);
|
||||
}
|
||||
|
||||
template <>
|
||||
inline typename remove_symint<c10::SymIntArrayRef>::type unpackSymInt(
|
||||
inline remove_symint<c10::SymIntArrayRef>::type unpackSymInt(
|
||||
c10::SymIntArrayRef x) {
|
||||
return C10_AS_INTARRAYREF_SLOW(x);
|
||||
}
|
||||
|
||||
template <>
|
||||
inline typename remove_symint<std::optional<c10::SymInt>>::type unpackSymInt(
|
||||
inline remove_symint<std::optional<c10::SymInt>>::type unpackSymInt(
|
||||
std::optional<c10::SymInt> x) {
|
||||
return x.has_value() ? std::make_optional(x->guard_int(__FILE__, __LINE__))
|
||||
: std::nullopt;
|
||||
}
|
||||
|
||||
template <>
|
||||
inline typename remove_symint<at::OptionalSymIntArrayRef>::type unpackSymInt(
|
||||
inline remove_symint<at::OptionalSymIntArrayRef>::type unpackSymInt(
|
||||
at::OptionalSymIntArrayRef x) {
|
||||
return x.has_value() ? std::make_optional(C10_AS_INTARRAYREF_SLOW(*x))
|
||||
: std::nullopt;
|
||||
|
||||
@ -631,8 +631,8 @@ call_functor_with_args_from_stack_(
|
||||
Stack* stack,
|
||||
std::index_sequence<ivalue_arg_indices...> /*unused*/,
|
||||
guts::typelist::typelist<ArgTypes...>* /*unused*/) {
|
||||
(void)(stack); // when sizeof...(ivalue_arg_indices) == 0, this argument would
|
||||
// be unused and we have to silence the compiler warning.
|
||||
(void)stack; // when sizeof...(ivalue_arg_indices) == 0, this argument would
|
||||
// be unused and we have to silence the compiler warning.
|
||||
|
||||
// We're explicitly filtering out DispatchKeySet from the argument list.
|
||||
// Some kernels take a DispatchKeySet as their first argument in order to
|
||||
|
||||
@ -18,6 +18,7 @@ struct TORCH_API EnumType : public NamedType {
|
||||
TypePtr value,
|
||||
std::vector<EnumNameValue> enum_names_values,
|
||||
std::weak_ptr<::torch::jit::CompilationUnit> cu) {
|
||||
C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum")
|
||||
switch (value->kind()) {
|
||||
case TypeKind::IntType:
|
||||
case TypeKind::FloatType:
|
||||
@ -34,6 +35,7 @@ struct TORCH_API EnumType : public NamedType {
|
||||
value->str(),
|
||||
"', only int, float and string are supported");
|
||||
}
|
||||
C10_DIAGNOSTIC_POP()
|
||||
}
|
||||
|
||||
std::string str() const override {
|
||||
|
||||
@ -601,8 +601,8 @@ std::ostream& IValue::repr(
|
||||
double d = v.toDouble();
|
||||
int c = std::fpclassify(d);
|
||||
if ((c == FP_NORMAL || c == FP_ZERO ) && std::abs(d) < 1e10) {
|
||||
int64_t i = int64_t(d);
|
||||
if (double(i) == d) {
|
||||
int64_t i = static_cast<int64_t>(d);
|
||||
if (static_cast<double>(i) == d) {
|
||||
// -0.0 (signed zero) needs to be parsed as -0.
|
||||
if (i == 0 && std::signbit(d)) {
|
||||
return out << "-" << i << ".";
|
||||
@ -799,8 +799,8 @@ std::ostream& operator<<(std::ostream & out, const IValue & v) {
|
||||
double d = v.toDouble();
|
||||
int c = std::fpclassify(d);
|
||||
if (c == FP_NORMAL || c == FP_ZERO) {
|
||||
int64_t i = int64_t(d);
|
||||
if (double(i) == d) {
|
||||
int64_t i = static_cast<int64_t>(d);
|
||||
if (static_cast<double>(i) == d) {
|
||||
return out << i << ".";
|
||||
}
|
||||
}
|
||||
|
||||
@ -41,7 +41,7 @@ void standardizeVectorForUnion(std::vector<TypePtr>* to_flatten);
|
||||
inline bool is_contiguous_strides(
|
||||
const IntArrayRef sizes,
|
||||
const IntArrayRef strides) {
|
||||
int n_dim = static_cast<int>(sizes.size());
|
||||
size_t n_dim = sizes.size();
|
||||
if (n_dim == 0) {
|
||||
return true;
|
||||
}
|
||||
@ -50,7 +50,7 @@ inline bool is_contiguous_strides(
|
||||
return false;
|
||||
}
|
||||
|
||||
for (int i = n_dim - 2; i >= 0; i--) {
|
||||
for (int i = static_cast<int>(n_dim) - 2; i >= 0; i--) {
|
||||
if (strides[i] != strides[i + 1] * sizes[i + 1]) {
|
||||
return false;
|
||||
}
|
||||
@ -922,6 +922,7 @@ struct TORCH_API DictType : public SharedType {
|
||||
if (auto dyn = key->castRaw<DynamicType>()) {
|
||||
kind = dyn->dynamicKind();
|
||||
}
|
||||
C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum")
|
||||
switch (kind) {
|
||||
case TypeKind::AnyType:
|
||||
case TypeKind::IntType:
|
||||
@ -938,6 +939,7 @@ struct TORCH_API DictType : public SharedType {
|
||||
key->str(),
|
||||
"', only int, float, complex, Tensor, device and string keys are supported");
|
||||
}
|
||||
C10_DIAGNOSTIC_POP()
|
||||
}
|
||||
|
||||
// aligned with the format in FunctionSchema
|
||||
@ -2371,7 +2373,7 @@ private:
|
||||
};
|
||||
|
||||
template<>
|
||||
inline typename detail::CastReturnType<NamedType>::type Type::cast<NamedType>() {
|
||||
inline detail::CastReturnType<NamedType>::type Type::cast<NamedType>() {
|
||||
if (kind() == TypeKind::TupleType || kind() == TypeKind::FunctionType ||
|
||||
kind() == TypeKind::ClassType || kind() == TypeKind::InterfaceType) {
|
||||
return std::static_pointer_cast<NamedType>(static_cast<NamedType *>(this)->shared_from_this());
|
||||
@ -2380,7 +2382,7 @@ inline typename detail::CastReturnType<NamedType>::type Type::cast<NamedType>()
|
||||
}
|
||||
|
||||
template<>
|
||||
inline typename detail::CastConstReturnType<NamedType>::type Type::cast<NamedType>() const {
|
||||
inline detail::CastConstReturnType<NamedType>::type Type::cast<NamedType>() const {
|
||||
if (kind() == TypeKind::TupleType || kind() == TypeKind::FunctionType ||
|
||||
kind() == TypeKind::ClassType || kind() == TypeKind::InterfaceType) {
|
||||
return std::static_pointer_cast<const NamedType>(static_cast<const NamedType *>(this)->shared_from_this());
|
||||
|
||||
@ -514,7 +514,7 @@ struct Vectorized<c10::qint8> : public Vectorizedqi {
|
||||
|
||||
using float_vec_return_type = std::array<Vectorized<float>, kFloatNumVecs>;
|
||||
using int_vec_return_type = std::array<Vectorized<c10::qint32>, kIntNumVecs>;
|
||||
using value_type = typename c10::qint8::underlying;
|
||||
using value_type = c10::qint8::underlying;
|
||||
|
||||
public:
|
||||
using Vectorizedqi::Vectorizedqi;
|
||||
@ -727,7 +727,7 @@ struct Vectorized<c10::quint8> : public Vectorizedqi {
|
||||
|
||||
using float_vec_return_type = std::array<Vectorized<float>, kFloatNumVecs>;
|
||||
using int_vec_return_type = std::array<Vectorized<c10::qint32>, kIntNumVecs>;
|
||||
using value_type = typename c10::quint8::underlying;
|
||||
using value_type = c10::quint8::underlying;
|
||||
|
||||
public:
|
||||
using Vectorizedqi::Vectorizedqi;
|
||||
|
||||
@ -567,7 +567,7 @@ struct Vectorized<c10::qint8> : public Vectorizedqi {
|
||||
|
||||
using float_vec_return_type = std::array<Vectorized<float>, 4>;
|
||||
using int_vec_return_type = std::array<Vectorized<c10::qint32>, 4>;
|
||||
using value_type = typename c10::qint8::underlying;
|
||||
using value_type = c10::qint8::underlying;
|
||||
|
||||
public:
|
||||
using Vectorizedqi::Vectorizedqi;
|
||||
@ -804,7 +804,7 @@ struct Vectorized<c10::quint8> : public Vectorizedqi {
|
||||
|
||||
using float_vec_return_type = std::array<Vectorized<float>, 4>;
|
||||
using int_vec_return_type = std::array<Vectorized<c10::qint32>, 4>;
|
||||
using value_type = typename c10::quint8::underlying;
|
||||
using value_type = c10::quint8::underlying;
|
||||
|
||||
public:
|
||||
using Vectorizedqi::Vectorizedqi;
|
||||
|
||||
@ -672,7 +672,7 @@ struct Vectorized {
|
||||
return map(std::sqrt);
|
||||
}
|
||||
Vectorized<T> reciprocal() const {
|
||||
return map([](T x) { return (T)(1) / x; });
|
||||
return map([](T x) { return (T)1 / x; });
|
||||
}
|
||||
Vectorized<T> rsqrt() const {
|
||||
return map([](T x) { return (T)1 / std::sqrt(x); });
|
||||
|
||||
@ -46,7 +46,7 @@ inline void vrsqrt(scalar_t* out, scalar_t* in, int64_t size) {
|
||||
parallel_for(0, size, 2048, [out, in](int64_t begin, int64_t end) {
|
||||
map(
|
||||
[](const Vectorized<scalar_t>& x) {
|
||||
return Vectorized<scalar_t>((scalar_t)(1)) / x.sqrt();
|
||||
return Vectorized<scalar_t>((scalar_t)1) / x.sqrt();
|
||||
},
|
||||
out + begin,
|
||||
in + begin,
|
||||
|
||||
@ -194,8 +194,8 @@ void CUDAGeneratorState::unregister_graph(cuda::CUDAGraph* graph) {
|
||||
void CUDAGeneratorState::capture_prologue() {
|
||||
capturing_ = true;
|
||||
offset_intragraph_ = 0;
|
||||
seed_extragraph_.fill_(int64_t(seed_));
|
||||
offset_extragraph_.fill_(int64_t(0));
|
||||
seed_extragraph_.fill_(static_cast<int64_t>(seed_));
|
||||
offset_extragraph_.fill_(0);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -216,8 +216,8 @@ void CUDAGeneratorState::replay_prologue(uint64_t wholegraph_increment) {
|
||||
at::cuda::assertNotCapturing(
|
||||
"Cannot prepare for replay during capturing stage.");
|
||||
if (wholegraph_increment) {
|
||||
seed_extragraph_.fill_(int64_t(seed_));
|
||||
offset_extragraph_.fill_(int64_t(philox_offset_per_thread_));
|
||||
seed_extragraph_.fill_(static_cast<int64_t>(seed_));
|
||||
offset_extragraph_.fill_(static_cast<int64_t>(philox_offset_per_thread_));
|
||||
// Applies the total increment achieved during previous captures to update the
|
||||
// offset.
|
||||
increase(wholegraph_increment);
|
||||
@ -329,7 +329,7 @@ c10::intrusive_ptr<c10::TensorImpl> CUDAGeneratorImpl::get_state() const {
|
||||
constexpr size_t offset_size = sizeof(int64_t);
|
||||
constexpr size_t total_size = seed_size + offset_size;
|
||||
|
||||
auto state_tensor = at::detail::empty_cpu({(int64_t)total_size}, ScalarType::Byte, std::nullopt, std::nullopt, std::nullopt, std::nullopt);
|
||||
auto state_tensor = at::detail::empty_cpu({static_cast<int64_t>(total_size)}, ScalarType::Byte, std::nullopt, std::nullopt, std::nullopt, std::nullopt);
|
||||
auto rng_state = state_tensor.data_ptr<uint8_t>();
|
||||
auto current_seed = this->current_seed();
|
||||
auto offset = static_cast<int64_t>(this->philox_offset_per_thread()); // Note that old THCGeneratorState had offset as std::atomic<int64_t>
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
#include <ATen/cuda/CUDAGreenContext.h>
|
||||
|
||||
#if defined(CUDA_VERSION) && !defined(USE_ROCM) && defined(PYTORCH_C10_DRIVER_API_SUPPORTED)
|
||||
#if defined(CUDA_VERSION) && (CUDA_VERSION >= 12030) && !defined(USE_ROCM) && defined(PYTORCH_C10_DRIVER_API_SUPPORTED)
|
||||
#include <c10/cuda/driver_api.h>
|
||||
#include <stdexcept>
|
||||
#include <vector>
|
||||
|
||||
@ -155,8 +155,8 @@ size_t parseChosenWorkspaceSize() {
|
||||
while (next != end) {
|
||||
std::smatch match = *next;
|
||||
TORCH_CHECK(match.size() == 3, "Expected CUBLAS_WORKSPACE_SPACE_CONFIG match of size 3 (Format :SIZE:COUNT)");
|
||||
size_t curr_size = (size_t) std::stoi(match.str(1));
|
||||
size_t count = (size_t) std::stoi(match.str(2));
|
||||
size_t curr_size = std::stoull(match.str(1));
|
||||
size_t count = std::stoull(match.str(2));
|
||||
total_size += curr_size * 1024 * count;
|
||||
next++;
|
||||
}
|
||||
|
||||
@ -3,6 +3,7 @@
|
||||
#include <ATen/ATen.h>
|
||||
#include <c10/util/irange.h>
|
||||
|
||||
#include <array>
|
||||
#include <iostream>
|
||||
#include <sstream>
|
||||
|
||||
@ -136,9 +137,9 @@ void FilterDescriptor::set(const at::Tensor &t, const at::MemoryFormat memory_fo
|
||||
"Weight strides: ", t.strides(), "\n",
|
||||
"cuDNN suggested memory_format: ", memory_format);
|
||||
|
||||
int size[CUDNN_DIM_MAX];
|
||||
std::array<int, CUDNN_DIM_MAX> size;
|
||||
for (const auto i : c10::irange(dim)) {
|
||||
size[i] = (int) t.size(i);
|
||||
size[i] = static_cast<int>(t.size(i));
|
||||
}
|
||||
for (const auto i : c10::irange(dim, pad)) {
|
||||
size[i] = 1;
|
||||
@ -156,7 +157,7 @@ void FilterDescriptor::set(const at::Tensor &t, const at::MemoryFormat memory_fo
|
||||
default:
|
||||
TORCH_INTERNAL_ASSERT(false, "unsupported memory_format for cuDNN filters");
|
||||
}
|
||||
set(getDataType(t), static_cast<int>(dim), size, filter_format);
|
||||
set(getDataType(t), static_cast<int>(dim), size.data(), filter_format);
|
||||
}
|
||||
|
||||
std::string cudnnMemoryFormatToString(cudnnTensorFormat_t tformat) {
|
||||
|
||||
@ -198,7 +198,7 @@ static void autogradBasedTransformSendToNext(
|
||||
}
|
||||
|
||||
// Step 6
|
||||
stack->erase(stack->end() - std::ptrdiff_t(args_size + ret_size), stack->end() - std::ptrdiff_t(ret_size));
|
||||
stack->erase(stack->end() - static_cast<std::ptrdiff_t>(args_size + ret_size), stack->end() - static_cast<std::ptrdiff_t>(ret_size));
|
||||
}
|
||||
|
||||
void GradInterpreterPtr::processImpl(
|
||||
|
||||
@ -443,14 +443,14 @@ static bool has_same_shape(
|
||||
if (!tensor.defined()) {
|
||||
return true;
|
||||
}
|
||||
if (rankWithoutBatchDim(tensor, tensor_bdim) != (int64_t) normalized_shape.size()) {
|
||||
if (rankWithoutBatchDim(tensor, tensor_bdim) != static_cast<int64_t>(normalized_shape.size())) {
|
||||
return false;
|
||||
}
|
||||
const auto tensor_shape = tensor.sizes();
|
||||
for (const auto i : c10::irange(normalized_shape.size())) {
|
||||
auto j = i;
|
||||
// (0, 1, 2), 1 -> (0, 2, 3)
|
||||
if (tensor_bdim.has_value() && (int64_t)i >= tensor_bdim.value()) {
|
||||
if (tensor_bdim.has_value() && static_cast<int64_t>(i) >= tensor_bdim.value()) {
|
||||
j = j + 1;
|
||||
}
|
||||
if (normalized_shape[i] != tensor_shape[j]) {
|
||||
|
||||
@ -135,7 +135,7 @@ static void boxed_reduction_batch_rule(const c10::OperatorHandle& op, torch::jit
|
||||
reduction_case = ReductionCase::DimArray;
|
||||
dims = arguments[dim_arg_pos].toIntList().vec();
|
||||
if (dims.empty()) {
|
||||
auto all_dims = range(0, std::max((int64_t)1, logical_dim));
|
||||
auto all_dims = range(0, std::max(static_cast<int64_t>(1), logical_dim));
|
||||
dims = std::vector<int64_t>(all_dims.begin(), all_dims.end());
|
||||
}
|
||||
} else if (arguments[dim_arg_pos].isInt()) {
|
||||
|
||||
@ -432,7 +432,7 @@ namespace {
|
||||
// Eg. Given `indexed_shape.size()` is 5 and
|
||||
// shape of `values` is (N, 2, 3), then following block
|
||||
// will reshape `values` to (N, 1, 1, 2, 3).
|
||||
if ( (int64_t) indexed_shape.size() > values_.dim()) {
|
||||
if ( static_cast<int64_t>(indexed_shape.size()) > values_.dim()) {
|
||||
auto values_sizes = values_.sym_sizes();
|
||||
|
||||
// number of unit dims (for broadcasting value to indexed_shape)
|
||||
|
||||
@ -109,7 +109,7 @@ std::tuple<Tensor, std::optional<int64_t>> repeat_batch_rule(
|
||||
SymDimVector sizes_with_bdim = { sizes.begin(), sizes.end() };
|
||||
sizes_with_bdim.insert(sizes_with_bdim.begin(), 1);
|
||||
auto self_ = moveBatchDimToFront(self, self_bdim);
|
||||
while (self_.dim() < (int64_t)sizes_with_bdim.size()) {
|
||||
while (self_.dim() < static_cast<int64_t>(sizes_with_bdim.size())) {
|
||||
self_ = self_.unsqueeze(1);
|
||||
}
|
||||
return std::make_tuple(self_.repeat_symint(sizes_with_bdim), 0);
|
||||
|
||||
@ -191,7 +191,7 @@ static void batchedTensorInplaceForLoopFallback(const c10::OperatorHandle& op, t
|
||||
// simplicity. When that is not the case, this code should be updated.
|
||||
const auto& argument = (*stack)[arguments_begin + arg_idx];
|
||||
if (batched_tensor_inputs_pos_iter == batched_tensor_inputs_position.end()
|
||||
|| (int64_t)arg_idx != *batched_tensor_inputs_pos_iter) {
|
||||
|| static_cast<int64_t>(arg_idx) != *batched_tensor_inputs_pos_iter) {
|
||||
// argument isn't a BatchedTensor
|
||||
torch::jit::push(stack, argument);
|
||||
continue;
|
||||
@ -345,7 +345,7 @@ void batchedTensorForLoopFallback(const c10::OperatorHandle& op, torch::jit::Sta
|
||||
// simplicity. When that is not the case, this code should be updated.
|
||||
const auto& argument = (*stack)[arguments_begin + arg_idx];
|
||||
if (batched_tensor_inputs_pos_iter == batched_tensor_inputs_position.end()
|
||||
|| (int64_t)arg_idx != *batched_tensor_inputs_pos_iter) {
|
||||
|| static_cast<int64_t>(arg_idx) != *batched_tensor_inputs_pos_iter) {
|
||||
// argument isn't a BatchedTensor
|
||||
torch::jit::push(stack, argument);
|
||||
continue;
|
||||
@ -473,7 +473,7 @@ void batchedNestedTensorForLoopFallback(const c10::OperatorHandle& op, torch::ji
|
||||
// simplicity. When that is not the case, this code should be updated.
|
||||
const auto& argument = (*stack)[arguments_begin + arg_idx];
|
||||
if (batched_tensor_inputs_pos_iter == batched_tensor_inputs_position.end()
|
||||
|| (int64_t)arg_idx != *batched_tensor_inputs_pos_iter) {
|
||||
|| static_cast<int64_t>(arg_idx) != *batched_tensor_inputs_pos_iter) {
|
||||
// argument isn't a BatchedTensor
|
||||
torch::jit::push(stack, argument);
|
||||
continue;
|
||||
|
||||
@ -157,7 +157,7 @@ Tensor& squeeze__batching_rule(Tensor& self) {
|
||||
const auto physical_shape = batched->value().sizes();
|
||||
auto how_many_dims_of_size_1_before_bdim = 0;
|
||||
for (const auto i : c10::irange(0, physical_shape.size())) {
|
||||
if ((int64_t)i == bdim) {
|
||||
if (static_cast<int64_t>(i) == bdim) {
|
||||
break;
|
||||
}
|
||||
if (physical_shape[i] == 1) {
|
||||
@ -573,7 +573,7 @@ Tensor cat_batching_rule(const ITensorListRef& tensors, int64_t dim) {
|
||||
}
|
||||
|
||||
auto new_dim = bdim_size.has_value() ? dim + 1 : dim;
|
||||
std::optional<int64_t> new_bdim = bdim_size.has_value() ? std::make_optional((int64_t)0) : std::nullopt;
|
||||
std::optional<int64_t> new_bdim = bdim_size.has_value() ? std::make_optional(static_cast<int64_t>(0)) : std::nullopt;
|
||||
auto result = at::cat(tensors_to_cat, new_dim);
|
||||
return makeBatched(result, new_bdim, get_current_level());
|
||||
}
|
||||
|
||||
@ -198,9 +198,9 @@ void avg_pool3d_out_frame(
|
||||
int64_t hend = std::min(hstart + kH, iheight + padH);
|
||||
int64_t wend = std::min(wstart + kW, iwidth + padW);
|
||||
int64_t pool_size = (tend - tstart) * (hend - hstart) * (wend - wstart);
|
||||
tstart = std::max(tstart, (int64_t) 0);
|
||||
hstart = std::max(hstart, (int64_t) 0);
|
||||
wstart = std::max(wstart, (int64_t) 0);
|
||||
tstart = std::max(tstart, static_cast<int64_t>(0));
|
||||
hstart = std::max(hstart, static_cast<int64_t>(0));
|
||||
wstart = std::max(wstart, static_cast<int64_t>(0));
|
||||
tend = std::min(tend, itime);
|
||||
hend = std::min(hend, iheight);
|
||||
wend = std::min(wend, iwidth);
|
||||
@ -377,9 +377,9 @@ void avg_pool3d_backward_out_frame(
|
||||
int64_t hend = std::min(hstart + kH, iheight + padH);
|
||||
int64_t wend = std::min(wstart + kW, iwidth + padW);
|
||||
int64_t pool_size = (tend -tstart) * (hend - hstart) * (wend - wstart);
|
||||
tstart = std::max(tstart, (int64_t) 0);
|
||||
hstart = std::max(hstart, (int64_t) 0);
|
||||
wstart = std::max(wstart, (int64_t) 0);
|
||||
tstart = std::max(tstart, static_cast<int64_t>(0));
|
||||
hstart = std::max(hstart, static_cast<int64_t>(0));
|
||||
wstart = std::max(wstart, static_cast<int64_t>(0));
|
||||
tend = std::min(tend, itime);
|
||||
hend = std::min(hend, iheight);
|
||||
wend = std::min(wend, iwidth);
|
||||
|
||||
@ -946,10 +946,10 @@ void apply_lu_factor(const Tensor& input, const Tensor& pivots, const Tensor& in
|
||||
}
|
||||
};
|
||||
// avoid overflow
|
||||
float matrix_rank = float(std::min(m, n));
|
||||
auto matrix_rank = std::min(m, n);
|
||||
// A heuristic tested on a 32 core/socket ICX system
|
||||
// https://github.com/pytorch/pytorch/pull/93037#discussion_r1090112948
|
||||
int64_t chunk_size_per_thread = int64_t(
|
||||
int64_t chunk_size_per_thread = static_cast<int64_t>(
|
||||
std::min(1.0, 3200.0 / (matrix_rank * matrix_rank * matrix_rank)));
|
||||
int64_t grain_size = chunk_size_per_thread * at::get_num_threads();
|
||||
at::parallel_for(0, batch_size, grain_size, loop);
|
||||
|
||||
@ -267,7 +267,7 @@ _scaled_mm_out_cpu_emulated(const Tensor& mat1, const Tensor& mat2,
|
||||
|
||||
float input_scale = scale_a.item<float>();
|
||||
float weight_scale = scale_b.item<float>();
|
||||
float output_scale = float(1.0);
|
||||
float output_scale = 1.0f;
|
||||
if (scale_result.has_value() &&
|
||||
(*out_dtype == ScalarType::Float8_e4m3fn ||
|
||||
*out_dtype == ScalarType::Float8_e5m2)) {
|
||||
|
||||
@ -331,7 +331,7 @@ bool gemv_use_fast_path<double>(
|
||||
[[maybe_unused]] double beta,
|
||||
int64_t incy) {
|
||||
return gemv_use_fast_path<float>(
|
||||
trans, m, n, (float)alpha, lda, incx, (float)beta, incy);
|
||||
trans, m, n, static_cast<float>(alpha), lda, incx, static_cast<float>(beta), incy);
|
||||
}
|
||||
|
||||
template <>
|
||||
@ -523,8 +523,8 @@ static inline void scal(int64_t n, scalar_t a, scalar_t *x, int64_t incx)
|
||||
if (n == 1) incx = 1;
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if (blas_impl::scal_use_fast_path<scalar_t>(n, incx)) {
|
||||
int i_n = (int)n;
|
||||
int i_incx = (int)incx;
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
blas_impl::scal_fast_path<scalar_t>(&i_n, &a, x, &i_incx);
|
||||
return;
|
||||
}
|
||||
@ -545,11 +545,11 @@ void gemv(char trans, int64_t m, int64_t n, scalar_t alpha, const scalar_t *a, i
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if (blas_impl::gemv_use_fast_path<scalar_t>(trans, m, n, alpha, lda, incx, beta, incy)) {
|
||||
TORCH_CHECK(lda >= std::max<int64_t>(1L, m), "lda should be at least max(1,", m, "), but have ", lda);
|
||||
int i_m = (int)m;
|
||||
int i_n = (int)n;
|
||||
int i_lda = (int)lda;
|
||||
int i_incx = (int)incx;
|
||||
int i_incy = (int)incy;
|
||||
int i_m = static_cast<int>(m);
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_lda = static_cast<int>(lda);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
int i_incy = static_cast<int>(incy);
|
||||
blas_impl::gemv_fast_path<scalar_t>(&trans, &i_m, &i_n, &alpha, a, &i_lda, x, &i_incx, &beta, y, &i_incy);
|
||||
return;
|
||||
}
|
||||
|
||||
@ -680,9 +680,9 @@ void axpy(int64_t n, double a, const double *x, int64_t incx, double *y, int64_t
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if( (n <= INT_MAX) && (incx <= INT_MAX) && (incy <= INT_MAX) )
|
||||
{
|
||||
int i_n = (int)n;
|
||||
int i_incx = (int)incx;
|
||||
int i_incy = (int)incy;
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
int i_incy = static_cast<int>(incy);
|
||||
#if C10_IOS
|
||||
cblas_daxpy(i_n, a, x, i_incx, y, i_incy);
|
||||
#else
|
||||
@ -705,9 +705,9 @@ void axpy(int64_t n, float a, const float *x, int64_t incx, float *y, int64_t in
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if( (n <= INT_MAX) && (incx <= INT_MAX) && (incy <= INT_MAX) )
|
||||
{
|
||||
int i_n = (int)n;
|
||||
int i_incx = (int)incx;
|
||||
int i_incy = (int)incy;
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
int i_incy = static_cast<int>(incy);
|
||||
#if C10_IOS
|
||||
cblas_saxpy(i_n, a, x, i_incx, y, i_incy);
|
||||
#else
|
||||
@ -730,9 +730,9 @@ void axpy(int64_t n, c10::complex<double> a, const c10::complex<double> *x, int6
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if( (n <= INT_MAX) && (incx <= INT_MAX) && (incy <= INT_MAX) )
|
||||
{
|
||||
int i_n = (int)n;
|
||||
int i_incx = (int)incx;
|
||||
int i_incy = (int)incy;
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
int i_incy = static_cast<int>(incy);
|
||||
#if C10_IOS
|
||||
cblas_zaxpy(i_n, &a, x, i_incx, y, i_incy);
|
||||
#else
|
||||
@ -755,9 +755,9 @@ void axpy(int64_t n, c10::complex<float> a, const c10::complex<float> *x, int64_
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if( (n <= INT_MAX) && (incx <= INT_MAX) && (incy <= INT_MAX) )
|
||||
{
|
||||
int i_n = (int)n;
|
||||
int i_incx = (int)incx;
|
||||
int i_incy = (int)incy;
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
int i_incy = static_cast<int>(incy);
|
||||
#if C10_IOS
|
||||
cblas_caxpy(i_n, &a, x, i_incx, y, i_incy);
|
||||
#else
|
||||
@ -781,9 +781,9 @@ void copy(int64_t n, const double *x, int64_t incx, double *y, int64_t incy) {
|
||||
}
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if( (n <= INT_MAX) && (incx <= INT_MAX) && (incy <= INT_MAX) ) {
|
||||
int i_n = (int)n;
|
||||
int i_incx = (int)incx;
|
||||
int i_incy = (int)incy;
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
int i_incy = static_cast<int>(incy);
|
||||
#if C10_IOS
|
||||
cblas_dcopy(i_n, x, i_incx, y, i_incy);
|
||||
#else
|
||||
@ -805,9 +805,9 @@ void copy(int64_t n, const float *x, int64_t incx, float *y, int64_t incy) {
|
||||
}
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if( (n <= INT_MAX) && (incx <= INT_MAX) && (incy <= INT_MAX) ) {
|
||||
int i_n = (int)n;
|
||||
int i_incx = (int)incx;
|
||||
int i_incy = (int)incy;
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
int i_incy = static_cast<int>(incy);
|
||||
#if C10_IOS
|
||||
cblas_scopy(i_n, x, i_incx, y, i_incy);
|
||||
#else
|
||||
@ -829,9 +829,9 @@ void copy(int64_t n, const c10::complex<double> *x, int64_t incx, c10::complex<d
|
||||
}
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if( (n <= INT_MAX) && (incx <= INT_MAX) && (incy <= INT_MAX) ) {
|
||||
int i_n = (int)n;
|
||||
int i_incx = (int)incx;
|
||||
int i_incy = (int)incy;
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
int i_incy = static_cast<int>(incy);
|
||||
#if C10_IOS
|
||||
cblas_zcopy(i_n, x, i_incx, y, i_incy);
|
||||
#else
|
||||
@ -853,9 +853,9 @@ void copy(int64_t n, const c10::complex<float> *x, int64_t incx, c10::complex<fl
|
||||
}
|
||||
#if AT_BUILD_WITH_BLAS()
|
||||
if( (n <= INT_MAX) && (incx <= INT_MAX) && (incy <= INT_MAX) ) {
|
||||
int i_n = (int)n;
|
||||
int i_incx = (int)incx;
|
||||
int i_incy = (int)incy;
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_incx = static_cast<int>(incx);
|
||||
int i_incy = static_cast<int>(incy);
|
||||
#if C10_IOS
|
||||
cblas_ccopy(i_n, &x, i_incx, y, i_incy);
|
||||
#else
|
||||
@ -1082,7 +1082,7 @@ struct Brgemm : public KernelCache <BrgemmKey, GemmHelper> {
|
||||
M,
|
||||
N,
|
||||
K,
|
||||
int64_t(1),
|
||||
1,
|
||||
ld_a,
|
||||
ld_b,
|
||||
ld_c,
|
||||
@ -1096,7 +1096,7 @@ struct Brgemm : public KernelCache <BrgemmKey, GemmHelper> {
|
||||
M,
|
||||
N,
|
||||
K,
|
||||
int64_t(1),
|
||||
1,
|
||||
ld_a,
|
||||
ld_b,
|
||||
ld_c,
|
||||
|
||||
@ -487,17 +487,17 @@ static Tensor _grid_sampler_2d_cpu_quantized(
|
||||
int64_t out_sC = output.stride(1);
|
||||
int64_t out_sH = output.stride(2);
|
||||
int64_t out_sW = output.stride(3);
|
||||
uint8_t* inp_ptr = (uint8_t*)input.data_ptr<quint8>();
|
||||
uint8_t* out_ptr = (uint8_t*)output.data_ptr<quint8>();
|
||||
float* grid_ptr = grid.data_ptr<float>();
|
||||
const uint8_t* inp_ptr = input.const_data_ptr<uint8_t>();
|
||||
uint8_t* out_ptr = output.data_ptr<uint8_t>();
|
||||
const float* grid_ptr = grid.const_data_ptr<float>();
|
||||
at::parallel_for(0, N, 0, [&](int64_t start, int64_t end) {
|
||||
for (const auto n : c10::irange(start, end)) {
|
||||
float* grid_ptr_N = grid_ptr + n * grid_sN;
|
||||
uint8_t* inp_ptr_N = inp_ptr + n * inp_sN;
|
||||
const float* grid_ptr_N = grid_ptr + n * grid_sN;
|
||||
const uint8_t* inp_ptr_N = inp_ptr + n * inp_sN;
|
||||
for (const auto h : c10::irange(out_H)) {
|
||||
for (const auto w : c10::irange(out_W)) {
|
||||
// get the corresponding input x, y, z coordinates from grid
|
||||
float* grid_ptr_NHW = grid_ptr_N + h * grid_sH + w * grid_sW;
|
||||
const float* grid_ptr_NHW = grid_ptr_N + h * grid_sH + w * grid_sW;
|
||||
float x = *grid_ptr_NHW;
|
||||
float y = grid_ptr_NHW[grid_sCoor];
|
||||
|
||||
@ -527,7 +527,7 @@ static Tensor _grid_sampler_2d_cpu_quantized(
|
||||
float se = (ix - ix_nw) * (iy - iy_nw);
|
||||
|
||||
// calculate bilinear weighted pixel value and set output pixel
|
||||
uint8_t* inp_ptr_NC = inp_ptr_N;
|
||||
const uint8_t* inp_ptr_NC = inp_ptr_N;
|
||||
uint8_t* out_ptr_NCHW =
|
||||
out_ptr + n * out_sN + h * out_sH + w * out_sW;
|
||||
for (int64_t c = 0; c < C;
|
||||
|
||||
@ -318,7 +318,7 @@ static std::vector<Tensor>& histogramdd_bin_edges_out(const Tensor& self, IntArr
|
||||
|
||||
const int64_t N = self.size(-1);
|
||||
const int64_t M = std::accumulate(self.sizes().begin(), self.sizes().end() - 1,
|
||||
(int64_t)1, std::multiplies<int64_t>());
|
||||
static_cast<int64_t>(1), std::multiplies<int64_t>());
|
||||
Tensor reshaped_self = self.reshape({ M, N });
|
||||
|
||||
auto outer_bin_edges = select_outer_bin_edges(reshaped_self, range);
|
||||
|
||||
@ -40,7 +40,7 @@ Tensor do_trapezoid(const Tensor& y, const Tensor& dx, int64_t dim) {
|
||||
// When dx is constant, the above formula simplifies
|
||||
// to dx * [(\sum_{i=1}^n y_i) - (y_1 + y_n)/2]
|
||||
Tensor do_trapezoid(const Tensor& y, double dx, int64_t dim) {
|
||||
return (y.sum(dim) - (y.select(dim, 0) + y.select(dim, -1)) * (0.5)) * dx;
|
||||
return (y.sum(dim) - (y.select(dim, 0) + y.select(dim, -1)) * 0.5) * dx;
|
||||
}
|
||||
|
||||
Tensor zeros_like_except(const Tensor& y, int64_t dim) {
|
||||
|
||||
@ -201,7 +201,7 @@ static Tensor sumproduct_pair(const Tensor& left_, const Tensor& right_, IntArra
|
||||
out_size.reserve(out_num_dim);
|
||||
for (auto& d : lro) out_size.push_back(left.sym_size(d));
|
||||
for (auto& d : lo) out_size.push_back(left.sym_size(d));
|
||||
for (auto& d : sum_dims_) { out_size.emplace_back(1); (void)(d); }; // avoid warning about not using d
|
||||
for (auto& d : sum_dims_) { out_size.emplace_back(1); (void)d; }; // avoid warning about not using d
|
||||
for (auto& d : ro) out_size.push_back(right.sym_size(d));
|
||||
|
||||
std::vector<int64_t> lpermutation(lro);
|
||||
@ -640,7 +640,7 @@ Tensor einsum(std::string_view equation, TensorList operands, at::OptionalIntArr
|
||||
}
|
||||
}
|
||||
|
||||
return ops[0];
|
||||
return std::move(ops[0]);
|
||||
}
|
||||
|
||||
// _trilinear computes a trilinear einstein sum with an unrolled dimension
|
||||
@ -805,7 +805,7 @@ Tensor tensordot(const Tensor& input1, const Tensor& input2, IntArrayRef dims1,
|
||||
std::vector<SymInt> rsizes; // rsizes: sizes of the result
|
||||
p1.reserve(input1.dim());
|
||||
p2.reserve(input2.dim());
|
||||
rsizes.reserve(input1.dim() + input2.dim() - (int64_t) dims1.size());
|
||||
rsizes.reserve(input1.dim() + input2.dim() - static_cast<int64_t>(dims1.size()));
|
||||
SymInt size1 = 1; // number of non-contracted elements in input1
|
||||
SymInt size2 = 1; // number of non-contracted elements in input2
|
||||
|
||||
|
||||
@ -1655,7 +1655,7 @@ static inline void baddbmm_cpu_kernel(const Tensor& result, const Tensor& self,
|
||||
auto s0 = self.accessor<const scalar_t, 3>();
|
||||
auto m0 = mat2.accessor<const scalar_t, 3>();
|
||||
|
||||
int64_t grain_size = std::max(internal::GRAIN_SIZE / (is * js * ks), (int64_t)1);
|
||||
int64_t grain_size = std::max(internal::GRAIN_SIZE / (is * js * ks), static_cast<int64_t>(1));
|
||||
using opmath_t = at::opmath_type<scalar_t>;
|
||||
parallel_for(0, bs, grain_size, [&](int64_t b_begin, int64_t b_end) {
|
||||
for (const auto b : c10::irange(b_begin, b_end)) {
|
||||
|
||||
@ -235,7 +235,7 @@ void nll_loss_out_frame(
|
||||
|
||||
constexpr int64_t cascade_sum_num_levels = 8;
|
||||
const int64_t level_power =
|
||||
std::max(int64_t(4), utils::CeilLog2(batch_size) / cascade_sum_num_levels);
|
||||
std::max(static_cast<int64_t>(4), utils::CeilLog2(batch_size) / cascade_sum_num_levels);
|
||||
const int64_t level_step = (1 << level_power);
|
||||
const int64_t level_mask = level_step - 1;
|
||||
|
||||
|
||||
@ -129,7 +129,7 @@ void nll_loss2d_forward_out_frame(
|
||||
for (const auto b : c10::irange(start, end)) {
|
||||
for (const auto h : c10::irange(H)) {
|
||||
for (const auto w : c10::irange(W)) {
|
||||
const int64_t cur_target = (int64_t)target_acc[b][h][w];
|
||||
const int64_t cur_target = target_acc[b][h][w];
|
||||
|
||||
if (cur_target == ignore_index) {
|
||||
output_acc[b][h][w] = static_cast<scalar_t>(0);
|
||||
@ -188,7 +188,7 @@ void nll_loss2d_forward_out_frame(
|
||||
// NOLINTNEXTLINE(cppcoreguidelines-avoid-c-arrays,modernize-avoid-c-arrays)
|
||||
scalar_t loss_partial_sums[cascade_sum_num_levels] = {0};
|
||||
const int64_t level_power =
|
||||
std::max(int64_t(4), utils::CeilLog2(numiter) / cascade_sum_num_levels);
|
||||
std::max(static_cast<int64_t>(4), utils::CeilLog2(numiter) / cascade_sum_num_levels);
|
||||
const int64_t level_step = (1 << level_power);
|
||||
const int64_t level_mask = level_step - 1;
|
||||
|
||||
|
||||
@ -192,7 +192,7 @@ Date: February 1996
|
||||
x = x - (std::erf(x) - y) / ((static_cast<T>(2.0)/static_cast<T>(std::sqrt(c10::pi<double>)))*std::exp(-x*x));
|
||||
x = x - (std::erf(x) - y) / ((static_cast<T>(2.0)/static_cast<T>(std::sqrt(c10::pi<double>)))*std::exp(-x*x));
|
||||
|
||||
return(x);
|
||||
return x;
|
||||
}
|
||||
|
||||
#undef CENTRAL_RANGE
|
||||
@ -3819,7 +3819,7 @@ inline C10_HOST_DEVICE T shifted_chebyshev_polynomial_v_forward(T x, int64_t n)
|
||||
|
||||
if ((n > 6) && (std::abs(x + x - T(1.0)) < T(1.0))) {
|
||||
if (std::sin(std::acos(x + x - T(1.0)) / T(2.0)) != T(1.0)) {
|
||||
return std::cos(((n) + T(0.5)) * std::acos(x + x - T(1.0))) / std::cos(std::acos(x + x - T(1.0)) / T(2.0));
|
||||
return std::cos((n + T(0.5)) * std::acos(x + x - T(1.0))) / std::cos(std::acos(x + x - T(1.0)) / T(2.0));
|
||||
}
|
||||
|
||||
if (n % 2 == 0) {
|
||||
|
||||
@ -193,22 +193,22 @@ Tensor _nnpack_spatial_convolution(
|
||||
const size_t input_channels = input.size(1);
|
||||
const size_t output_channels = weight.size(0);
|
||||
const struct nnp_size input_size = {
|
||||
.width = (size_t)input.size(3),
|
||||
.height = (size_t)input.size(2),
|
||||
.width = static_cast<size_t>(input.size(3)),
|
||||
.height = static_cast<size_t>(input.size(2)),
|
||||
};
|
||||
const struct nnp_padding input_padding = {
|
||||
.top = (size_t)padding[0],
|
||||
.right = (size_t)padding[1],
|
||||
.bottom = (size_t)padding[0],
|
||||
.left = (size_t)padding[1],
|
||||
.top = static_cast<size_t>(padding[0]),
|
||||
.right = static_cast<size_t>(padding[1]),
|
||||
.bottom = static_cast<size_t>(padding[0]),
|
||||
.left = static_cast<size_t>(padding[1]),
|
||||
};
|
||||
const struct nnp_size kernel_size = {
|
||||
.width = (size_t)weight.size(3),
|
||||
.height = (size_t)weight.size(2),
|
||||
.width = static_cast<size_t>(weight.size(3)),
|
||||
.height = static_cast<size_t>(weight.size(2)),
|
||||
};
|
||||
const struct nnp_size output_size = {
|
||||
.width = (size_t)output.size(3),
|
||||
.height = (size_t)output.size(2),
|
||||
.width = static_cast<size_t>(output.size(3)),
|
||||
.height = static_cast<size_t>(output.size(2)),
|
||||
};
|
||||
const nnp_size output_subsample = {
|
||||
.width = static_cast<std::size_t>(stride[1]),
|
||||
|
||||
@ -248,8 +248,8 @@ void slow_conv_transpose3d_out_cpu_template(
|
||||
Tensor weight = weight_.contiguous();
|
||||
Tensor bias = bias_.defined() ? bias_.contiguous() : bias_;
|
||||
|
||||
const int n_input_plane = (int)weight.size(0);
|
||||
const int n_output_plane = (int)weight.size(1);
|
||||
const auto n_input_plane = weight.size(0);
|
||||
const auto n_output_plane = weight.size(1);
|
||||
|
||||
bool is_batch = false;
|
||||
if (input.dim() == 4) {
|
||||
|
||||
@ -84,8 +84,8 @@ static std::vector<int64_t> aligned_size(
|
||||
DimnameList aligned_names,
|
||||
bool is_aligning_two_tensors) {
|
||||
std::vector<int64_t> expanded_sizes(aligned_names.size(), 1);
|
||||
ptrdiff_t dim = (ptrdiff_t)tensor_sizes.size() - 1;
|
||||
ptrdiff_t idx = (ptrdiff_t)aligned_names.size() - 1;
|
||||
ptrdiff_t dim = static_cast<ptrdiff_t>(tensor_sizes.size()) - 1;
|
||||
ptrdiff_t idx = static_cast<ptrdiff_t>(aligned_names.size()) - 1;
|
||||
for (; idx >= 0 && dim >= 0; --idx) {
|
||||
if (tensor_names[dim] != aligned_names[idx]) {
|
||||
continue;
|
||||
|
||||
@ -25,7 +25,7 @@ std::tuple<Tensor, Tensor> _rowwise_prune_helper(
|
||||
auto mask_contig = mask.contiguous();
|
||||
auto mask_data = mask_contig.data_ptr<bool>();
|
||||
for (const auto i : c10::irange(mask.numel())) {
|
||||
num_non_masked_rows += (((mask_data[i] == true)) ? 1 : 0);
|
||||
num_non_masked_rows += ((mask_data[i] == true) ? 1 : 0);
|
||||
}
|
||||
int num_cols = weights.size(1);
|
||||
auto pruned_2d_tensor = at::empty({num_non_masked_rows, num_cols},
|
||||
|
||||
@ -176,7 +176,7 @@ void host_softmax(
|
||||
scalar_t* input_data_base = input.data_ptr<scalar_t>();
|
||||
scalar_t* output_data_base = output.data_ptr<scalar_t>();
|
||||
bool* mask_data_base = mask;
|
||||
int64_t grain_size = std::min(internal::GRAIN_SIZE / dim_size, (int64_t)1);
|
||||
int64_t grain_size = std::min(internal::GRAIN_SIZE / dim_size, static_cast<int64_t>(1));
|
||||
parallel_for(
|
||||
0, outer_size * inner_size, grain_size,
|
||||
[&](int64_t begin, int64_t end) {
|
||||
@ -265,7 +265,7 @@ void host_softmax_backward(
|
||||
scalar_t* output_data_base = output.data_ptr<scalar_t>();
|
||||
scalar_t* gradOutput_data_base = grad.data_ptr<scalar_t>();
|
||||
bool* mask_data_base = mask;
|
||||
int64_t grain_size = std::min(internal::GRAIN_SIZE / dim_size, (int64_t)1);
|
||||
int64_t grain_size = std::min(internal::GRAIN_SIZE / dim_size, static_cast<int64_t>(1));
|
||||
parallel_for(
|
||||
0, outer_size * inner_size, grain_size, [&](int64_t begin, int64_t end) {
|
||||
for (const auto i : c10::irange(begin, end)) {
|
||||
|
||||
@ -1701,13 +1701,13 @@ Tensor& index_select_out_cpu_(
|
||||
TORCH_CHECK_INDEX(
|
||||
(self_i >= 0) && (self_i < self_dim_size),
|
||||
"index out of range in self");
|
||||
auto self_data = static_cast<const char*>(selfSlice_data) +
|
||||
auto self_data = const_cast<char*>(static_cast<const char*>(
|
||||
selfSlice_data)) +
|
||||
self_i * self_stride_bytes;
|
||||
auto result_data = static_cast<char*>(resultSlice_data) +
|
||||
i * result_stride_bytes;
|
||||
sub_iter.unsafe_replace_operand(0, result_data);
|
||||
sub_iter.unsafe_replace_operand(
|
||||
1, const_cast<char*>(self_data));
|
||||
sub_iter.unsafe_replace_operand(1, self_data);
|
||||
copy_stub(sub_iter.device_type(), sub_iter, false);
|
||||
};
|
||||
});
|
||||
|
||||
@ -1382,7 +1382,7 @@ void randperm_cpu(Tensor& result, int64_t n, CPUGeneratorImpl* generator) {
|
||||
// use no-initialization Fischer-Yates variant
|
||||
// https://en.wikipedia.org/wiki/Fisher%E2%80%93Yates_shuffle#The_.22inside-out.22_algorithm
|
||||
for (int64_t i = 0; i < n; i++) {
|
||||
int64_t z = (int64_t)(generator->random64() % (i + 1));
|
||||
int64_t z = static_cast<int64_t>(generator->random64() % (i + 1));
|
||||
r__data[i * r__stride_0] = i;
|
||||
r__data[i * r__stride_0] = r__data[z * r__stride_0];
|
||||
r__data[z * r__stride_0] = i;
|
||||
|
||||
@ -40,7 +40,7 @@ at::Tensor PackedLinearWeightQnnp::apply_dynamic_impl<false>(
|
||||
"quantized_sparse_linear(): Input tensor rank should be >= 2");
|
||||
|
||||
const auto rows_input = c10::multiply_integers(input.sizes().begin(), input.sizes().end() - 1);
|
||||
const auto cols_input = static_cast<int64_t>(input.size(input.dim() - 1));
|
||||
const auto cols_input = input.size(input.dim() - 1);
|
||||
TORCH_CHECK(
|
||||
cols_input == input_channels_,
|
||||
"quantized_sparse_linear: Input tensor's last and weight tensor's"
|
||||
|
||||
@ -65,8 +65,8 @@ LinearPackedSerializationType PackedLinearWeight::unpack() {
|
||||
#ifdef USE_PYTORCH_QNNPACK
|
||||
|
||||
LinearPackedSerializationType PackedLinearWeightQnnp::unpack() {
|
||||
const int64_t N = static_cast<int64_t>(output_channels_);
|
||||
const int64_t K = static_cast<int64_t>(input_channels_);
|
||||
const int64_t N = output_channels_;
|
||||
const int64_t K = input_channels_;
|
||||
|
||||
float* w_scales_ptr = w_scales_.data_ptr<float>();
|
||||
|
||||
|
||||
@ -998,7 +998,7 @@ void softplus_backward_kernel(TensorIteratorBase& iter, const Scalar& beta_, con
|
||||
auto threshold = threshold_.to<float>();
|
||||
const Vec beta_vec(beta);
|
||||
const Vec threshold_vec(threshold);
|
||||
const Vec one_vec(static_cast<float>(1.0));
|
||||
const Vec one_vec(1.0f);
|
||||
cpu_kernel_vec(
|
||||
iter,
|
||||
[beta, threshold](scalar_t a, scalar_t b) -> scalar_t {
|
||||
|
||||
@ -17,7 +17,7 @@ static inline void cpu_atomic_add_float(float* dst, float fvalue)
|
||||
} uf32_t;
|
||||
|
||||
uf32_t new_value, old_value;
|
||||
std::atomic<unsigned>* dst_intV = (std::atomic<unsigned>*)(dst);
|
||||
std::atomic<unsigned>* dst_intV = (std::atomic<unsigned>*)dst;
|
||||
|
||||
old_value.floatV = *dst;
|
||||
new_value.floatV = old_value.floatV + fvalue;
|
||||
|
||||
@ -851,7 +851,7 @@ void sigmoid_backward_kernel(TensorIteratorBase& iter) {
|
||||
});
|
||||
});
|
||||
} else if (iter.dtype() == kBFloat16) {
|
||||
auto one_vec = Vectorized<float>((float)(1));
|
||||
auto one_vec = Vectorized<float>((float)1);
|
||||
cpu_kernel_vec(
|
||||
iter,
|
||||
[=](BFloat16 a, BFloat16 b) -> BFloat16 {
|
||||
|
||||
@ -77,9 +77,7 @@ static void reduced_float_copy_kernel(TensorIteratorBase &iter, bool requires_ne
|
||||
|
||||
int64_t grain_size = at::internal::GRAIN_SIZE;
|
||||
|
||||
auto loop = [strides_in, requires_neg](char** base, const int64_t* strides, int64_t size0, int64_t size1) {
|
||||
std::array<char*, 2> data;
|
||||
std::copy_n(base, 2, data.data());
|
||||
auto loop = [strides_in, requires_neg](char** data, const int64_t* strides, int64_t size0, int64_t size1) {
|
||||
const int64_t *outer_strides = &strides[2];
|
||||
|
||||
for ([[maybe_unused]] const auto it : c10::irange(size1)) {
|
||||
@ -146,9 +144,7 @@ static void reduced_float_copy_kernel(TensorIteratorBase &iter, bool requires_ne
|
||||
|
||||
int64_t grain_size = at::internal::GRAIN_SIZE;
|
||||
|
||||
auto loop = [strides_in, requires_neg](char** base, const int64_t* strides, int64_t size0, int64_t size1) {
|
||||
std::array<char*, 2> data;
|
||||
std::copy_n(base, 2, data.data());
|
||||
auto loop = [strides_in, requires_neg](char** data, const int64_t* strides, int64_t size0, int64_t size1) {
|
||||
const int64_t *outer_strides = &strides[2];
|
||||
|
||||
for ([[maybe_unused]] const auto it : c10::irange(size1)) {
|
||||
|
||||
@ -493,40 +493,33 @@ void cpu_hflip_vec(at::TensorIterator& iter) {
|
||||
|
||||
for ([[maybe_unused]] const auto j : c10::irange(size1)) {
|
||||
// vectorized loop with negative stride for output
|
||||
char** C10_RESTRICT data_ = data_arr.data();
|
||||
int64_t n = size0;
|
||||
|
||||
char* C10_RESTRICT data[ntensors];
|
||||
for (const auto arg : c10::irange(ntensors)) {
|
||||
data[arg] = data_[arg];
|
||||
}
|
||||
|
||||
int64_t i = 0;
|
||||
|
||||
// data[0] unaligned pre-pass
|
||||
// data_arr[0] unaligned pre-pass
|
||||
int64_t offset = (j * n + (n - i - Vec::size())) % 32;
|
||||
offset = (offset >= n) ? n : offset;
|
||||
for (; i < offset; i++) {
|
||||
scalar_t* out_ptr = (scalar_t*)(data[0] - i * stride);
|
||||
*out_ptr = c10::load((scalar_t *)(data[1] + i * stride));
|
||||
scalar_t* out_ptr = (scalar_t*)(data_arr[0] - i * stride);
|
||||
*out_ptr = c10::load((scalar_t *)(data_arr[1] + i * stride));
|
||||
}
|
||||
// Empirically found that it is faster to process 3 data items together vs 2 or 4
|
||||
for (; i <= n - 3 * Vec::size(); i += 3 * Vec::size()) {
|
||||
auto out1 = Vec::loadu(data[1] + i * stride);
|
||||
auto out2 = Vec::loadu(data[1] + (i + Vec::size()) * stride);
|
||||
auto out3 = Vec::loadu(data[1] + (i + 2 * Vec::size()) * stride);
|
||||
auto out1 = Vec::loadu(data_arr[1] + i * stride);
|
||||
auto out2 = Vec::loadu(data_arr[1] + (i + Vec::size()) * stride);
|
||||
auto out3 = Vec::loadu(data_arr[1] + (i + 2 * Vec::size()) * stride);
|
||||
// flip the vector: 1234 -> 4321
|
||||
out1 = flip(out1);
|
||||
out2 = flip(out2);
|
||||
out3 = flip(out3);
|
||||
out1.store(data[0] - (i + Vec::size() - 1) * stride);
|
||||
out2.store(data[0] - (i + 2 * Vec::size() - 1) * stride);
|
||||
out3.store(data[0] - (i + 3 * Vec::size() - 1) * stride);
|
||||
out1.store(data_arr[0] - (i + Vec::size() - 1) * stride);
|
||||
out2.store(data_arr[0] - (i + 2 * Vec::size() - 1) * stride);
|
||||
out3.store(data_arr[0] - (i + 3 * Vec::size() - 1) * stride);
|
||||
}
|
||||
if (i < n) {
|
||||
for (; i < n; i++) {
|
||||
scalar_t* out_ptr = (scalar_t*)(data[0] - i * stride);
|
||||
*out_ptr = c10::load((scalar_t *)(data[1] + i * stride));
|
||||
scalar_t* out_ptr = (scalar_t*)(data_arr[0] - i * stride);
|
||||
*out_ptr = c10::load((scalar_t *)(data_arr[1] + i * stride));
|
||||
}
|
||||
}
|
||||
|
||||
@ -560,15 +553,8 @@ void cpu_vflip_memcpy(at::TensorIterator& iter) {
|
||||
const int64_t stride = strides[0];
|
||||
|
||||
for ([[maybe_unused]] const auto j : c10::irange(size1)) {
|
||||
char** C10_RESTRICT data_ = data_arr.data();
|
||||
int64_t n = size0;
|
||||
|
||||
char* C10_RESTRICT data[ntensors];
|
||||
for (const auto arg : c10::irange(ntensors)) {
|
||||
data[arg] = data_[arg];
|
||||
}
|
||||
|
||||
memcpy(data[0], data[1], n * stride);
|
||||
memcpy(data_arr[0], data_arr[1], n * stride);
|
||||
|
||||
// advance:
|
||||
for (const auto arg : c10::irange(data_arr.size())) {
|
||||
|
||||
@ -139,7 +139,7 @@ void smooth_l1_backward_cpu_kernel(TensorIterator& iter, const Scalar& norm, dou
|
||||
}
|
||||
);
|
||||
} else {
|
||||
AT_DISPATCH_ALL_TYPES_AND(kHalf, dtype, "smooth_l1_backward_cpu_out", [&] {
|
||||
AT_DISPATCH_ALL_TYPES(dtype, "smooth_l1_backward_cpu_out", [&] {
|
||||
auto norm_val = norm.to<scalar_t>();
|
||||
scalar_t beta_val(beta);
|
||||
auto norm_val_vec = Vectorized<scalar_t>(norm_val);
|
||||
|
||||
@ -298,7 +298,7 @@ void unfolded2d_copy(
|
||||
memcpy(
|
||||
dst + (size_t)y * output_width + x,
|
||||
src + (size_t)iy * input_width + ix,
|
||||
sizeof(scalar_t) * (1));
|
||||
sizeof(scalar_t) * 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
@ -317,7 +317,7 @@ void unfolded2d_copy(
|
||||
memcpy(
|
||||
dst + (size_t)y * output_width + x,
|
||||
src + (size_t)iy * input_width + ix + x * dW,
|
||||
sizeof(scalar_t) * (1));
|
||||
sizeof(scalar_t) * 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -342,7 +342,7 @@ void upsample_avx_bilinear_bicubic_uint8(
|
||||
|
||||
if (need_horizontal) {
|
||||
int interp_dim = 3;
|
||||
auto stride = (skip_unpacking) ? num_channels : 4;
|
||||
auto stride = skip_unpacking ? num_channels : 4;
|
||||
std::tie(horiz_indices_weights, ksize_horiz, horiz_weights_precision) =
|
||||
F::compute_index_ranges_int16_weights(
|
||||
/*input_size=*/xin,
|
||||
@ -358,7 +358,7 @@ void upsample_avx_bilinear_bicubic_uint8(
|
||||
|
||||
if (need_vertical) {
|
||||
int interp_dim = 2;
|
||||
auto stride = (skip_unpacking) ? num_channels * xout : 4 * xout;
|
||||
auto stride = skip_unpacking ? num_channels * xout : 4 * xout;
|
||||
std::tie(vert_indices_weights, ksize_vert, vert_weights_precision) =
|
||||
F::compute_index_ranges_int16_weights(
|
||||
/*input_size=*/yin,
|
||||
@ -377,17 +377,17 @@ void upsample_avx_bilinear_bicubic_uint8(
|
||||
// horizontal-only or vertical-only interpolation, and if the tensor doesn't
|
||||
// need repacking
|
||||
if (need_horizontal && (need_vertical || !skip_packing)) {
|
||||
auto c = (skip_unpacking) ? num_channels : 4;
|
||||
auto c = skip_unpacking ? num_channels : 4;
|
||||
buffer_horiz = at::empty({c, yin, xout}, input.options());
|
||||
}
|
||||
if (need_vertical && !skip_packing) {
|
||||
auto c = (skip_unpacking) ? num_channels : 4;
|
||||
auto c = skip_unpacking ? num_channels : 4;
|
||||
buffer_vert = at::empty({c, yout, xout}, input.options());
|
||||
}
|
||||
|
||||
for (const auto i : c10::irange(batch_size)) {
|
||||
|
||||
at::Tensor unpacked_input = (skip_unpacking) ? input[i] : unpack_rgb(input[i]);
|
||||
at::Tensor unpacked_input = skip_unpacking ? input[i] : unpack_rgb(input[i]);
|
||||
at::Tensor unpacked_output;
|
||||
|
||||
if (need_horizontal) {
|
||||
@ -411,7 +411,7 @@ void upsample_avx_bilinear_bicubic_uint8(
|
||||
unpacked_output = unpacked_input = unpacked_output_temp;
|
||||
}
|
||||
if (need_vertical) {
|
||||
unpacked_output = (skip_packing) ? output[i] : buffer_vert;
|
||||
unpacked_output = skip_packing ? output[i] : buffer_vert;
|
||||
|
||||
ImagingResampleVertical(
|
||||
unpacked_output,
|
||||
@ -502,7 +502,7 @@ void ImagingResampleHorizontalConvolution8u4x(
|
||||
// RGBA: b4_delta = b4_delta_soft = 3
|
||||
// RGB : b4_delta = 5
|
||||
// RGB : b4_delta_soft = 4
|
||||
const auto b4_delta = (stride == 4) ? 3 : ((is_last_line) ? 5 : 4);
|
||||
const auto b4_delta = (stride == 4) ? 3 : (is_last_line ? 5 : 4);
|
||||
|
||||
// In block 2 (2 means we process 2 weights values together), we read input data
|
||||
// with _mm_loadl_epi64, i.e. 8 bytes, per one line:
|
||||
@ -515,7 +515,7 @@ void ImagingResampleHorizontalConvolution8u4x(
|
||||
// RGBA: b2_delta = b2_delta_soft = 1
|
||||
// RGB : b2_delta = 2
|
||||
// RGB : b2_delta_soft = 1
|
||||
const auto b2_delta = (stride == 4) ? 1 : ((is_last_line) ? 2 : 1);
|
||||
const auto b2_delta = (stride == 4) ? 1 : (is_last_line ? 2 : 1);
|
||||
|
||||
const auto max_out_x_strided = out_xsize * stride;
|
||||
const auto max_in_x_strided = in_xsize * stride;
|
||||
@ -819,7 +819,7 @@ void ImagingResampleHorizontalConvolution8u(
|
||||
// RGBA: b8_delta = b8_delta_soft = 7
|
||||
// RGB : b8_delta = 10
|
||||
// RGB : b8_delta_soft = 9
|
||||
const auto b8_delta = (stride == 4) ? 7 : ((is_last_line) ? 10 : 9);
|
||||
const auto b8_delta = (stride == 4) ? 7 : (is_last_line ? 10 : 9);
|
||||
|
||||
// In block 4 (4 means we process 4 weight values together), we read
|
||||
// 16 bytes of input data.
|
||||
@ -832,7 +832,7 @@ void ImagingResampleHorizontalConvolution8u(
|
||||
// RGBA: b4_delta = b4_delta_soft = 3
|
||||
// RGB : b4_delta = 5
|
||||
// RGB : b4_delta_soft = 4
|
||||
const auto b4_delta = (stride == 4) ? 3 : ((is_last_line) ? 5 : 4);
|
||||
const auto b4_delta = (stride == 4) ? 3 : (is_last_line ? 5 : 4);
|
||||
|
||||
// In block 2 (2 means we process 2 weight values together), we read
|
||||
// 8 bytes of input data.
|
||||
@ -845,7 +845,7 @@ void ImagingResampleHorizontalConvolution8u(
|
||||
// RGBA: b2_delta = b2_delta_soft = 1
|
||||
// RGB : b2_delta = 2
|
||||
// RGB : b2_delta_soft = 1
|
||||
const auto b2_delta = (stride == 4) ? 1 : ((is_last_line) ? 2 : 1);
|
||||
const auto b2_delta = (stride == 4) ? 1 : (is_last_line ? 2 : 1);
|
||||
|
||||
const auto max_out_x_strided = out_xsize * stride;
|
||||
const auto max_in_x_strided = in_xsize * stride;
|
||||
|
||||
@ -644,8 +644,8 @@ void weight_to_int4pack_kernel(
|
||||
int32_t val2 = src[(d + 32) * K + k];
|
||||
int32_t val3 = src[(d + 48) * K + k];
|
||||
|
||||
uint8_t packed02 = (((uint8_t)(val2) << 4)) | ((uint8_t)(val0));
|
||||
uint8_t packed13 = (((uint8_t)(val3) << 4)) | ((uint8_t)(val1));
|
||||
uint8_t packed02 = ((uint8_t)val2 << 4) | ((uint8_t)val0);
|
||||
uint8_t packed13 = ((uint8_t)val3 << 4) | ((uint8_t)val1);
|
||||
|
||||
dst[k * 32 + d] = packed02;
|
||||
dst[k * 32 + 16 + d] = packed13;
|
||||
@ -656,7 +656,7 @@ void weight_to_int4pack_kernel(
|
||||
int32_t val0 = src[n * K + k];
|
||||
int32_t val1 = src[n * K + K + k];
|
||||
|
||||
uint8_t packed = (((uint8_t)(val1) << 4)) | ((uint8_t)(val0));
|
||||
uint8_t packed = ((uint8_t)val1 << 4) | ((uint8_t)val0);
|
||||
dst[k * nb_size / 2 + n / 2] = packed;
|
||||
}
|
||||
}
|
||||
@ -667,7 +667,7 @@ void weight_to_int4pack_kernel(
|
||||
int32_t val0 = src[(d + 0) * K + k];
|
||||
int32_t val1 = src[(d + 16) * K + k];
|
||||
|
||||
uint8_t packed01 = (((uint8_t)(val1) << 4)) | ((uint8_t)(val0));
|
||||
uint8_t packed01 = ((uint8_t)val1 << 4) | ((uint8_t)val0);
|
||||
dst[k * 16 + d] = packed01;
|
||||
}
|
||||
} else {
|
||||
@ -676,7 +676,7 @@ void weight_to_int4pack_kernel(
|
||||
int32_t val0 = src[n * K + k];
|
||||
int32_t val1 = src[n * K + K + k];
|
||||
|
||||
uint8_t packed = (((uint8_t)(val1) << 4)) | ((uint8_t)(val0));
|
||||
uint8_t packed = ((uint8_t)val1 << 4) | ((uint8_t)val0);
|
||||
dst[k * nb_size / 2 + n / 2] = packed;
|
||||
}
|
||||
}
|
||||
@ -685,7 +685,7 @@ void weight_to_int4pack_kernel(
|
||||
int32_t val0 = src[n * K + k];
|
||||
int32_t val1 = src[n * K + K + k];
|
||||
|
||||
uint8_t packed = (((uint8_t)(val1) << 4)) | ((uint8_t)(val0));
|
||||
uint8_t packed = ((uint8_t)val1 << 4) | ((uint8_t)val0);
|
||||
dst[k * nb_size / 2 + n / 2] = packed;
|
||||
}
|
||||
#endif
|
||||
@ -872,16 +872,16 @@ void ref_dyn_quant_matmul_4bit_channelwise_kernel(
|
||||
for (size_t k_idx = 0; k_idx < k; ++k_idx) {
|
||||
const float src0_0 = src_ptr[k_idx];
|
||||
|
||||
max0 = (std::max)(src0_0, max0);
|
||||
min0 = (std::min)(src0_0, min0);
|
||||
max0 = std::max(src0_0, max0);
|
||||
min0 = std::min(src0_0, min0);
|
||||
}
|
||||
|
||||
// Maximum/minimum int8 values
|
||||
const float qmin = (float)INT8_MIN;
|
||||
const float qmax = (float)INT8_MAX;
|
||||
|
||||
const float rmin0 = (std::min)(0.0f, min0);
|
||||
const float rmax0 = (std::max)(0.0f, max0);
|
||||
const float rmin0 = std::min(0.0f, min0);
|
||||
const float rmax0 = std::max(0.0f, max0);
|
||||
|
||||
const float scale0 =
|
||||
rmin0 == rmax0 ? 1.f : (qmax - qmin) / (rmax0 - rmin0);
|
||||
@ -900,8 +900,8 @@ void ref_dyn_quant_matmul_4bit_channelwise_kernel(
|
||||
? qmin - descaled_min0
|
||||
: qmax - descaled_max0;
|
||||
|
||||
zero_point0 = (std::max)(zero_point0, qmin);
|
||||
zero_point0 = (std::min)(zero_point0, qmax);
|
||||
zero_point0 = std::max(zero_point0, qmin);
|
||||
zero_point0 = std::min(zero_point0, qmax);
|
||||
|
||||
// Round to nearest integer
|
||||
const int32_t nudged_zero_point0 = lrintf(zero_point0);
|
||||
@ -909,9 +909,9 @@ void ref_dyn_quant_matmul_4bit_channelwise_kernel(
|
||||
int8_t* dst_ptr = lhs_qa8dx + m_idx * dst_stride;
|
||||
|
||||
// LHS offset at the beginning of the row
|
||||
*((float*)(dst_ptr)) = recip_scale0;
|
||||
*((float*)dst_ptr) = recip_scale0;
|
||||
dst_ptr += sizeof(float);
|
||||
*((int32_t*)(dst_ptr)) = -nudged_zero_point0;
|
||||
*((int32_t*)dst_ptr) = -nudged_zero_point0;
|
||||
dst_ptr += sizeof(int32_t);
|
||||
|
||||
// Quantize the channels
|
||||
@ -922,8 +922,8 @@ void ref_dyn_quant_matmul_4bit_channelwise_kernel(
|
||||
int32_t v0_s32 = (int32_t)(std::round(src0_0 * scale0));
|
||||
|
||||
v0_s32 = v0_s32 + nudged_zero_point0;
|
||||
v0_s32 = (std::max)(v0_s32, static_cast<int32_t>(INT8_MIN));
|
||||
v0_s32 = (std::min)(v0_s32, static_cast<int32_t>(INT8_MAX));
|
||||
v0_s32 = std::max(v0_s32, static_cast<int32_t>(INT8_MIN));
|
||||
v0_s32 = std::min(v0_s32, static_cast<int32_t>(INT8_MAX));
|
||||
dst_ptr[0] = (int8_t)v0_s32;
|
||||
dst_ptr += sizeof(int8_t);
|
||||
}
|
||||
@ -988,8 +988,8 @@ void ref_dyn_quant_matmul_4bit_channelwise_kernel(
|
||||
main_acc = main_acc * lhs_scale;
|
||||
|
||||
// Clamp (min-max) operation
|
||||
main_acc = (std::max)(main_acc, scalar_min);
|
||||
main_acc = (std::min)(main_acc, scalar_max);
|
||||
main_acc = std::max(main_acc, scalar_min);
|
||||
main_acc = std::min(main_acc, scalar_max);
|
||||
|
||||
dst_f32[0] = main_acc;
|
||||
dst_f32 += 1;
|
||||
@ -1024,15 +1024,15 @@ void ref_dyn_quant_matmul_4bit_groupwise_kernel(
|
||||
|
||||
for (size_t k_idx = 0; k_idx < k; ++k_idx) {
|
||||
const float src0_0 = src_ptr[k_idx];
|
||||
max0 = (std::max)(src0_0, max0);
|
||||
min0 = (std::min)(src0_0, min0);
|
||||
max0 = std::max(src0_0, max0);
|
||||
min0 = std::min(src0_0, min0);
|
||||
}
|
||||
|
||||
const float qmin = (float)INT8_MIN;
|
||||
const float qmax = (float)INT8_MAX;
|
||||
|
||||
const float rmin0 = (std::min)(0.0f, min0);
|
||||
const float rmax0 = (std::max)(0.0f, max0);
|
||||
const float rmin0 = std::min(0.0f, min0);
|
||||
const float rmax0 = std::max(0.0f, max0);
|
||||
const float scale0 =
|
||||
(rmin0 == rmax0) ? 1.f : (qmax - qmin) / (rmax0 - rmin0);
|
||||
const float recip_scale0 = scale0 ? 1.0f / scale0 : 0.0f;
|
||||
@ -1044,22 +1044,22 @@ void ref_dyn_quant_matmul_4bit_groupwise_kernel(
|
||||
? qmin - descaled_min0
|
||||
: qmax - descaled_max0;
|
||||
|
||||
zero_point0 = (std::max)(zero_point0, qmin);
|
||||
zero_point0 = (std::min)(zero_point0, qmax);
|
||||
zero_point0 = std::max(zero_point0, qmin);
|
||||
zero_point0 = std::min(zero_point0, qmax);
|
||||
const int32_t nudged_zero_point0 = lrintf(zero_point0);
|
||||
|
||||
int8_t* dst_ptr = lhs_qa8dx + row_idx * dst_stride;
|
||||
|
||||
*((float*)(dst_ptr)) = recip_scale0;
|
||||
*((float*)dst_ptr) = recip_scale0;
|
||||
dst_ptr += sizeof(float);
|
||||
*((int32_t*)(dst_ptr)) = -nudged_zero_point0;
|
||||
*((int32_t*)dst_ptr) = -nudged_zero_point0;
|
||||
dst_ptr += sizeof(int32_t);
|
||||
|
||||
for (size_t k_idx = 0; k_idx < k; ++k_idx) {
|
||||
const float src0_0 = src_ptr[k_idx];
|
||||
int32_t v0_s32 = (int32_t)(std::round(src0_0 * scale0));
|
||||
v0_s32 = (std::max)(
|
||||
(std::min)(
|
||||
v0_s32 = std::max(
|
||||
std::min(
|
||||
v0_s32 + nudged_zero_point0, static_cast<int32_t>(INT8_MAX)),
|
||||
static_cast<int32_t>(INT8_MIN));
|
||||
dst_ptr[0] = (int8_t)v0_s32;
|
||||
@ -1118,8 +1118,8 @@ void ref_dyn_quant_matmul_4bit_groupwise_kernel(
|
||||
}
|
||||
|
||||
main_acc = main_acc * lhs_scale;
|
||||
main_acc = (std::max)(main_acc, scalar_min);
|
||||
main_acc = (std::min)(main_acc, scalar_max);
|
||||
main_acc = std::max(main_acc, scalar_min);
|
||||
main_acc = std::min(main_acc, scalar_max);
|
||||
|
||||
dst_f32[0] = main_acc;
|
||||
dst_f32 += 1;
|
||||
|
||||
@ -4,7 +4,6 @@
|
||||
#include <c10/util/SmallVector.h>
|
||||
#include <c10/core/Scalar.h>
|
||||
#include <c10/core/ScalarType.h>
|
||||
#include <c10/util/Exception.h>
|
||||
#define TORCH_ASSERT_ONLY_METHOD_OPERATORS
|
||||
#include <ATen/core/Tensor.h>
|
||||
#include <ATen/core/NamedTensor.h>
|
||||
|
||||
@ -753,8 +753,8 @@ static void apply_cholesky_cusolver_potrf_looped(const Tensor& self_working_copy
|
||||
handle, params, uplo, n, datatype,
|
||||
self_working_copy_ptr + i * matrix_stride,
|
||||
lda, datatype,
|
||||
(char*)workdata_device_ptr + i * worksize_device, worksize_device,
|
||||
(char*)workdata_host_ptr + i * worksize_host, worksize_host,
|
||||
static_cast<char*>(workdata_device_ptr) + i * worksize_device, worksize_device,
|
||||
static_cast<char*>(workdata_host_ptr) + i * worksize_host, worksize_host,
|
||||
infos_ptr + i
|
||||
);
|
||||
}
|
||||
|
||||
@ -119,8 +119,8 @@ void setConvolutionParams(
|
||||
params->input_dim = input.dim();
|
||||
params->memory_format = memory_format;
|
||||
for (int i = 0; i != params->input_dim; ++i) {
|
||||
params->input_size[i] = (int)input.sizes()[i];
|
||||
params->weight_size[i] = (int)weight.sizes()[i];
|
||||
params->input_size[i] = static_cast<int>(input.sizes()[i]);
|
||||
params->weight_size[i] = static_cast<int>(weight.sizes()[i]);
|
||||
}
|
||||
// ASSERT(padding.size() == stride.size())
|
||||
// ASSERT(padding.size() == dilation.size())
|
||||
|
||||
@ -64,7 +64,7 @@
|
||||
// fastest algorithm combination with a sub optimal mathType.
|
||||
|
||||
constexpr size_t operator"" _TiB(unsigned long long n) {
|
||||
return size_t(n) * 1024 * 1024 * 1024 * 1024;
|
||||
return static_cast<size_t>(n) * 1024 * 1024 * 1024 * 1024;
|
||||
}
|
||||
|
||||
namespace at {
|
||||
|
||||
@ -46,7 +46,7 @@ namespace {
|
||||
|
||||
// TODO: remove duplicate code in Conv_v7.cpp
|
||||
constexpr int64_t operator"" _TiB(unsigned long long n) {
|
||||
return size_t(n) << 40;
|
||||
return static_cast<size_t>(n) << 40;
|
||||
}
|
||||
|
||||
uint8_t getAlignment(const Tensor& t) {
|
||||
@ -93,7 +93,10 @@ cudnn_frontend::Tensor getTensorDescriptorWithTypeVirtual(
|
||||
|
||||
std::vector<int64_t> strides_copy(std::begin(strides), std::end(strides));
|
||||
fixSizeOneDimStride<int64_t>(
|
||||
sizes.size(), &sizes[0], (int64_t*)&strides_copy[0], channels_last);
|
||||
sizes.size(),
|
||||
&sizes[0],
|
||||
static_cast<int64_t*>(&strides_copy[0]),
|
||||
channels_last);
|
||||
auto r = cudnn_frontend::TensorBuilder()
|
||||
.setDim(sizes.size(), sizes.data())
|
||||
.setStrides(strides_copy.size(), strides_copy.data())
|
||||
|
||||
@ -44,6 +44,7 @@ std::tuple<Tensor, Tensor> cudnn_grid_sampler_backward(
|
||||
#include <ATen/cudnn/Descriptors.h>
|
||||
#include <ATen/cudnn/Types.h>
|
||||
#include <ATen/cudnn/Utils.h>
|
||||
#include <array>
|
||||
|
||||
#include <ATen/TensorUtils.h>
|
||||
#include <c10/util/irange.h>
|
||||
@ -59,11 +60,11 @@ void setSamplerDescriptor(
|
||||
SpatialTransformerDescriptor& desc,
|
||||
cudnnDataType_t dataType,
|
||||
const at::Tensor& tensor) {
|
||||
int inputSize[4] = {0};
|
||||
std::array<int, 4> inputSize{0};
|
||||
for (const auto i : c10::irange(tensor.dim())) {
|
||||
inputSize[i] = (int)tensor.size(i);
|
||||
inputSize[i] = static_cast<int>(tensor.size(i));
|
||||
}
|
||||
desc.set(dataType, 4, inputSize);
|
||||
desc.set(dataType, 4, inputSize.data());
|
||||
}
|
||||
|
||||
void checkGridSize(CheckedFrom c, TensorArg grid, TensorArg input) {
|
||||
|
||||
@ -656,7 +656,8 @@ void add_projection_weights(
|
||||
TORCH_INTERNAL_ASSERT(
|
||||
nb_dims <= min_dim, "nb_dims = ", nb_dims, "; min_dim = ", min_dim);
|
||||
auto elem_size = dataSize(getCudnnDataType(weight_buf));
|
||||
auto offset_bytes = (char*)matrix_pointer - (char*)weight_buf.data_ptr();
|
||||
auto offset_bytes = static_cast<const char*>(matrix_pointer) -
|
||||
static_cast<const char*>(weight_buf.data_ptr());
|
||||
TORCH_INTERNAL_ASSERT(
|
||||
offset_bytes % elem_size == 0,
|
||||
"offset_bytes = ",
|
||||
@ -794,8 +795,8 @@ get_parameters(
|
||||
"; min_dim = ",
|
||||
min_dim);
|
||||
auto elem_size = dataSize(getCudnnDataType(weight_buf));
|
||||
auto offset_bytes =
|
||||
(char*)matrix_pointer - (char*)weight_buf.data_ptr();
|
||||
auto offset_bytes = static_cast<const char*>(matrix_pointer) -
|
||||
static_cast<const char*>(weight_buf.data_ptr());
|
||||
TORCH_INTERNAL_ASSERT(
|
||||
offset_bytes % elem_size == 0,
|
||||
"offset_bytes = ",
|
||||
|
||||
@ -330,7 +330,6 @@ Tensor _fft_c2c_mkl(const Tensor& self, IntArrayRef dim, int64_t normalization,
|
||||
}
|
||||
|
||||
#elif AT_MKL_ENABLED()
|
||||
#include <ATen/Dispatch.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <numeric>
|
||||
|
||||
@ -535,7 +535,7 @@ mkldnn_scaled_mm(const Tensor& mat1, const Tensor& mat2,
|
||||
|
||||
float input_scale = scale_a.item<float>();
|
||||
float weight_scale = scale_b.item<float>();
|
||||
float output_scale = float(1.0);
|
||||
float output_scale = 1.0f;
|
||||
if (scale_result.has_value() &&
|
||||
(*out_dtype == ScalarType::Float8_e4m3fn ||
|
||||
*out_dtype == ScalarType::Float8_e5m2)) {
|
||||
|
||||
@ -530,7 +530,7 @@ static Tensor get_mkldnn_serialized_md(const Tensor& self) {
|
||||
#else
|
||||
TORCH_CHECK(false, "Unexpected IDeep version to do weight serialization.");
|
||||
#endif
|
||||
Tensor serialized_md = at::from_blob((void*)serialized_wei_desc.data(), {(int64_t)serialized_wei_desc.size()}, at::TensorOptions(at::kByte));
|
||||
Tensor serialized_md = at::from_blob((void*)serialized_wei_desc.data(), {static_cast<int64_t>(serialized_wei_desc.size())}, at::TensorOptions(at::kByte));
|
||||
auto res = at::empty_like(serialized_md);
|
||||
// serialized_md shares the buffer with serialized_wei_desc,
|
||||
// which will be released outside of this function thus invalidating the buffer of serialized_md.
|
||||
|
||||
@ -576,14 +576,14 @@ static void _mkldnn_gemm_i8i8i32_with_blas(
|
||||
n,
|
||||
k,
|
||||
alpha,
|
||||
(int8_t*)self.data_ptr(),
|
||||
static_cast<int8_t*>(self.data_ptr()),
|
||||
lda,
|
||||
ao,
|
||||
(int8_t*)mat2.data_ptr(),
|
||||
static_cast<int8_t*>(mat2.data_ptr()),
|
||||
ldb,
|
||||
bo,
|
||||
beta,
|
||||
(int32_t*)result.data_ptr(),
|
||||
static_cast<int32_t*>(result.data_ptr()),
|
||||
ldc,
|
||||
&co);
|
||||
}
|
||||
|
||||
@ -41,7 +41,7 @@ void woq_matmul_int4_impl(
|
||||
dst_usr_dims;
|
||||
dnnl::memory::dims m1_usr_strides, m2_usr_strides, scale_usr_strides,
|
||||
zp_usr_strides, dst_usr_strides;
|
||||
int compressed_k = (int)(k / 8);
|
||||
int compressed_k = k / 8;
|
||||
int num_groups = (int)(k / group_size);
|
||||
m1_usr_dims = {m, k};
|
||||
m1_usr_strides = {m1.stride(0), m1.stride(1)};
|
||||
|
||||
@ -370,7 +370,7 @@ static void nllnd_loss_backward_impl(Tensor& grad_input_arg,
|
||||
onValue:-1.0f
|
||||
offValue:0.0f
|
||||
name:nil];
|
||||
oneHotTensor = castMPSTensor(mpsGraph, oneHotTensor, [inputTensor dataType]);
|
||||
oneHotTensor = castMPSTensor(mpsGraph, oneHotTensor, inputTensor.dataType);
|
||||
if (isWeightsArrayValid) {
|
||||
oneHotTensor = [mpsGraph multiplicationWithPrimaryTensor:oneHotTensor
|
||||
secondaryTensor:weightTensor
|
||||
@ -705,7 +705,6 @@ static void smooth_l1_loss_template(const Tensor& input,
|
||||
TORCH_CHECK(beta >= 0, "smooth_l1_loss does not support negative values for beta.");
|
||||
TORCH_CHECK(input.is_mps());
|
||||
TORCH_CHECK(target.is_mps());
|
||||
TORCH_CHECK_NOT_IMPLEMENTED(input.scalar_type() != kLong, "MPS doesn't know how to do square_i64");
|
||||
if ((input.numel() == 0) || (target.numel() == 0)) {
|
||||
reduction == Reduction::Mean ? output.fill_(std::numeric_limits<float>::quiet_NaN()) : output.zero_();
|
||||
return;
|
||||
@ -772,7 +771,7 @@ static void smooth_l1_loss_backward_impl(const Tensor& grad_output,
|
||||
MPSGraphTensor* targetTensor = mpsGraphRankedPlaceHolder(mpsGraph, target);
|
||||
MPSGraphTensor* gradOutputTensor = mpsGraphRankedPlaceHolder(mpsGraph, grad_output);
|
||||
|
||||
MPSGraphTensor* betaTensor = [mpsGraph constantWithScalar:beta dataType:[inputTensor dataType]];
|
||||
MPSGraphTensor* betaTensor = [mpsGraph constantWithScalar:beta dataType:MPSDataTypeFloat32];
|
||||
// xn - yn
|
||||
MPSGraphTensor* diffTensor = [mpsGraph subtractionWithPrimaryTensor:inputTensor
|
||||
secondaryTensor:targetTensor
|
||||
@ -798,8 +797,7 @@ static void smooth_l1_loss_backward_impl(const Tensor& grad_output,
|
||||
name:@"lossTensor"];
|
||||
MPSGraphTensor* outputTensor = lossTensor;
|
||||
if (reduction == Reduction::Mean) {
|
||||
MPSGraphTensor* numelTensor = [mpsGraph constantWithScalar:(double)input.numel()
|
||||
dataType:[lossTensor dataType]];
|
||||
MPSGraphTensor* numelTensor = [mpsGraph constantWithScalar:(double)input.numel() dataType:MPSDataTypeFloat32];
|
||||
outputTensor = [mpsGraph divisionWithPrimaryTensor:lossTensor secondaryTensor:numelTensor name:nil];
|
||||
}
|
||||
MPSGraphTensor* gradInputTensor = [mpsGraph multiplicationWithPrimaryTensor:outputTensor
|
||||
|
||||
@ -10,6 +10,7 @@
|
||||
#include <ATen/Functions.h>
|
||||
#include <ATen/NativeFunctions.h>
|
||||
#else
|
||||
#include <ATen/ops/aminmax.h>
|
||||
#include <ATen/ops/avg_pool2d.h>
|
||||
#include <ATen/ops/avg_pool2d_backward.h>
|
||||
#include <ATen/ops/avg_pool2d_backward_native.h>
|
||||
@ -544,8 +545,9 @@ static void max_unpool_out_mps_template(const Tensor& input,
|
||||
if (indices.defined() && indices.numel() > 0) {
|
||||
auto output_image_size = c10::multiply_integers(output_size_);
|
||||
|
||||
int64_t min_idx = indices.min().item<int64_t>();
|
||||
int64_t max_idx = indices.max().item<int64_t>();
|
||||
auto [min_idx_tensor, max_idx_tensor] = indices.aminmax();
|
||||
int64_t min_idx = min_idx_tensor.item<int64_t>();
|
||||
int64_t max_idx = max_idx_tensor.item<int64_t>();
|
||||
|
||||
if (min_idx < 0 || max_idx >= output_image_size) {
|
||||
int64_t error_idx = (min_idx < 0) ? min_idx : max_idx;
|
||||
|
||||
@ -83,6 +83,31 @@ std::string get_type_str<int32_t>() {
|
||||
return "int32_t";
|
||||
}
|
||||
|
||||
// If all tensors are contiguous with the same dtype and the cat dimension is 0,
|
||||
// then we can simply copy each tensor's underlying buffer contiguously into the
|
||||
// output.
|
||||
static void cat_out_mps_contiguous_impl(const ITensorListRef& inputs, const Tensor& output) {
|
||||
MPSStream* stream = getCurrentMPSStream();
|
||||
id<MTLBuffer> output_buffer = getMTLBufferStorage(output);
|
||||
size_t output_offset = output.storage_offset() * output.itemsize();
|
||||
|
||||
for (const Tensor& input : inputs) {
|
||||
if (cat_should_skip_tensor(input)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
id<MTLBuffer> input_buffer = getMTLBufferStorage(input);
|
||||
size_t input_offset = input.storage_offset() * input.itemsize();
|
||||
auto nbytes = input.nbytes();
|
||||
auto profile_id =
|
||||
getMPSProfiler().beginProfileCopy(input_buffer, output_buffer, input, output, nbytes, /*non_blocking=*/true);
|
||||
|
||||
stream->copy(input_buffer, output_buffer, nbytes, input_offset, output_offset, profile_id, SyncType::NONE);
|
||||
|
||||
output_offset += nbytes;
|
||||
}
|
||||
}
|
||||
|
||||
// NOTE: `output` is expected to already have the correct size.
|
||||
template <typename idx_type_t>
|
||||
static void cat_out_mps_impl(const ITensorListRef& inputs, int64_t dimension, const Tensor& output) {
|
||||
@ -105,7 +130,7 @@ static void cat_out_mps_impl(const ITensorListRef& inputs, int64_t dimension, co
|
||||
// copy all the input tensor data into a packed buffer, which would not be
|
||||
// ideal.
|
||||
for (const Tensor& input : inputs) {
|
||||
if (input.numel() == 0) {
|
||||
if (cat_should_skip_tensor(input)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
@ -243,101 +268,16 @@ TORCH_IMPL_FUNC(cat_out_mps)
|
||||
if (out.numel() == 0) {
|
||||
return;
|
||||
}
|
||||
|
||||
auto materialized_inputs = inputs.materialize();
|
||||
auto out_dtype = at::native::result_type(inputs);
|
||||
bool has_large_tensor =
|
||||
isTooLargeForMPSGraph(out) || std::any_of(materialized_inputs.begin(), materialized_inputs.end(), [](auto& t) {
|
||||
return !cat_should_skip_tensor(t) && isTooLargeForMPSGraph(t);
|
||||
});
|
||||
|
||||
int idx = 0;
|
||||
for (const Tensor& t : materialized_inputs) {
|
||||
TORCH_CHECK(t.dim() > 0, "zero-dimensional tensor (at position ", idx, ") cannot be concatenated");
|
||||
auto lap = at::get_overlap_status(out, t);
|
||||
TORCH_CHECK(lap != at::MemOverlapStatus::Partial && lap != at::MemOverlapStatus::Full,
|
||||
"torch.cat(): unsupported operation: the input tensors cannot refer to any "
|
||||
"of the output memory locations. Found overlap in input tensor ",
|
||||
idx);
|
||||
idx++;
|
||||
}
|
||||
// Check for type promotion
|
||||
TORCH_CHECK(canCast(out_dtype, out.scalar_type()),
|
||||
"torch.cat(): input types can't be cast to the desired output type ",
|
||||
out.scalar_type());
|
||||
TORCH_CHECK(!inputs.empty(), "torch.cat(): invalid number of inputs ", inputs.size());
|
||||
|
||||
dimension = legacy_cat_wrap_dim(dimension, materialized_inputs);
|
||||
TORCH_CHECK(dimension >= 0, "torch.cat(): invalid dimension ", dimension);
|
||||
|
||||
// previously, size [0] tensors were the only possible empty tensors; thus, it
|
||||
// wasn't possible to cat empty tensors unless all the other tensors were
|
||||
// 1-dimensional, so we allowed these tensors to be "skipped". We maintain
|
||||
// this behavior for backwards compatibility, but only for this specific size
|
||||
// (i.e. other empty sizes are not skipped).
|
||||
// FIXME: warn if this is the case
|
||||
auto should_skip = [](const Tensor& t) { return t.dim() == 1 && t.size(0) == 0; };
|
||||
at::assert_no_internal_overlap(out);
|
||||
|
||||
Tensor notSkippedTensor;
|
||||
// Indices of tensors to be skipped because they're empty
|
||||
std::vector<int64_t> skipped_tensor_indices;
|
||||
// Tensors to be read
|
||||
std::vector<Tensor> input_tensors;
|
||||
int tensor_idx = 0;
|
||||
for (const Tensor& t : materialized_inputs) {
|
||||
if (t.numel() == 0 || should_skip(t)) {
|
||||
skipped_tensor_indices.push_back(tensor_idx);
|
||||
tensor_idx++;
|
||||
continue;
|
||||
}
|
||||
input_tensors.push_back(t);
|
||||
// TODO: Is this OK?
|
||||
notSkippedTensor = t;
|
||||
tensor_idx++;
|
||||
}
|
||||
// If all inputs are empty tensors, return an empty tensor
|
||||
if (!notSkippedTensor.defined()) {
|
||||
return;
|
||||
}
|
||||
for (const Tensor& t : inputs) {
|
||||
TORCH_CHECK(t.device() == notSkippedTensor.device(),
|
||||
"torch.cat(): all input tensors must be on the same device. Received ",
|
||||
t.device(),
|
||||
" and ",
|
||||
notSkippedTensor.device());
|
||||
}
|
||||
TORCH_CHECK(out.device() == notSkippedTensor.device(),
|
||||
"torch.cat(): all input tensors and out must be on the same device, but inputs are on ",
|
||||
notSkippedTensor.device(),
|
||||
" and out is on ",
|
||||
out.device());
|
||||
|
||||
std::vector<int64_t> size(notSkippedTensor.sizes().vec());
|
||||
|
||||
// Compute size of the result in the cat dimension
|
||||
int64_t cat_dim_size = 0;
|
||||
idx = 0;
|
||||
bool has_large_tensor = false;
|
||||
for (const Tensor& tensor : materialized_inputs) {
|
||||
if (isTooLargeForMPSGraph(tensor)) {
|
||||
has_large_tensor |= true;
|
||||
}
|
||||
if (!should_skip(tensor)) {
|
||||
// TODO: Factor out `check_shape_except_dim`
|
||||
check_shape_except_dim(notSkippedTensor, tensor, dimension, idx);
|
||||
cat_dim_size += tensor.size(dimension);
|
||||
idx++;
|
||||
}
|
||||
}
|
||||
// Compute the size of the result
|
||||
size[dimension] = cat_dim_size;
|
||||
// skip resizing if size of result is same as expected
|
||||
if (out.sizes() != size) {
|
||||
out.resize_(size, MemoryFormat::Contiguous);
|
||||
}
|
||||
if (out.numel() == 0) {
|
||||
return;
|
||||
}
|
||||
|
||||
has_large_tensor |= isTooLargeForMPSGraph(out);
|
||||
|
||||
if (has_large_tensor) {
|
||||
if (all_contiguous && all_same_dtype && (memory_format == MemoryFormat::Contiguous) && (dimension == 0)) {
|
||||
return mps::cat_out_mps_contiguous_impl(materialized_inputs, out);
|
||||
} else if (has_large_tensor) {
|
||||
return mps::cat_out_mps_impl<int64_t>(materialized_inputs, dimension, out);
|
||||
} else {
|
||||
return mps::cat_out_mps_impl<int32_t>(materialized_inputs, dimension, out);
|
||||
|
||||
@ -2602,12 +2602,16 @@
|
||||
device_check: NoCheck # TensorIterator
|
||||
structured_delegate: exp.out
|
||||
variants: function, method
|
||||
dispatch:
|
||||
SparseCPU, SparseCUDA, SparseMPS: exp_sparse
|
||||
tags: [core, pointwise]
|
||||
|
||||
- func: exp_(Tensor(a!) self) -> Tensor(a!)
|
||||
device_check: NoCheck # TensorIterator
|
||||
structured_delegate: exp.out
|
||||
variants: function, method
|
||||
dispatch:
|
||||
SparseCPU, SparseCUDA, SparseMPS: exp_sparse_
|
||||
tags: pointwise
|
||||
|
||||
- func: exp.out(Tensor self, *, Tensor(a!) out) -> Tensor(a!)
|
||||
@ -2616,6 +2620,7 @@
|
||||
structured_inherits: TensorIteratorBase
|
||||
dispatch:
|
||||
CPU, CUDA, MPS, MTIA: exp_out
|
||||
SparseCPU, SparseCUDA, SparseMPS: exp_sparse_out
|
||||
tags: pointwise
|
||||
|
||||
- func: exp2(Tensor self) -> Tensor
|
||||
|
||||
@ -65,7 +65,7 @@ void quantize_vec(
|
||||
(typename T::underlying*)dst,
|
||||
count,
|
||||
fbgemm::TensorQuantizationParams{
|
||||
(float)scale, (int32_t)zero_point, precision});
|
||||
static_cast<float>(scale), static_cast<int32_t>(zero_point), precision});
|
||||
}
|
||||
|
||||
#if defined(__ARM_NEON__) || defined(__aarch64__)
|
||||
|
||||
@ -40,7 +40,7 @@ inline int start_index(int out_idx, int out_len, int in_len) {
|
||||
* This function computes the start index on input matrix.
|
||||
*/
|
||||
// NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions)
|
||||
return (int)std::floor((float)(out_idx * in_len) / out_len);
|
||||
return static_cast<int>(std::floor(static_cast<float>(out_idx * in_len) / out_len));
|
||||
}
|
||||
|
||||
inline int end_index(int out_idx, int out_len, int in_len) {
|
||||
@ -49,7 +49,7 @@ inline int end_index(int out_idx, int out_len, int in_len) {
|
||||
* This function computes the end index on input matrix.
|
||||
*/
|
||||
// NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions)
|
||||
return (int)std::ceil((float)((out_idx + 1) * in_len) / out_len);
|
||||
return static_cast<int>(std::ceil(static_cast<float>((out_idx + 1) * in_len) / out_len));
|
||||
}
|
||||
|
||||
// adaptive avg pool for 2D and 3D inputs
|
||||
|
||||
@ -71,8 +71,8 @@ void avg_pool2d_out_frame(
|
||||
int64_t hend = std::min(hstart + kH, inputHeight + padH);
|
||||
int64_t wend = std::min(wstart + kW, inputWidth + padW);
|
||||
int64_t pool_size = (hend - hstart) * (wend - wstart);
|
||||
hstart = std::max(hstart, (int64_t)0);
|
||||
wstart = std::max(wstart, (int64_t)0);
|
||||
hstart = std::max(hstart, static_cast<int64_t>(0));
|
||||
wstart = std::max(wstart, static_cast<int64_t>(0));
|
||||
hend = std::min(hend, inputHeight);
|
||||
wend = std::min(wend, inputWidth);
|
||||
|
||||
|
||||
@ -646,7 +646,7 @@ class QConvPackWeightInt8 final {
|
||||
torch::List<int64_t> output_padding;
|
||||
output_padding.reserve(kSpatialDim);
|
||||
for ([[maybe_unused]] const auto idx : c10::irange(kSpatialDim)) {
|
||||
output_padding.push_back((int64_t)0);
|
||||
output_padding.push_back(0);
|
||||
}
|
||||
return _run(weight, bias, stride, padding, output_padding, dilation, groups,
|
||||
/*transpose=*/false);
|
||||
|
||||
@ -301,6 +301,10 @@ def define_qnnpack(third_party, labels = []):
|
||||
"-DQNNP_PRIVATE=",
|
||||
"-DQNNP_INTERNAL=",
|
||||
],
|
||||
fbobjc_compiler_flags = [
|
||||
"-Wno-switch-enum",
|
||||
"-Wno-switch-default",
|
||||
],
|
||||
labels = [
|
||||
"supermodule:android/default/pytorch",
|
||||
"supermodule:ios/default/public.pytorch",
|
||||
|
||||
@ -134,7 +134,7 @@ class QConvPackWeightInt8Cudnn final {
|
||||
torch::List<int64_t> output_padding;
|
||||
output_padding.reserve(kSpatialDim);
|
||||
for ([[maybe_unused]] const auto idx : c10::irange(kSpatialDim)) {
|
||||
output_padding.push_back((int64_t)0);
|
||||
output_padding.push_back(0);
|
||||
}
|
||||
return _run(weight, bias, stride, padding, output_padding, dilation, groups,
|
||||
/*transpose=*/false);
|
||||
|
||||
@ -26,6 +26,8 @@
|
||||
#include <ATen/ops/erf_native.h>
|
||||
#include <ATen/ops/erfinv.h>
|
||||
#include <ATen/ops/erfinv_native.h>
|
||||
#include <ATen/ops/exp.h>
|
||||
#include <ATen/ops/exp_native.h>
|
||||
#include <ATen/ops/expm1.h>
|
||||
#include <ATen/ops/expm1_native.h>
|
||||
#include <ATen/ops/floor.h>
|
||||
@ -175,6 +177,7 @@ COALESCED_UNARY_UFUNC(atanh)
|
||||
COALESCED_UNARY_UFUNC(ceil)
|
||||
COALESCED_UNARY_UFUNC(deg2rad)
|
||||
COALESCED_UNARY_UFUNC(erf)
|
||||
COALESCED_UNARY_UFUNC(exp)
|
||||
COALESCED_UNARY_UFUNC(erfinv)
|
||||
COALESCED_UNARY_UFUNC(expm1)
|
||||
COALESCED_UNARY_UFUNC(floor)
|
||||
|
||||
@ -16,8 +16,8 @@ void Xcoo2csr(const int *coorowind, int64_t nnz, int64_t m, int *csrrowptr) {
|
||||
"cusparseXcoo2csr only supports m, nnz with the bound [val] <= ",
|
||||
INT_MAX);
|
||||
|
||||
int i_nnz = (int)nnz;
|
||||
int i_m = (int)m;
|
||||
int i_nnz = static_cast<int>(nnz);
|
||||
int i_m = static_cast<int>(m);
|
||||
|
||||
auto handle = at::cuda::getCurrentCUDASparseHandle();
|
||||
TORCH_CUDASPARSE_CHECK(cusparseXcoo2csr(handle, coorowind, i_nnz, i_m, csrrowptr, CUSPARSE_INDEX_BASE_ZERO));
|
||||
@ -202,7 +202,7 @@ void CreateIdentityPermutation(int64_t nnz, int *P) {
|
||||
TORCH_CHECK((nnz <= INT_MAX),
|
||||
"Xcsrsort_bufferSizeExt only supports m, n, nnz with the bound [val] <= ",
|
||||
INT_MAX);
|
||||
int i_nnz = (int)nnz;
|
||||
int i_nnz = static_cast<int>(nnz);
|
||||
|
||||
auto handle = at::cuda::getCurrentCUDASparseHandle();
|
||||
cusparseCreateIdentityPermutation(handle, i_nnz, P);
|
||||
@ -213,9 +213,9 @@ void Xcsrsort_bufferSizeExt(int64_t m, int64_t n, int64_t nnz, const int *csrRow
|
||||
TORCH_CHECK((m <= INT_MAX) && (n <= INT_MAX) && (nnz <= INT_MAX),
|
||||
"Xcsrsort_bufferSizeExt only supports m, n, nnz with the bound [val] <=",
|
||||
INT_MAX);
|
||||
int i_m = (int)m;
|
||||
int i_n = (int)n;
|
||||
int i_nnz = (int)nnz;
|
||||
int i_m = static_cast<int>(m);
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_nnz = static_cast<int>(nnz);
|
||||
|
||||
auto handle = at::cuda::getCurrentCUDASparseHandle();
|
||||
TORCH_CUDASPARSE_CHECK(cusparseXcsrsort_bufferSizeExt(handle, i_m, i_n, i_nnz, csrRowPtr, csrColInd, pBufferSizeInBytes));
|
||||
@ -226,9 +226,9 @@ void Xcsrsort(int64_t m, int64_t n, int64_t nnz, const int *csrRowPtr, int *csrC
|
||||
TORCH_CHECK((m <= INT_MAX) && (n <= INT_MAX) && (nnz <= INT_MAX),
|
||||
"Xcsrsort only supports m, n, nnz with the bound [val] <= ",
|
||||
INT_MAX);
|
||||
int i_m = (int)m;
|
||||
int i_n = (int)n;
|
||||
int i_nnz = (int)nnz;
|
||||
int i_m = static_cast<int>(m);
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_nnz = static_cast<int>(nnz);
|
||||
|
||||
auto handle = at::cuda::getCurrentCUDASparseHandle();
|
||||
cusparseMatDescr_t desc;
|
||||
@ -242,9 +242,9 @@ void Xcoosort_bufferSizeExt(int64_t m, int64_t n, int64_t nnz, const int *cooRow
|
||||
TORCH_CHECK((m <= INT_MAX) && (n <= INT_MAX) && (nnz <= INT_MAX),
|
||||
"Xcoosort_bufferSizeExt only supports m, n, nnz with the bound [val] <= ",
|
||||
INT_MAX);
|
||||
int i_m = (int)m;
|
||||
int i_n = (int)n;
|
||||
int i_nnz = (int)nnz;
|
||||
int i_m = static_cast<int>(m);
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_nnz = static_cast<int>(nnz);
|
||||
|
||||
auto handle = at::cuda::getCurrentCUDASparseHandle();
|
||||
TORCH_CUDASPARSE_CHECK(cusparseXcoosort_bufferSizeExt(handle, i_m, i_n, i_nnz, cooRows, cooCols, pBufferSizeInBytes));
|
||||
@ -255,9 +255,9 @@ void XcoosortByRow(int64_t m, int64_t n, int64_t nnz, int *cooRows, int *cooCols
|
||||
TORCH_CHECK((m <= INT_MAX) && (n <= INT_MAX) && (nnz <= INT_MAX),
|
||||
"XcoosortByRow only supports m, n, nnz with the bound [val] <= ",
|
||||
INT_MAX);
|
||||
int i_m = (int)m;
|
||||
int i_n = (int)n;
|
||||
int i_nnz = (int)nnz;
|
||||
int i_m = static_cast<int>(m);
|
||||
int i_n = static_cast<int>(n);
|
||||
int i_nnz = static_cast<int>(nnz);
|
||||
|
||||
auto handle = at::cuda::getCurrentCUDASparseHandle();
|
||||
TORCH_CUDASPARSE_CHECK(cusparseXcoosortByRow(handle, i_m, i_n, i_nnz, cooRows, cooCols, P, pBuffer));
|
||||
|
||||
@ -155,7 +155,7 @@ void set_params_fprop(Flash_fwd_params ¶ms,
|
||||
// [Minor] We want to round down since when we do the comparison we use <= instead of <
|
||||
// params.p_dropout_in_uint = uint32_t(std::floor(params.p_dropout * 4294967295.0));
|
||||
// params.p_dropout_in_uint16_t = uint16_t(std::floor(params.p_dropout * 65535.0));
|
||||
params.p_dropout_in_uint8_t = uint8_t(std::floor(params.p_dropout * 255.0));
|
||||
params.p_dropout_in_uint8_t = static_cast<uint8_t>(std::floor(params.p_dropout * 255.0));
|
||||
params.rp_dropout = 1.f / params.p_dropout;
|
||||
params.scale_softmax_rp_dropout = params.rp_dropout * params.scale_softmax;
|
||||
TORCH_CHECK(p_dropout < 1.f);
|
||||
@ -307,7 +307,7 @@ inline int num_splits_heuristic(int batch_nheads_mblocks, int num_SMs, int num_n
|
||||
if (!is_split_eligible(num_splits)) {
|
||||
efficiency.push_back(0.f);
|
||||
} else {
|
||||
float n_waves = float(batch_nheads_mblocks * num_splits) / num_SMs;
|
||||
float n_waves = static_cast<float>(batch_nheads_mblocks * num_splits) / num_SMs;
|
||||
float eff = n_waves / ceil(n_waves);
|
||||
// printf("num_splits = %d, eff = %f\n", num_splits, eff);
|
||||
if (eff > max_efficiency) { max_efficiency = eff; }
|
||||
|
||||
@ -341,7 +341,7 @@ inline bool check_grouped_query_attention(sdp_params const& params, bool debug)
|
||||
const auto v_num_heads = params.value.sym_size(-3);
|
||||
const bool same_kv_heads = k_num_heads == v_num_heads;
|
||||
|
||||
if (requires_same_num_heads && !(same_kv_heads)){
|
||||
if (requires_same_num_heads && !same_kv_heads){
|
||||
if (debug) {
|
||||
TORCH_WARN(
|
||||
"Both fused kernels require key and value to have the same num_heads and batch_size but got: ",
|
||||
|
||||
@ -202,6 +202,7 @@ supported:
|
||||
- select_backward
|
||||
- _trilinear
|
||||
- linalg_pinv.atol_rtol_tensor
|
||||
- svd
|
||||
- logsumexp.out
|
||||
symint:
|
||||
- empty.memory_format
|
||||
|
||||
@ -43,9 +43,9 @@ bool available(
|
||||
(kFloat == weight.scalar_type()) &&
|
||||
// Bias
|
||||
(bias_sizes_opt.has_value() ? ((1 == bias_sizes_opt->size()) &&
|
||||
((transposed ? (weight.size(Layout::Filter::input) ==
|
||||
(transposed ? (weight.size(Layout::Filter::input) ==
|
||||
((*bias_sizes_opt)[0] / groups))
|
||||
: (weight.size(Layout::Filter::output) == ((*bias_sizes_opt)[0])))))
|
||||
: (weight.size(Layout::Filter::output) == ((*bias_sizes_opt)[0]))))
|
||||
: true) &&
|
||||
// Padding
|
||||
(padding[Layout::Parameter::height] >= 0) &&
|
||||
@ -133,10 +133,10 @@ const Tensor reorder_weights_for_transpose_conv(const Tensor& weight_nhwc,
|
||||
int kernel_height = weight_nhwc.size(2);
|
||||
|
||||
int o_offset = 1;
|
||||
int h_offset = (output_channels_per_group);
|
||||
int w_offset = (output_channels_per_group)*(kernel_height);
|
||||
int i_offset = (output_channels_per_group)*(kernel_height)*(kernel_width);
|
||||
int g_offset = (output_channels_per_group)*(kernel_height)*(kernel_width)*(input_channels_per_group);
|
||||
int h_offset = output_channels_per_group;
|
||||
int w_offset = output_channels_per_group*kernel_height;
|
||||
int i_offset = output_channels_per_group*kernel_height*kernel_width;
|
||||
int g_offset = output_channels_per_group*kernel_height*kernel_width*input_channels_per_group;
|
||||
|
||||
Tensor reordered = mobile::empty_with_tail_padding(
|
||||
weight_nhwc.sizes(),
|
||||
|
||||
@ -28,6 +28,7 @@
|
||||
#include <c10/util/OptionalArrayRef.h>
|
||||
#include <c10/util/intrusive_ptr.h>
|
||||
#include <c10/macros/Export.h>
|
||||
#include <c10/macros/Macros.h>
|
||||
#include <ATen/core/CheckMemoryFormat.h>
|
||||
#include <ATen/core/DeprecatedTypePropertiesRegistry.h>
|
||||
#include <ATen/core/DeprecatedTypeProperties.h>
|
||||
@ -129,6 +130,7 @@ class TORCH_API Tensor: public TensorBase {
|
||||
return *this;
|
||||
}
|
||||
|
||||
C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum")
|
||||
switch (this->layout()) {
|
||||
case at::kSparse:
|
||||
case at::kSparseCsr:
|
||||
@ -139,6 +141,7 @@ class TORCH_API Tensor: public TensorBase {
|
||||
default:
|
||||
return this->_conj();
|
||||
}
|
||||
C10_DIAGNOSTIC_POP()
|
||||
}
|
||||
|
||||
// Aliased by Dimname overloads, so need explicit using
|
||||
|
||||
@ -3,10 +3,13 @@
|
||||
#include <c10/core/DeviceType.h>
|
||||
#include <c10/core/DispatchKey.h>
|
||||
#include <c10/core/DispatchKeySet.h>
|
||||
#include <c10/macros/Macros.h>
|
||||
#include <c10/util/Exception.h>
|
||||
|
||||
#include <stdexcept>
|
||||
|
||||
C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum")
|
||||
|
||||
namespace c10 {
|
||||
|
||||
/**
|
||||
@ -402,3 +405,5 @@ inline bool isSparseCsr(Backend b) {
|
||||
}
|
||||
|
||||
} // namespace c10
|
||||
|
||||
C10_DIAGNOSTIC_POP()
|
||||
|
||||
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user