[llvm] r292306 - [NVPTX] Standardize asm printer on "foo \tbar".
Justin Lebar via llvm-commits
llvm-commits at lists.llvm.org
Tue Jan 17 16:09:36 PST 2017
Author: jlebar
Date: Tue Jan 17 18:09:36 2017
New Revision: 292306
URL: http://llvm.org/viewvc/llvm-project?rev=292306&view=rev
Log:
[NVPTX] Standardize asm printer on "foo \tbar".
Some instructions were printed as "foo\tbar", but most are printed as
"foo \bar". Standardize on the latter form.
Modified:
llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td?rev=292306&r1=292305&r2=292306&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td Tue Jan 17 18:09:36 2017
@@ -389,57 +389,57 @@ let hasSideEffects = 0 in {
NVPTXInst<(outs RC:$dst),
(ins Int16Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".s8\t$dst, $src;"), []>;
+ FromName, ".s8 \t$dst, $src;"), []>;
def _u8 :
NVPTXInst<(outs RC:$dst),
(ins Int16Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".u8\t$dst, $src;"), []>;
+ FromName, ".u8 \t$dst, $src;"), []>;
def _s16 :
NVPTXInst<(outs RC:$dst),
(ins Int16Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".s16\t$dst, $src;"), []>;
+ FromName, ".s16 \t$dst, $src;"), []>;
def _u16 :
NVPTXInst<(outs RC:$dst),
(ins Int16Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".u16\t$dst, $src;"), []>;
+ FromName, ".u16 \t$dst, $src;"), []>;
def _s32 :
NVPTXInst<(outs RC:$dst),
(ins Int32Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".s32\t$dst, $src;"), []>;
+ FromName, ".s32 \t$dst, $src;"), []>;
def _u32 :
NVPTXInst<(outs RC:$dst),
(ins Int32Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".u32\t$dst, $src;"), []>;
+ FromName, ".u32 \t$dst, $src;"), []>;
def _s64 :
NVPTXInst<(outs RC:$dst),
(ins Int64Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".s64\t$dst, $src;"), []>;
+ FromName, ".s64 \t$dst, $src;"), []>;
def _u64 :
NVPTXInst<(outs RC:$dst),
(ins Int64Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".u64\t$dst, $src;"), []>;
+ FromName, ".u64 \t$dst, $src;"), []>;
def _f16 :
NVPTXInst<(outs RC:$dst),
(ins Float16Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".f16\t$dst, $src;"), []>;
+ FromName, ".f16 \t$dst, $src;"), []>;
def _f32 :
NVPTXInst<(outs RC:$dst),
(ins Float32Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".f32\t$dst, $src;"), []>;
+ FromName, ".f32 \t$dst, $src;"), []>;
def _f64 :
NVPTXInst<(outs RC:$dst),
(ins Float64Regs:$src, CvtMode:$mode),
!strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
- FromName, ".f64\t$dst, $src;"), []>;
+ FromName, ".f64 \t$dst, $src;"), []>;
}
// Generate cvts from all types to all types.
@@ -1373,15 +1373,15 @@ let hasSideEffects = 0 in {
def rr :
NVPTXInst<(outs Int1Regs:$dst), (ins RC:$a, RC:$b, CmpMode:$cmp),
!strconcat("setp${cmp:base}${cmp:ftz}.", TypeStr,
- "\t$dst, $a, $b;"), []>;
+ " \t$dst, $a, $b;"), []>;
def ri :
NVPTXInst<(outs Int1Regs:$dst), (ins RC:$a, ImmCls:$b, CmpMode:$cmp),
!strconcat("setp${cmp:base}${cmp:ftz}.", TypeStr,
- "\t$dst, $a, $b;"), []>;
+ " \t$dst, $a, $b;"), []>;
def ir :
NVPTXInst<(outs Int1Regs:$dst), (ins ImmCls:$a, RC:$b, CmpMode:$cmp),
!strconcat("setp${cmp:base}${cmp:ftz}.", TypeStr,
- "\t$dst, $a, $b;"), []>;
+ " \t$dst, $a, $b;"), []>;
}
}
@@ -1410,13 +1410,13 @@ let hasSideEffects = 0 in {
multiclass SET<string TypeStr, RegisterClass RC, Operand ImmCls> {
def rr : NVPTXInst<(outs Int32Regs:$dst),
(ins RC:$a, RC:$b, CmpMode:$cmp),
- !strconcat("set$cmp.", TypeStr, "\t$dst, $a, $b;"), []>;
+ !strconcat("set$cmp.", TypeStr, " \t$dst, $a, $b;"), []>;
def ri : NVPTXInst<(outs Int32Regs:$dst),
(ins RC:$a, ImmCls:$b, CmpMode:$cmp),
- !strconcat("set$cmp.", TypeStr, "\t$dst, $a, $b;"), []>;
+ !strconcat("set$cmp.", TypeStr, " \t$dst, $a, $b;"), []>;
def ir : NVPTXInst<(outs Int32Regs:$dst),
(ins ImmCls:$a, RC:$b, CmpMode:$cmp),
- !strconcat("set$cmp.", TypeStr, "\t$dst, $a, $b;"), []>;
+ !strconcat("set$cmp.", TypeStr, " \t$dst, $a, $b;"), []>;
}
}
@@ -1445,16 +1445,16 @@ let hasSideEffects = 0 in {
multiclass SELP<string TypeStr, RegisterClass RC, Operand ImmCls> {
def rr : NVPTXInst<(outs RC:$dst),
(ins RC:$a, RC:$b, Int1Regs:$p),
- !strconcat("selp.", TypeStr, "\t$dst, $a, $b, $p;"), []>;
+ !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"), []>;
def ri : NVPTXInst<(outs RC:$dst),
(ins RC:$a, ImmCls:$b, Int1Regs:$p),
- !strconcat("selp.", TypeStr, "\t$dst, $a, $b, $p;"), []>;
+ !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"), []>;
def ir : NVPTXInst<(outs RC:$dst),
(ins ImmCls:$a, RC:$b, Int1Regs:$p),
- !strconcat("selp.", TypeStr, "\t$dst, $a, $b, $p;"), []>;
+ !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"), []>;
def ii : NVPTXInst<(outs RC:$dst),
(ins ImmCls:$a, ImmCls:$b, Int1Regs:$p),
- !strconcat("selp.", TypeStr, "\t$dst, $a, $b, $p;"), []>;
+ !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"), []>;
}
multiclass SELP_PATTERN<string TypeStr, RegisterClass RC, Operand ImmCls,
@@ -1462,22 +1462,22 @@ let hasSideEffects = 0 in {
def rr :
NVPTXInst<(outs RC:$dst),
(ins RC:$a, RC:$b, Int1Regs:$p),
- !strconcat("selp.", TypeStr, "\t$dst, $a, $b, $p;"),
+ !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"),
[(set RC:$dst, (select Int1Regs:$p, RC:$a, RC:$b))]>;
def ri :
NVPTXInst<(outs RC:$dst),
(ins RC:$a, ImmCls:$b, Int1Regs:$p),
- !strconcat("selp.", TypeStr, "\t$dst, $a, $b, $p;"),
+ !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"),
[(set RC:$dst, (select Int1Regs:$p, RC:$a, ImmNode:$b))]>;
def ir :
NVPTXInst<(outs RC:$dst),
(ins ImmCls:$a, RC:$b, Int1Regs:$p),
- !strconcat("selp.", TypeStr, "\t$dst, $a, $b, $p;"),
+ !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"),
[(set RC:$dst, (select Int1Regs:$p, ImmNode:$a, RC:$b))]>;
def ii :
NVPTXInst<(outs RC:$dst),
(ins ImmCls:$a, ImmCls:$b, Int1Regs:$p),
- !strconcat("selp.", TypeStr, "\t$dst, $a, $b, $p;"),
+ !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"),
[(set RC:$dst, (select Int1Regs:$p, ImmNode:$a, ImmNode:$b))]>;
}
}
@@ -1960,33 +1960,33 @@ let mayLoad = 1 in {
class LoadParamV2MemInst<NVPTXRegClass regclass, string opstr> :
NVPTXInst<(outs regclass:$dst, regclass:$dst2), (ins i32imm:$b),
!strconcat("ld.param.v2", opstr,
- "\t{{$dst, $dst2}}, [retval0+$b];"), []>;
+ " \t{{$dst, $dst2}}, [retval0+$b];"), []>;
class LoadParamV4MemInst<NVPTXRegClass regclass, string opstr> :
NVPTXInst<(outs regclass:$dst, regclass:$dst2, regclass:$dst3,
regclass:$dst4),
(ins i32imm:$b),
!strconcat("ld.param.v4", opstr,
- "\t{{$dst, $dst2, $dst3, $dst4}}, [retval0+$b];"),
+ " \t{{$dst, $dst2, $dst3, $dst4}}, [retval0+$b];"),
[]>;
}
class LoadParamRegInst<NVPTXRegClass regclass, string opstr> :
NVPTXInst<(outs regclass:$dst), (ins i32imm:$b),
- !strconcat("mov", opstr, "\t$dst, retval$b;"),
+ !strconcat("mov", opstr, " \t$dst, retval$b;"),
[(set regclass:$dst, (LoadParam (i32 0), (i32 imm:$b)))]>;
let mayStore = 1 in {
class StoreParamInst<NVPTXRegClass regclass, string opstr> :
NVPTXInst<(outs), (ins regclass:$val, i32imm:$a, i32imm:$b),
- !strconcat("st.param", opstr, "\t[param$a+$b], $val;"),
+ !strconcat("st.param", opstr, " \t[param$a+$b], $val;"),
[]>;
class StoreParamV2Inst<NVPTXRegClass regclass, string opstr> :
NVPTXInst<(outs), (ins regclass:$val, regclass:$val2,
i32imm:$a, i32imm:$b),
!strconcat("st.param.v2", opstr,
- "\t[param$a+$b], {{$val, $val2}};"),
+ " \t[param$a+$b], {{$val, $val2}};"),
[]>;
class StoreParamV4Inst<NVPTXRegClass regclass, string opstr> :
@@ -1994,18 +1994,18 @@ let mayStore = 1 in {
regclass:$val4, i32imm:$a,
i32imm:$b),
!strconcat("st.param.v4", opstr,
- "\t[param$a+$b], {{$val, $val2, $val3, $val4}};"),
+ " \t[param$a+$b], {{$val, $val2, $val3, $val4}};"),
[]>;
class StoreRetvalInst<NVPTXRegClass regclass, string opstr> :
NVPTXInst<(outs), (ins regclass:$val, i32imm:$a),
- !strconcat("st.param", opstr, "\t[func_retval0+$a], $val;"),
+ !strconcat("st.param", opstr, " \t[func_retval0+$a], $val;"),
[]>;
class StoreRetvalV2Inst<NVPTXRegClass regclass, string opstr> :
NVPTXInst<(outs), (ins regclass:$val, regclass:$val2, i32imm:$a),
!strconcat("st.param.v2", opstr,
- "\t[func_retval0+$a], {{$val, $val2}};"),
+ " \t[func_retval0+$a], {{$val, $val2}};"),
[]>;
class StoreRetvalV4Inst<NVPTXRegClass regclass, string opstr> :
@@ -2013,7 +2013,7 @@ let mayStore = 1 in {
(ins regclass:$val, regclass:$val2, regclass:$val3,
regclass:$val4, i32imm:$a),
!strconcat("st.param.v4", opstr,
- "\t[func_retval0+$a], {{$val, $val2, $val3, $val4}};"),
+ " \t[func_retval0+$a], {{$val, $val2, $val3, $val4}};"),
[]>;
}
@@ -2188,14 +2188,14 @@ def DeclareScalarRegInst :
class MoveParamInst<NVPTXRegClass regclass, string asmstr> :
NVPTXInst<(outs regclass:$dst), (ins regclass:$src),
- !strconcat("mov", asmstr, "\t$dst, $src;"),
+ !strconcat("mov", asmstr, " \t$dst, $src;"),
[(set regclass:$dst, (MoveParam regclass:$src))]>;
def MoveParamI64 : MoveParamInst<Int64Regs, ".b64">;
def MoveParamI32 : MoveParamInst<Int32Regs, ".b32">;
def MoveParamI16 :
NVPTXInst<(outs Int16Regs:$dst), (ins Int16Regs:$src),
- "cvt.u16.u32\t$dst, $src;",
+ "cvt.u16.u32 \t$dst, $src;",
[(set Int16Regs:$dst, (MoveParam Int16Regs:$src))]>;
def MoveParamF64 : MoveParamInst<Float64Regs, ".f64">;
def MoveParamF32 : MoveParamInst<Float32Regs, ".f32">;
@@ -2763,39 +2763,39 @@ let hasSideEffects = 0 in {
def V4I16toI64 : NVPTXInst<(outs Int64Regs:$d),
(ins Int16Regs:$s1, Int16Regs:$s2,
Int16Regs:$s3, Int16Regs:$s4),
- "mov.b64\t$d, {{$s1, $s2, $s3, $s4}};", []>;
+ "mov.b64 \t$d, {{$s1, $s2, $s3, $s4}};", []>;
def V2I16toI32 : NVPTXInst<(outs Int32Regs:$d),
(ins Int16Regs:$s1, Int16Regs:$s2),
- "mov.b32\t$d, {{$s1, $s2}};", []>;
+ "mov.b32 \t$d, {{$s1, $s2}};", []>;
def V2I32toI64 : NVPTXInst<(outs Int64Regs:$d),
(ins Int32Regs:$s1, Int32Regs:$s2),
- "mov.b64\t$d, {{$s1, $s2}};", []>;
+ "mov.b64 \t$d, {{$s1, $s2}};", []>;
def V2F32toF64 : NVPTXInst<(outs Float64Regs:$d),
(ins Float32Regs:$s1, Float32Regs:$s2),
- "mov.b64\t$d, {{$s1, $s2}};", []>;
+ "mov.b64 \t$d, {{$s1, $s2}};", []>;
// unpack a larger int register to a set of smaller int registers
def I64toV4I16 : NVPTXInst<(outs Int16Regs:$d1, Int16Regs:$d2,
Int16Regs:$d3, Int16Regs:$d4),
(ins Int64Regs:$s),
- "mov.b64\t{{$d1, $d2, $d3, $d4}}, $s;", []>;
+ "mov.b64 \t{{$d1, $d2, $d3, $d4}}, $s;", []>;
def I32toV2I16 : NVPTXInst<(outs Int16Regs:$d1, Int16Regs:$d2),
(ins Int32Regs:$s),
- "mov.b32\t{{$d1, $d2}}, $s;", []>;
+ "mov.b32 \t{{$d1, $d2}}, $s;", []>;
def I64toV2I32 : NVPTXInst<(outs Int32Regs:$d1, Int32Regs:$d2),
(ins Int64Regs:$s),
- "mov.b64\t{{$d1, $d2}}, $s;", []>;
+ "mov.b64 \t{{$d1, $d2}}, $s;", []>;
def F64toV2F32 : NVPTXInst<(outs Float32Regs:$d1, Float32Regs:$d2),
(ins Float64Regs:$s),
- "mov.b64\t{{$d1, $d2}}, $s;", []>;
+ "mov.b64 \t{{$d1, $d2}}, $s;", []>;
}
// Count leading zeros
let hasSideEffects = 0 in {
def CLZr32 : NVPTXInst<(outs Int32Regs:$d), (ins Int32Regs:$a),
- "clz.b32\t$d, $a;", []>;
+ "clz.b32 \t$d, $a;", []>;
def CLZr64 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
- "clz.b64\t$d, $a;", []>;
+ "clz.b64 \t$d, $a;", []>;
}
// 32-bit has a direct PTX instruction
@@ -2831,9 +2831,9 @@ def : Pat<(i32 (zext (ctlz Int16Regs:$a)
// Population count
let hasSideEffects = 0 in {
def POPCr32 : NVPTXInst<(outs Int32Regs:$d), (ins Int32Regs:$a),
- "popc.b32\t$d, $a;", []>;
+ "popc.b32 \t$d, $a;", []>;
def POPCr64 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
- "popc.b64\t$d, $a;", []>;
+ "popc.b64 \t$d, $a;", []>;
}
// 32-bit has a direct PTX instruction
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td?rev=292306&r1=292305&r2=292306&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td Tue Jan 17 18:09:36 2017
@@ -62,7 +62,7 @@ def INT_BARRIER0_OR : NVPTXInst<(outs In
"}}"),
[(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
@@ -1919,7 +1919,7 @@ def ISSPACEP_SHARED_64
// 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)>;
@@ -2098,19 +2098,19 @@ def TEX_1D_F32_S32
: 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
@@ -2118,27 +2118,27 @@ 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
@@ -2146,27 +2146,27 @@ 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
@@ -2174,7 +2174,7 @@ 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\\};",
[]>;
@@ -2182,14 +2182,14 @@ def TEX_1D_ARRAY_F32_S32
: 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
@@ -2197,7 +2197,7 @@ 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
@@ -2205,21 +2205,21 @@ 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
@@ -2227,7 +2227,7 @@ 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
@@ -2235,21 +2235,21 @@ 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
@@ -2257,7 +2257,7 @@ 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
@@ -2265,7 +2265,7 @@ 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\\};",
[]>;
@@ -2273,14 +2273,14 @@ def TEX_2D_F32_S32
: 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
@@ -2288,7 +2288,7 @@ 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
@@ -2297,7 +2297,7 @@ 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\\};",
[]>;
@@ -2305,14 +2305,14 @@ def TEX_2D_S32_S32
: 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
@@ -2320,7 +2320,7 @@ 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
@@ -2329,7 +2329,7 @@ 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\\};",
[]>;
@@ -2337,14 +2337,14 @@ def TEX_2D_U32_S32
: 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
@@ -2352,7 +2352,7 @@ 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
@@ -2361,7 +2361,7 @@ 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\\};",
[]>;
@@ -2371,7 +2371,7 @@ def TEX_2D_ARRAY_F32_S32
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
@@ -2379,7 +2379,7 @@ 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
@@ -2387,7 +2387,7 @@ 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
@@ -2396,7 +2396,7 @@ 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\\};",
[]>;
@@ -2405,7 +2405,7 @@ def TEX_2D_ARRAY_S32_S32
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
@@ -2413,7 +2413,7 @@ 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
@@ -2421,7 +2421,7 @@ 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
@@ -2431,7 +2431,7 @@ 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\\};",
[]>;
@@ -2440,7 +2440,7 @@ def TEX_2D_ARRAY_U32_S32
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
@@ -2448,7 +2448,7 @@ 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
@@ -2456,7 +2456,7 @@ 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
@@ -2466,7 +2466,7 @@ 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\\};",
[]>;
@@ -2476,7 +2476,7 @@ def TEX_3D_F32_S32
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
@@ -2484,7 +2484,7 @@ 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
@@ -2492,7 +2492,7 @@ 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
@@ -2503,7 +2503,7 @@ 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\\};",
@@ -2513,7 +2513,7 @@ def TEX_3D_S32_S32
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
@@ -2521,7 +2521,7 @@ 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
@@ -2529,7 +2529,7 @@ 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
@@ -2540,7 +2540,7 @@ 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\\};",
@@ -2550,7 +2550,7 @@ def TEX_3D_U32_S32
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
@@ -2558,7 +2558,7 @@ 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
@@ -2566,7 +2566,7 @@ 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
@@ -2577,7 +2577,7 @@ 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\\};",
@@ -2588,7 +2588,7 @@ def TEX_CUBE_F32_F32
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
@@ -2597,7 +2597,7 @@ 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
@@ -2605,7 +2605,7 @@ 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
@@ -2614,7 +2614,7 @@ 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
@@ -2622,7 +2622,7 @@ 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
@@ -2631,7 +2631,7 @@ 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;",
[]>;
@@ -2640,7 +2640,7 @@ def TEX_CUBE_ARRAY_F32_F32
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
@@ -2649,7 +2649,7 @@ 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
@@ -2657,7 +2657,7 @@ 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
@@ -2666,7 +2666,7 @@ 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
@@ -2674,7 +2674,7 @@ 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
@@ -2683,7 +2683,7 @@ 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;",
[]>;
@@ -2691,84 +2691,84 @@ def TLD4_R_2D_F32_F32
: 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\\}];",
[]>;
}
@@ -2781,19 +2781,19 @@ def TEX_UNIFIED_1D_F32_S32
: 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
@@ -2801,27 +2801,27 @@ 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
@@ -2829,27 +2829,27 @@ 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
@@ -2857,7 +2857,7 @@ 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\\};",
[]>;
@@ -2865,14 +2865,14 @@ def TEX_UNIFIED_1D_ARRAY_F32_S32
: 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
@@ -2880,7 +2880,7 @@ 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
@@ -2888,21 +2888,21 @@ 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
@@ -2910,7 +2910,7 @@ 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
@@ -2918,21 +2918,21 @@ 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
@@ -2940,7 +2940,7 @@ 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
@@ -2948,7 +2948,7 @@ 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\\};",
[]>;
@@ -2956,14 +2956,14 @@ def TEX_UNIFIED_2D_F32_S32
: 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
@@ -2971,7 +2971,7 @@ 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
@@ -2980,7 +2980,7 @@ 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\\};",
[]>;
@@ -2988,14 +2988,14 @@ def TEX_UNIFIED_2D_S32_S32
: 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
@@ -3003,7 +3003,7 @@ 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
@@ -3012,7 +3012,7 @@ 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\\};",
[]>;
@@ -3020,14 +3020,14 @@ def TEX_UNIFIED_2D_U32_S32
: 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
@@ -3035,7 +3035,7 @@ 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
@@ -3044,7 +3044,7 @@ 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\\};",
[]>;
@@ -3054,7 +3054,7 @@ def TEX_UNIFIED_2D_ARRAY_F32_S32
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
@@ -3062,7 +3062,7 @@ 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
@@ -3070,7 +3070,7 @@ 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
@@ -3079,7 +3079,7 @@ 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\\};",
[]>;
@@ -3088,7 +3088,7 @@ def TEX_UNIFIED_2D_ARRAY_S32_S32
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
@@ -3096,7 +3096,7 @@ 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
@@ -3104,7 +3104,7 @@ 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
@@ -3114,7 +3114,7 @@ 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\\};",
[]>;
@@ -3123,7 +3123,7 @@ def TEX_UNIFIED_2D_ARRAY_U32_S32
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
@@ -3131,7 +3131,7 @@ 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
@@ -3139,7 +3139,7 @@ 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
@@ -3149,7 +3149,7 @@ 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\\};",
[]>;
@@ -3159,7 +3159,7 @@ def TEX_UNIFIED_3D_F32_S32
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
@@ -3167,7 +3167,7 @@ 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
@@ -3175,7 +3175,7 @@ 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
@@ -3186,7 +3186,7 @@ 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\\};",
@@ -3196,7 +3196,7 @@ def TEX_UNIFIED_3D_S32_S32
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
@@ -3204,7 +3204,7 @@ 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
@@ -3212,7 +3212,7 @@ 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
@@ -3223,7 +3223,7 @@ 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\\};",
@@ -3233,7 +3233,7 @@ def TEX_UNIFIED_3D_U32_S32
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
@@ -3241,7 +3241,7 @@ 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
@@ -3249,7 +3249,7 @@ 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
@@ -3260,7 +3260,7 @@ 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\\};",
@@ -3271,7 +3271,7 @@ def TEX_UNIFIED_CUBE_F32_F32
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
@@ -3280,7 +3280,7 @@ 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
@@ -3288,7 +3288,7 @@ 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
@@ -3297,7 +3297,7 @@ 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
@@ -3305,7 +3305,7 @@ 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
@@ -3314,7 +3314,7 @@ 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;",
[]>;
@@ -3323,7 +3323,7 @@ def TEX_UNIFIED_CUBE_ARRAY_F32_F32
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
@@ -3332,7 +3332,7 @@ 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
@@ -3340,7 +3340,7 @@ 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
@@ -3349,7 +3349,7 @@ 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
@@ -3357,7 +3357,7 @@ 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
@@ -3366,7 +3366,7 @@ 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;",
[]>;
@@ -3374,84 +3374,84 @@ def TLD4_UNIFIED_R_2D_F32_F32
: 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\\}];",
[]>;
}
More information about the llvm-commits
mailing list