"}}"),
[(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>;
-def INT_BAR_SYNC : NVPTXInst<(outs), (ins i32imm:$i), "bar.sync\t$i;",
+def INT_BAR_SYNC : NVPTXInst<(outs), (ins i32imm:$i), "bar.sync \t$i;",
[(int_nvvm_bar_sync imm:$i)]>;
// shfl.{up,down,bfly,idx}.b32
// Special register reads
def MOV_SPECIAL : NVPTXInst<(outs Int32Regs:$d),
(ins SpecialRegs:$r),
- "mov.b32\t$d, $r;", []>;
+ "mov.b32 \t$d, $r;", []>;
def : Pat<(int_nvvm_read_ptx_sreg_envreg0), (MOV_SPECIAL ENVREG0)>;
def : Pat<(int_nvvm_read_ptx_sreg_envreg1), (MOV_SPECIAL ENVREG1)>;
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
- "tex.1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
+ "tex.1d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
[]>;
def TEX_1D_F32_F32
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x),
- "tex.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
+ "tex.1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
[]>;
def TEX_1D_F32_F32_LEVEL
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$lod),
- "tex.level.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x\\}], $lod;",
[]>;
def TEX_1D_F32_F32_GRAD
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
def TEX_1D_S32_S32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
- "tex.1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
+ "tex.1d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
[]>;
def TEX_1D_S32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x),
- "tex.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
+ "tex.1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
[]>;
def TEX_1D_S32_F32_LEVEL
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x\\}], $lod;",
[]>;
def TEX_1D_S32_F32_GRAD
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
def TEX_1D_U32_S32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
- "tex.1d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
+ "tex.1d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
[]>;
def TEX_1D_U32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x),
- "tex.1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
+ "tex.1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
[]>;
def TEX_1D_U32_F32_LEVEL
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x\\}], $lod;",
[]>;
def TEX_1D_U32_F32_GRAD
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
- "tex.a1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}];",
[]>;
def TEX_1D_ARRAY_F32_F32
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x),
- "tex.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}];",
[]>;
def TEX_1D_ARRAY_F32_F32_LEVEL
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}], $lod;",
[]>;
def TEX_1D_ARRAY_F32_F32_GRAD
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
def TEX_1D_ARRAY_S32_S32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
- "tex.a1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}];",
[]>;
def TEX_1D_ARRAY_S32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x),
- "tex.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}];",
[]>;
def TEX_1D_ARRAY_S32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}], $lod;",
[]>;
def TEX_1D_ARRAY_S32_F32_GRAD
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
def TEX_1D_ARRAY_U32_S32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
- "tex.a1d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}];",
[]>;
def TEX_1D_ARRAY_U32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x),
- "tex.a1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}];",
[]>;
def TEX_1D_ARRAY_U32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.a1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}], $lod;",
[]>;
def TEX_1D_ARRAY_U32_F32_GRAD
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.a1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
- "tex.2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TEX_2D_F32_F32
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tex.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TEX_2D_F32_F32_LEVEL
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$lod),
- "tex.level.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}], $lod;",
[]>;
def TEX_2D_F32_F32_GRAD
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
- "tex.2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TEX_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tex.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TEX_2D_S32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$lod),
- "tex.level.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}], $lod;",
[]>;
def TEX_2D_S32_F32_GRAD
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
- "tex.2d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TEX_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tex.2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TEX_2D_U32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$lod),
- "tex.level.2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}], $lod;",
[]>;
def TEX_2D_U32_F32_GRAD
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
Int32Regs:$y),
- "tex.a2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_2D_ARRAY_F32_F32
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y),
- "tex.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_2D_ARRAY_F32_F32_LEVEL
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y, Float32Regs:$lod),
- "tex.level.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}], $lod;",
[]>;
def TEX_2D_ARRAY_F32_F32_GRAD
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y, Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
Int32Regs:$y),
- "tex.a2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_2D_ARRAY_S32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y),
- "tex.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_2D_ARRAY_S32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y, Float32Regs:$lod),
- "tex.level.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}], $lod;",
[]>;
def TEX_2D_ARRAY_S32_F32_GRAD
Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
Int32Regs:$y),
- "tex.a2d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_2D_ARRAY_U32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y),
- "tex.a2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_2D_ARRAY_U32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y, Float32Regs:$lod),
- "tex.level.a2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}], $lod;",
[]>;
def TEX_2D_ARRAY_U32_F32_GRAD
Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.a2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
Int32Regs:$z),
- "tex.3d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_3D_F32_F32
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z),
- "tex.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_3D_F32_F32_LEVEL
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z, Float32Regs:$lod),
- "tex.level.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.3d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_3D_F32_F32_GRAD
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$gradx2, Float32Regs:$grady0,
Float32Regs:$grady1, Float32Regs:$grady2),
- "tex.grad.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.3d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}], "
"\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
"\\{$grady0, $grady1, $grady2, $grady2\\};",
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
Int32Regs:$z),
- "tex.3d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_3D_S32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z),
- "tex.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_3D_S32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z, Float32Regs:$lod),
- "tex.level.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.3d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_3D_S32_F32_GRAD
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$gradx2, Float32Regs:$grady0,
Float32Regs:$grady1, Float32Regs:$grady2),
- "tex.grad.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.3d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}], "
"\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
"\\{$grady0, $grady1, $grady2, $grady2\\};",
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
Int32Regs:$z),
- "tex.3d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_3D_U32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z),
- "tex.3d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_3D_U32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z, Float32Regs:$lod),
- "tex.level.3d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.3d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_3D_U32_F32_GRAD
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$gradx2, Float32Regs:$grady0,
Float32Regs:$grady1, Float32Regs:$grady2),
- "tex.grad.3d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.3d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}], "
"\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
"\\{$grady0, $grady1, $grady2, $grady2\\};",
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.cube.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.cube.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_CUBE_F32_F32_LEVEL
(ins Int64Regs:$t, Int64Regs:$s,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.cube.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.cube.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_CUBE_S32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.cube.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.cube.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_CUBE_S32_F32_LEVEL
(ins Int64Regs:$t, Int64Regs:$s,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.cube.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.cube.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_CUBE_U32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.cube.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.cube.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_CUBE_U32_F32_LEVEL
(ins Int64Regs:$t, Int64Regs:$s,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.cube.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.cube.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.acube.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.acube.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $z\\}];",
[]>;
def TEX_CUBE_ARRAY_F32_F32_LEVEL
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.acube.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.acube.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $z\\}], $lod;",
[]>;
def TEX_CUBE_ARRAY_S32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.acube.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.acube.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $z\\}];",
[]>;
def TEX_CUBE_ARRAY_S32_F32_LEVEL
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.acube.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.acube.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $z\\}], $lod;",
[]>;
def TEX_CUBE_ARRAY_U32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.acube.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.acube.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $z\\}];",
[]>;
def TEX_CUBE_ARRAY_U32_F32_LEVEL
(ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.acube.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.acube.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, $s, \\{$l, $x, $y, $z\\}], $lod;",
[]>;
: NVPTXInst<(outs Float32Regs:$v0, Float32Regs:$v1,
Float32Regs:$v2, Float32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.r.2d.v4.f32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.r.2d.v4.f32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_G_2D_F32_F32
: NVPTXInst<(outs Float32Regs:$v0, Float32Regs:$v1,
Float32Regs:$v2, Float32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.g.2d.v4.f32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.g.2d.v4.f32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_B_2D_F32_F32
: NVPTXInst<(outs Float32Regs:$v0, Float32Regs:$v1,
Float32Regs:$v2, Float32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.b.2d.v4.f32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.b.2d.v4.f32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_A_2D_F32_F32
: NVPTXInst<(outs Float32Regs:$v0, Float32Regs:$v1,
Float32Regs:$v2, Float32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.a.2d.v4.f32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.a.2d.v4.f32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_R_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.r.2d.v4.s32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.r.2d.v4.s32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_G_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.g.2d.v4.s32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.g.2d.v4.s32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_B_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.b.2d.v4.s32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.b.2d.v4.s32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_A_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.a.2d.v4.s32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.a.2d.v4.s32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_R_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.r.2d.v4.u32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.r.2d.v4.u32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_G_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.g.2d.v4.u32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.g.2d.v4.u32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_B_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.b.2d.v4.u32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.b.2d.v4.u32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
def TLD4_A_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
- "tld4.a.2d.v4.u32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.a.2d.v4.u32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, $s, \\{$x, $y\\}];",
[]>;
}
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$x),
- "tex.1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
+ "tex.1d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
[]>;
def TEX_UNIFIED_1D_F32_F32
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x),
- "tex.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
+ "tex.1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
[]>;
def TEX_UNIFIED_1D_F32_F32_LEVEL
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$lod),
- "tex.level.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x\\}], $lod;",
[]>;
def TEX_UNIFIED_1D_F32_F32_GRAD
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
def TEX_UNIFIED_1D_S32_S32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$x),
- "tex.1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
+ "tex.1d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
[]>;
def TEX_UNIFIED_1D_S32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x),
- "tex.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
+ "tex.1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
[]>;
def TEX_UNIFIED_1D_S32_F32_LEVEL
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x\\}], $lod;",
[]>;
def TEX_UNIFIED_1D_S32_F32_GRAD
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
def TEX_UNIFIED_1D_U32_S32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$x),
- "tex.1d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
+ "tex.1d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
[]>;
def TEX_UNIFIED_1D_U32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x),
- "tex.1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
+ "tex.1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
[]>;
def TEX_UNIFIED_1D_U32_F32_LEVEL
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x\\}], $lod;",
[]>;
def TEX_UNIFIED_1D_U32_F32_GRAD
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Int32Regs:$x),
- "tex.a1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}];",
[]>;
def TEX_UNIFIED_1D_ARRAY_F32_F32
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x),
- "tex.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}];",
[]>;
def TEX_UNIFIED_1D_ARRAY_F32_F32_LEVEL
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}], $lod;",
[]>;
def TEX_UNIFIED_1D_ARRAY_F32_F32_GRAD
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a1d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
def TEX_UNIFIED_1D_ARRAY_S32_S32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Int32Regs:$x),
- "tex.a1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}];",
[]>;
def TEX_UNIFIED_1D_ARRAY_S32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x),
- "tex.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}];",
[]>;
def TEX_UNIFIED_1D_ARRAY_S32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}], $lod;",
[]>;
def TEX_UNIFIED_1D_ARRAY_S32_F32_GRAD
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a1d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
def TEX_UNIFIED_1D_ARRAY_U32_S32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Int32Regs:$x),
- "tex.a1d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}];",
[]>;
def TEX_UNIFIED_1D_ARRAY_U32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x),
- "tex.a1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}];",
[]>;
def TEX_UNIFIED_1D_ARRAY_U32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$lod),
- "tex.level.a1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}], $lod;",
[]>;
def TEX_UNIFIED_1D_ARRAY_U32_F32_GRAD
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$gradx, Float32Regs:$grady),
- "tex.grad.a1d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a1d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
[]>;
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$x, Int32Regs:$y),
- "tex.2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TEX_UNIFIED_2D_F32_F32
: NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tex.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TEX_UNIFIED_2D_F32_F32_LEVEL
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$lod),
- "tex.level.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}], $lod;",
[]>;
def TEX_UNIFIED_2D_F32_F32_GRAD
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$x, Int32Regs:$y),
- "tex.2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TEX_UNIFIED_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tex.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TEX_UNIFIED_2D_S32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$lod),
- "tex.level.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}], $lod;",
[]>;
def TEX_UNIFIED_2D_S32_F32_GRAD
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$x, Int32Regs:$y),
- "tex.2d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TEX_UNIFIED_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tex.2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TEX_UNIFIED_2D_U32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$lod),
- "tex.level.2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}], $lod;",
[]>;
def TEX_UNIFIED_2D_U32_F32_GRAD
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Int32Regs:$x,
Int32Regs:$y),
- "tex.a2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_UNIFIED_2D_ARRAY_F32_F32
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y),
- "tex.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_UNIFIED_2D_ARRAY_F32_F32_LEVEL
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y, Float32Regs:$lod),
- "tex.level.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}], $lod;",
[]>;
def TEX_UNIFIED_2D_ARRAY_F32_F32_GRAD
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y, Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a2d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Int32Regs:$x,
Int32Regs:$y),
- "tex.a2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_UNIFIED_2D_ARRAY_S32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y),
- "tex.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_UNIFIED_2D_ARRAY_S32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y, Float32Regs:$lod),
- "tex.level.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}], $lod;",
[]>;
def TEX_UNIFIED_2D_ARRAY_S32_F32_GRAD
Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a2d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Int32Regs:$x,
Int32Regs:$y),
- "tex.a2d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_UNIFIED_2D_ARRAY_U32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y),
- "tex.a2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.a2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}];",
[]>;
def TEX_UNIFIED_2D_ARRAY_U32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l, Float32Regs:$x,
Float32Regs:$y, Float32Regs:$lod),
- "tex.level.a2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.a2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}], $lod;",
[]>;
def TEX_UNIFIED_2D_ARRAY_U32_F32_GRAD
Float32Regs:$y,
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$grady0, Float32Regs:$grady1),
- "tex.grad.a2d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.a2d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
"\\{$grady0, $grady1\\};",
[]>;
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$x, Int32Regs:$y,
Int32Regs:$z),
- "tex.3d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_UNIFIED_3D_F32_F32
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z),
- "tex.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_UNIFIED_3D_F32_F32_LEVEL
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z, Float32Regs:$lod),
- "tex.level.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.3d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_UNIFIED_3D_F32_F32_GRAD
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$gradx2, Float32Regs:$grady0,
Float32Regs:$grady1, Float32Regs:$grady2),
- "tex.grad.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.3d.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}], "
"\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
"\\{$grady0, $grady1, $grady2, $grady2\\};",
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$x, Int32Regs:$y,
Int32Regs:$z),
- "tex.3d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.s32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_UNIFIED_3D_S32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z),
- "tex.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_UNIFIED_3D_S32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z, Float32Regs:$lod),
- "tex.level.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.3d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_UNIFIED_3D_S32_F32_GRAD
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$gradx2, Float32Regs:$grady0,
Float32Regs:$grady1, Float32Regs:$grady2),
- "tex.grad.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.3d.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}], "
"\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
"\\{$grady0, $grady1, $grady2, $grady2\\};",
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$x, Int32Regs:$y,
Int32Regs:$z),
- "tex.3d.v4.u32.s32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.u32.s32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_UNIFIED_3D_U32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z),
- "tex.3d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.3d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_UNIFIED_3D_U32_F32_LEVEL
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y,
Float32Regs:$z, Float32Regs:$lod),
- "tex.level.3d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.3d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_UNIFIED_3D_U32_F32_GRAD
Float32Regs:$gradx0, Float32Regs:$gradx1,
Float32Regs:$gradx2, Float32Regs:$grady0,
Float32Regs:$grady1, Float32Regs:$grady2),
- "tex.grad.3d.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.grad.3d.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}], "
"\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
"\\{$grady0, $grady1, $grady2, $grady2\\};",
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.cube.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.cube.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_UNIFIED_CUBE_F32_F32_LEVEL
(ins Int64Regs:$t,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.cube.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.cube.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_UNIFIED_CUBE_S32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.cube.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.cube.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_UNIFIED_CUBE_S32_F32_LEVEL
(ins Int64Regs:$t,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.cube.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.cube.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
def TEX_UNIFIED_CUBE_U32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.cube.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.cube.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}];",
[]>;
def TEX_UNIFIED_CUBE_U32_F32_LEVEL
(ins Int64Regs:$t,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.cube.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.cube.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$x, $y, $z, $z\\}], $lod;",
[]>;
Float32Regs:$b, Float32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.acube.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.acube.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $z\\}];",
[]>;
def TEX_UNIFIED_CUBE_ARRAY_F32_F32_LEVEL
(ins Int64Regs:$t, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.acube.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.acube.v4.f32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $z\\}], $lod;",
[]>;
def TEX_UNIFIED_CUBE_ARRAY_S32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.acube.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.acube.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $z\\}];",
[]>;
def TEX_UNIFIED_CUBE_ARRAY_S32_F32_LEVEL
(ins Int64Regs:$t, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.acube.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.acube.v4.s32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $z\\}], $lod;",
[]>;
def TEX_UNIFIED_CUBE_ARRAY_U32_F32
Int32Regs:$b, Int32Regs:$a),
(ins Int64Regs:$t, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z),
- "tex.acube.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.acube.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $z\\}];",
[]>;
def TEX_UNIFIED_CUBE_ARRAY_U32_F32_LEVEL
(ins Int64Regs:$t, Int32Regs:$l,
Float32Regs:$x, Float32Regs:$y, Float32Regs:$z,
Float32Regs:$lod),
- "tex.level.acube.v4.u32.f32\t\\{$r, $g, $b, $a\\}, "
+ "tex.level.acube.v4.u32.f32 \t\\{$r, $g, $b, $a\\}, "
"[$t, \\{$l, $x, $y, $z\\}], $lod;",
[]>;
: NVPTXInst<(outs Float32Regs:$v0, Float32Regs:$v1,
Float32Regs:$v2, Float32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.r.2d.v4.f32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.r.2d.v4.f32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_G_2D_F32_F32
: NVPTXInst<(outs Float32Regs:$v0, Float32Regs:$v1,
Float32Regs:$v2, Float32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.g.2d.v4.f32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.g.2d.v4.f32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_B_2D_F32_F32
: NVPTXInst<(outs Float32Regs:$v0, Float32Regs:$v1,
Float32Regs:$v2, Float32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.b.2d.v4.f32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.b.2d.v4.f32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_A_2D_F32_F32
: NVPTXInst<(outs Float32Regs:$v0, Float32Regs:$v1,
Float32Regs:$v2, Float32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.a.2d.v4.f32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.a.2d.v4.f32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_R_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.r.2d.v4.s32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.r.2d.v4.s32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_G_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.g.2d.v4.s32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.g.2d.v4.s32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_B_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.b.2d.v4.s32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.b.2d.v4.s32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_A_2D_S32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.a.2d.v4.s32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.a.2d.v4.s32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_R_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.r.2d.v4.u32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.r.2d.v4.u32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_G_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.g.2d.v4.u32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.g.2d.v4.u32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_B_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.b.2d.v4.u32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.b.2d.v4.u32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
def TLD4_UNIFIED_A_2D_U32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
Int32Regs:$v2, Int32Regs:$v3),
(ins Int64Regs:$t, Float32Regs:$x, Float32Regs:$y),
- "tld4.a.2d.v4.u32.f32\t\\{$v0, $v1, $v2, $v3\\}, "
+ "tld4.a.2d.v4.u32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
"[$t, \\{$x, $y\\}];",
[]>;
}