Environment
- TileOP API:
a795b973020de012fa2a8e2c29b26adfc83d5d29
- llvm-project:
611105f2be11fab9a8ef20bd02b740f2c5d786b3
- Reproducer consumer: SuperNPUBench
ops-20260823 (244d088fdc56dc0b41cc77b60e03dc974b67ea93)
Problem
Compiling any normalization testcase with its documented datatype define:
clang++ ... -DDType=__half ...
fails while parsing the TQUANT inline-assembly templates in include/jcore/template_asm.hpp:
error: unknown symbolic operand name in inline assembly string
"B.DATR %c[DType], RNONE, SAT\n"
^
The same error is emitted for the four B.DATR variants around lines 6539, 6557, 6575, and 6594.
Root cause
The inline asm uses DType as a symbolic operand label:
"B.DATR %c[DType], RNONE, SAT\n"
...
[DType] "i"(type_traits<typename tile_shape_out::DType>::TypeCode)
The consumer's public DType preprocessor macro rewrites the constraint label token [DType] to [__half]. Macro expansion does not occur inside the asm string, which still references %c[DType], so Clang correctly reports that the symbolic operand is missing.
This happens even when the consumer does not instantiate TQUANT because the template definition is parsed after preprocessing.
Verified minimal fix
Rename the asm operand label without changing the encoded value:
- "B.DATR %c[DType], RNONE, SAT\n"
- [DType] "i"(...)
+ "B.DATR %c[DstType], RNONE, SAT\n"
+ [DstType] "i"(...)
Applying this rename to all four variants makes all 10 default normalization ELFs compile successfully. No instruction encoding or runtime behavior changes.
Expected behavior
TileOP public headers should avoid generic symbolic asm labels that can collide with common consumer macros.
Suggested acceptance test
Compile a small TU that includes common/pto_tileop.hpp with -DDType=__half. The header should parse successfully, including when TQUANT is not instantiated.
Environment
a795b973020de012fa2a8e2c29b26adfc83d5d29611105f2be11fab9a8ef20bd02b740f2c5d786b3ops-20260823(244d088fdc56dc0b41cc77b60e03dc974b67ea93)Problem
Compiling any normalization testcase with its documented datatype define:
fails while parsing the TQUANT inline-assembly templates in
include/jcore/template_asm.hpp:The same error is emitted for the four B.DATR variants around lines 6539, 6557, 6575, and 6594.
Root cause
The inline asm uses
DTypeas a symbolic operand label:The consumer's public
DTypepreprocessor macro rewrites the constraint label token[DType]to[__half]. Macro expansion does not occur inside the asm string, which still references%c[DType], so Clang correctly reports that the symbolic operand is missing.This happens even when the consumer does not instantiate TQUANT because the template definition is parsed after preprocessing.
Verified minimal fix
Rename the asm operand label without changing the encoded value:
Applying this rename to all four variants makes all 10 default normalization ELFs compile successfully. No instruction encoding or runtime behavior changes.
Expected behavior
TileOP public headers should avoid generic symbolic asm labels that can collide with common consumer macros.
Suggested acceptance test
Compile a small TU that includes
common/pto_tileop.hppwith-DDType=__half. The header should parse successfully, including when TQUANT is not instantiated.