Remove asm_dialect=AD_ATT from all llvm.inline_asm calls (Dao-AILab#98)

Remove asm_dialect from all inline asm calls, not just the braced ones.
PTX is neither AT&T nor Intel syntax — the default (no dialect) passes
the asm string through to ptxas unmodified, which is correct.

Contributed by Second Nature Computing (https://joinsecondnature.com)

Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>

Author: blake-snc

quack/activation.py

@@ -30,7 +30,6 @@ def tanh(a: float | Float32, *, loc=None, ip=None) -> Float32:
             "=f,f",
             has_side_effects=False,
             is_align_stack=False,
-            asm_dialect=llvm.AsmDialect.AD_ATT,
         )
     )
 

quack/copy_utils.py

@@ -584,7 +584,6 @@ def cpasync_reduce_bulk_add_f32(
         # "l,r,r,l",
         has_side_effects=True,
         is_align_stack=False,
-        asm_dialect=llvm.AsmDialect.AD_ATT,
     )
 
 
@@ -715,7 +714,6 @@ def tma_gather4_load(
         "r,l,r,r,r,r,r,r",  # constraints: register, long, 6x register
         has_side_effects=True,
         is_align_stack=False,
-        asm_dialect=llvm.AsmDialect.AD_ATT,
         loc=loc,
         ip=ip,
     )

quack/rounding.py

@@ -52,7 +52,6 @@ def mul_wide_u32(a: Uint32, b: Uint32, *, loc=None, ip=None) -> tuple:
         "=r,=r,r,r",
         has_side_effects=False,
         is_align_stack=False,
-        asm_dialect=llvm.AsmDialect.AD_ATT,
     )
     i32_ty = T.i32()
     hi = cutlass.Uint32(llvm.extractvalue(i32_ty, result, [0], loc=loc, ip=ip))
@@ -85,7 +84,6 @@ def cvt_f32x2_bf16x2_rs(
             "=r,f,f,r",
             has_side_effects=False,
             is_align_stack=False,
-            asm_dialect=llvm.AsmDialect.AD_ATT,
         )
     )
 

quack/tensormap_manager.py

@@ -87,7 +87,6 @@ def update_tensormap_shape(
                         "r,r",
                         has_side_effects=True,
                         is_align_stack=False,
-                        asm_dialect=llvm.AsmDialect.AD_ATT,
                     )
             # wait until it's safe to update tensormap in global memory
             with cute.arch.elect_one():
@@ -109,7 +108,6 @@ def update_tensormap_shape(
                         "l,r",
                         has_side_effects=True,
                         is_align_stack=False,
-                        asm_dialect=llvm.AsmDialect.AD_ATT,
                     )
                 cute.arch.sync_warp()
                 cute.nvgpu.cpasync.fence_tma_desc_release()

quack/utils.py

@@ -39,7 +39,6 @@ def set_block_rank(
             "=r,r,r",
             has_side_effects=False,
             is_align_stack=False,
-            asm_dialect=llvm.AsmDialect.AD_ATT,
         )
     )
 
@@ -72,7 +71,6 @@ def store_shared_remote(
         f"r,{constraint},r",
         has_side_effects=True,
         is_align_stack=False,
-        asm_dialect=llvm.AsmDialect.AD_ATT,
     )
 
 
@@ -120,7 +118,6 @@ def store_shared_remote_x4(
         f"r,r,{constraint},{constraint},{constraint},{constraint}",
         has_side_effects=True,
         is_align_stack=False,
-        asm_dialect=llvm.AsmDialect.AD_ATT,
     )
 
 
@@ -156,7 +153,6 @@ def sqrt(a: float | Float32, *, loc=None, ip=None) -> Float32:
             "=f,f",
             has_side_effects=False,
             is_align_stack=False,
-            asm_dialect=llvm.AsmDialect.AD_ATT,
         )
     )
 
@@ -171,7 +167,6 @@ def ceil(a: float | Float32, *, loc=None, ip=None) -> Int32:
             "=r,f",
             has_side_effects=False,
             is_align_stack=False,
-            asm_dialect=llvm.AsmDialect.AD_ATT,
         )
     )