mirror of
https://github.com/vosen/ZLUDA.git
synced 2025-08-04 15:19:49 +00:00
Simplify mad_hi_cc and remove leftover comments
This commit is contained in:
NyanCatTW1
parent
c27cfbcfa3
commit
e2b9501a99
4 changed files with 74 additions and 142 deletions
|
|
@ -2099,21 +2099,6 @@ fn emit_inst_madcc(
|
||||||
mul_result,
|
mul_result,
|
||||||
arg.src3,
|
arg.src3,
|
||||||
)
|
)
|
||||||
/*
|
|
||||||
let builder = ctx.builder.get();
|
|
||||||
let src1 = ctx.names.value(arg.src1)?;
|
|
||||||
let src2 = ctx.names.value(arg.src2)?;
|
|
||||||
let mul_result = unsafe { LLVMBuildMul(builder, src1, src2, LLVM_UNNAMED) };
|
|
||||||
emit_inst_addsub_cc_impl(
|
|
||||||
ctx,
|
|
||||||
"add",
|
|
||||||
type_,
|
|
||||||
arg.dst,
|
|
||||||
arg.carry_out,
|
|
||||||
mul_result,
|
|
||||||
arg.src3,
|
|
||||||
)
|
|
||||||
*/
|
|
||||||
}
|
}
|
||||||
|
|
||||||
fn get_llvm_type_struct<'a>(
|
fn get_llvm_type_struct<'a>(
|
||||||
|
|
@ -2192,29 +2177,6 @@ fn emit_inst_madc(
|
||||||
mul_result,
|
mul_result,
|
||||||
args.src3,
|
args.src3,
|
||||||
)
|
)
|
||||||
/*
|
|
||||||
let src3 = ctx.names.value(args.src3)?;
|
|
||||||
let add_no_carry = unsafe { LLVMBuildAdd(builder, mul_result, src3, LLVM_UNNAMED) };
|
|
||||||
let carry_flag = ctx.names.value(args.carry_in)?;
|
|
||||||
let llvm_type = get_llvm_type(ctx, &ast::Type::Scalar(type_))?;
|
|
||||||
let carry_flag = unsafe { LLVMBuildZExt(builder, carry_flag, llvm_type, LLVM_UNNAMED) };
|
|
||||||
if let Some(carry_out) = args.carry_out {
|
|
||||||
emit_inst_addsub_cc_impl(
|
|
||||||
ctx,
|
|
||||||
"add",
|
|
||||||
type_,
|
|
||||||
args.dst,
|
|
||||||
carry_out,
|
|
||||||
add_no_carry,
|
|
||||||
carry_flag,
|
|
||||||
)?;
|
|
||||||
} else {
|
|
||||||
ctx.names.register_result(args.dst, |dst| unsafe {
|
|
||||||
LLVMBuildAdd(builder, add_no_carry, carry_flag, dst)
|
|
||||||
});
|
|
||||||
}
|
|
||||||
Ok(())
|
|
||||||
*/
|
|
||||||
}
|
}
|
||||||
|
|
||||||
fn emit_inst_add_c(
|
fn emit_inst_add_c(
|
||||||
|
|
|
||||||
|
|
@ -1,12 +1,12 @@
|
||||||
target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7"
|
target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7"
|
||||||
target triple = "amdgcn-amd-amdhsa"
|
target triple = "amdgcn-amd-amdhsa"
|
||||||
|
|
||||||
define protected amdgpu_kernel void @mad_hi_cc(ptr addrspace(4) byref(i64) %"76", ptr addrspace(4) byref(i64) %"77") #0 {
|
define protected amdgpu_kernel void @mad_hi_cc(ptr addrspace(4) byref(i64) %"61", ptr addrspace(4) byref(i64) %"62") #0 {
|
||||||
"99":
|
"78":
|
||||||
|
%"14" = alloca i1, align 1, addrspace(5)
|
||||||
|
store i1 false, ptr addrspace(5) %"14", align 1
|
||||||
%"15" = alloca i1, align 1, addrspace(5)
|
%"15" = alloca i1, align 1, addrspace(5)
|
||||||
store i1 false, ptr addrspace(5) %"15", align 1
|
store i1 false, ptr addrspace(5) %"15", align 1
|
||||||
%"16" = alloca i1, align 1, addrspace(5)
|
|
||||||
store i1 false, ptr addrspace(5) %"16", align 1
|
|
||||||
%"4" = alloca i64, align 8, addrspace(5)
|
%"4" = alloca i64, align 8, addrspace(5)
|
||||||
%"5" = alloca i64, align 8, addrspace(5)
|
%"5" = alloca i64, align 8, addrspace(5)
|
||||||
%"6" = alloca i32, align 4, addrspace(5)
|
%"6" = alloca i32, align 4, addrspace(5)
|
||||||
|
|
@ -17,90 +17,69 @@ define protected amdgpu_kernel void @mad_hi_cc(ptr addrspace(4) byref(i64) %"76"
|
||||||
%"11" = alloca i32, align 4, addrspace(5)
|
%"11" = alloca i32, align 4, addrspace(5)
|
||||||
%"12" = alloca i32, align 4, addrspace(5)
|
%"12" = alloca i32, align 4, addrspace(5)
|
||||||
%"13" = alloca i32, align 4, addrspace(5)
|
%"13" = alloca i32, align 4, addrspace(5)
|
||||||
%"14" = alloca i32, align 4, addrspace(5)
|
%"16" = load i64, ptr addrspace(4) %"61", align 8
|
||||||
%"17" = load i64, ptr addrspace(4) %"76", align 8
|
store i64 %"16", ptr addrspace(5) %"4", align 8
|
||||||
store i64 %"17", ptr addrspace(5) %"4", align 8
|
%"17" = load i64, ptr addrspace(4) %"62", align 8
|
||||||
%"18" = load i64, ptr addrspace(4) %"77", align 8
|
store i64 %"17", ptr addrspace(5) %"5", align 8
|
||||||
store i64 %"18", ptr addrspace(5) %"5", align 8
|
%"19" = load i64, ptr addrspace(5) %"4", align 8
|
||||||
%"20" = load i64, ptr addrspace(5) %"4", align 8
|
%"64" = inttoptr i64 %"19" to ptr
|
||||||
%"79" = inttoptr i64 %"20" to ptr
|
%"63" = load i32, ptr %"64", align 4
|
||||||
%"78" = load i32, ptr %"79", align 4
|
store i32 %"63", ptr addrspace(5) %"8", align 4
|
||||||
store i32 %"78", ptr addrspace(5) %"8", align 4
|
%"21" = load i64, ptr addrspace(5) %"4", align 8
|
||||||
%"22" = load i64, ptr addrspace(5) %"4", align 8
|
%"65" = inttoptr i64 %"21" to ptr
|
||||||
%"80" = inttoptr i64 %"22" to ptr
|
%"80" = getelementptr inbounds i8, ptr %"65", i64 4
|
||||||
%"101" = getelementptr inbounds i8, ptr %"80", i64 4
|
%"66" = load i32, ptr %"80", align 4
|
||||||
%"81" = load i32, ptr %"101", align 4
|
store i32 %"66", ptr addrspace(5) %"9", align 4
|
||||||
store i32 %"81", ptr addrspace(5) %"9", align 4
|
%"23" = load i64, ptr addrspace(5) %"4", align 8
|
||||||
%"24" = load i64, ptr addrspace(5) %"4", align 8
|
%"67" = inttoptr i64 %"23" to ptr
|
||||||
%"82" = inttoptr i64 %"24" to ptr
|
%"82" = getelementptr inbounds i8, ptr %"67", i64 8
|
||||||
%"103" = getelementptr inbounds i8, ptr %"82", i64 8
|
%"22" = load i32, ptr %"82", align 4
|
||||||
%"23" = load i32, ptr %"103", align 4
|
store i32 %"22", ptr addrspace(5) %"10", align 4
|
||||||
store i32 %"23", ptr addrspace(5) %"10", align 4
|
%"26" = load i32, ptr addrspace(5) %"8", align 4
|
||||||
%"27" = load i32, ptr addrspace(5) %"8", align 4
|
%"27" = load i32, ptr addrspace(5) %"9", align 4
|
||||||
%"28" = load i32, ptr addrspace(5) %"9", align 4
|
%"28" = load i32, ptr addrspace(5) %"10", align 4
|
||||||
%"29" = load i32, ptr addrspace(5) %"10", align 4
|
%0 = sext i32 %"26" to i64
|
||||||
%0 = sext i32 %"27" to i64
|
%1 = sext i32 %"27" to i64
|
||||||
%1 = sext i32 %"28" to i64
|
|
||||||
%2 = mul nsw i64 %0, %1
|
%2 = mul nsw i64 %0, %1
|
||||||
%3 = lshr i64 %2, 32
|
%3 = lshr i64 %2, 32
|
||||||
%4 = trunc i64 %3 to i32
|
%4 = trunc i64 %3 to i32
|
||||||
%5 = call { i32, i1 } @llvm.uadd.with.overflow.i32(i32 %4, i32 %"29")
|
%5 = call { i32, i1 } @llvm.uadd.with.overflow.i32(i32 %4, i32 %"28")
|
||||||
%"25" = extractvalue { i32, i1 } %5, 0
|
%"24" = extractvalue { i32, i1 } %5, 0
|
||||||
%"26" = extractvalue { i32, i1 } %5, 1
|
%"25" = extractvalue { i32, i1 } %5, 1
|
||||||
store i32 %"25", ptr addrspace(5) %"7", align 4
|
store i32 %"24", ptr addrspace(5) %"7", align 4
|
||||||
store i1 %"26", ptr addrspace(5) %"15", align 1
|
store i1 %"25", ptr addrspace(5) %"14", align 1
|
||||||
%"30" = load i64, ptr addrspace(5) %"5", align 8
|
%6 = call { i32, i1 } @llvm.uadd.with.overflow.i32(i32 1, i32 -2)
|
||||||
%"31" = load i32, ptr addrspace(5) %"7", align 4
|
%"29" = extractvalue { i32, i1 } %6, 0
|
||||||
%"86" = inttoptr i64 %"30" to ptr
|
%"30" = extractvalue { i32, i1 } %6, 1
|
||||||
store i32 %"31", ptr %"86", align 4
|
store i32 %"29", ptr addrspace(5) %"6", align 4
|
||||||
%6 = call { i32, i1 } @llvm.uadd.with.overflow.i32(i32 0, i32 -1)
|
store i1 %"30", ptr addrspace(5) %"14", align 1
|
||||||
%"32" = extractvalue { i32, i1 } %6, 0
|
%"32" = load i1, ptr addrspace(5) %"14", align 1
|
||||||
%"33" = extractvalue { i32, i1 } %6, 1
|
%7 = zext i1 %"32" to i32
|
||||||
store i32 %"32", ptr addrspace(5) %"6", align 4
|
%"71" = add i32 0, %7
|
||||||
store i1 %"33", ptr addrspace(5) %"15", align 1
|
store i32 %"71", ptr addrspace(5) %"12", align 4
|
||||||
%"36" = load i1, ptr addrspace(5) %"15", align 1
|
|
||||||
%7 = zext i1 %"36" to i32
|
|
||||||
%8 = call { i32, i1 } @llvm.uadd.with.overflow.i32(i32 1, i32 -1)
|
%8 = call { i32, i1 } @llvm.uadd.with.overflow.i32(i32 1, i32 -1)
|
||||||
%9 = extractvalue { i32, i1 } %8, 0
|
%"33" = extractvalue { i32, i1 } %8, 0
|
||||||
%10 = extractvalue { i32, i1 } %8, 1
|
%"34" = extractvalue { i32, i1 } %8, 1
|
||||||
%11 = call { i32, i1 } @llvm.uadd.with.overflow.i32(i32 %9, i32 %7)
|
store i32 %"33", ptr addrspace(5) %"6", align 4
|
||||||
%"87" = extractvalue { i32, i1 } %11, 0
|
store i1 %"34", ptr addrspace(5) %"14", align 1
|
||||||
%12 = extractvalue { i32, i1 } %11, 1
|
%"36" = load i1, ptr addrspace(5) %"14", align 1
|
||||||
%"35" = xor i1 %10, %12
|
%9 = zext i1 %"36" to i32
|
||||||
store i32 %"87", ptr addrspace(5) %"11", align 4
|
%"72" = add i32 0, %9
|
||||||
store i1 %"35", ptr addrspace(5) %"15", align 1
|
store i32 %"72", ptr addrspace(5) %"13", align 4
|
||||||
%"38" = load i1, ptr addrspace(5) %"15", align 1
|
%"37" = load i64, ptr addrspace(5) %"5", align 8
|
||||||
%13 = zext i1 %"38" to i32
|
%"38" = load i32, ptr addrspace(5) %"7", align 4
|
||||||
%"88" = add i32 0, %13
|
%"73" = inttoptr i64 %"37" to ptr
|
||||||
store i32 %"88", ptr addrspace(5) %"12", align 4
|
store i32 %"38", ptr %"73", align 4
|
||||||
%"40" = load i1, ptr addrspace(5) %"15", align 1
|
%"39" = load i64, ptr addrspace(5) %"5", align 8
|
||||||
%14 = zext i1 %"40" to i32
|
%"40" = load i32, ptr addrspace(5) %"12", align 4
|
||||||
%"89" = add i32 0, %14
|
%"74" = inttoptr i64 %"39" to ptr
|
||||||
store i32 %"89", ptr addrspace(5) %"13", align 4
|
%"84" = getelementptr inbounds i8, ptr %"74", i64 4
|
||||||
%"42" = load i1, ptr addrspace(5) %"16", align 1
|
store i32 %"40", ptr %"84", align 4
|
||||||
%15 = zext i1 %"42" to i32
|
%"41" = load i64, ptr addrspace(5) %"5", align 8
|
||||||
%"90" = sub i32 2, %15
|
%"42" = load i32, ptr addrspace(5) %"13", align 4
|
||||||
store i32 %"90", ptr addrspace(5) %"14", align 4
|
%"76" = inttoptr i64 %"41" to ptr
|
||||||
%"43" = load i64, ptr addrspace(5) %"5", align 8
|
%"86" = getelementptr inbounds i8, ptr %"76", i64 8
|
||||||
%"44" = load i32, ptr addrspace(5) %"11", align 4
|
store i32 %"42", ptr %"86", align 4
|
||||||
%"91" = inttoptr i64 %"43" to ptr
|
|
||||||
%"105" = getelementptr inbounds i8, ptr %"91", i64 4
|
|
||||||
store i32 %"44", ptr %"105", align 4
|
|
||||||
%"45" = load i64, ptr addrspace(5) %"5", align 8
|
|
||||||
%"46" = load i32, ptr addrspace(5) %"12", align 4
|
|
||||||
%"93" = inttoptr i64 %"45" to ptr
|
|
||||||
%"107" = getelementptr inbounds i8, ptr %"93", i64 8
|
|
||||||
store i32 %"46", ptr %"107", align 4
|
|
||||||
%"47" = load i64, ptr addrspace(5) %"5", align 8
|
|
||||||
%"48" = load i32, ptr addrspace(5) %"13", align 4
|
|
||||||
%"95" = inttoptr i64 %"47" to ptr
|
|
||||||
%"109" = getelementptr inbounds i8, ptr %"95", i64 12
|
|
||||||
store i32 %"48", ptr %"109", align 4
|
|
||||||
%"49" = load i64, ptr addrspace(5) %"5", align 8
|
|
||||||
%"50" = load i32, ptr addrspace(5) %"14", align 4
|
|
||||||
%"97" = inttoptr i64 %"49" to ptr
|
|
||||||
%"111" = getelementptr inbounds i8, ptr %"97", i64 16
|
|
||||||
store i32 %"50", ptr %"111", align 4
|
|
||||||
ret void
|
ret void
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -17,9 +17,8 @@
|
||||||
.reg .b32 src3;
|
.reg .b32 src3;
|
||||||
|
|
||||||
.reg .b32 result_1;
|
.reg .b32 result_1;
|
||||||
.reg .b32 carry_out_1_1;
|
.reg .b32 carry_out_1;
|
||||||
.reg .b32 carry_out_1_2;
|
.reg .b32 carry_out_2;
|
||||||
.reg .b32 carry_out_1_3;
|
|
||||||
|
|
||||||
ld.param.u64 in_addr, [input];
|
ld.param.u64 in_addr, [input];
|
||||||
ld.param.u64 out_addr, [output];
|
ld.param.u64 out_addr, [output];
|
||||||
|
|
@ -29,22 +28,14 @@
|
||||||
ld.s32 src2, [in_addr+4];
|
ld.s32 src2, [in_addr+4];
|
||||||
ld.b32 src3, [in_addr+8];
|
ld.b32 src3, [in_addr+8];
|
||||||
mad.hi.cc.s32 dst1, src1, src2, src3;
|
mad.hi.cc.s32 dst1, src1, src2, src3;
|
||||||
|
|
||||||
|
mad.hi.cc.u32 unused, 65536, 65536, 4294967294; // non-overflowing
|
||||||
|
addc.u32 carry_out_1, 0, 0; // carry_out_1 should be 0
|
||||||
|
mad.hi.cc.u32 unused, 65536, 65536, 4294967295; // overflowing
|
||||||
|
addc.u32 carry_out_2, 0, 0; // carry_out_2 should be 1
|
||||||
|
|
||||||
st.s32 [out_addr], dst1;
|
st.s32 [out_addr], dst1;
|
||||||
|
st.s32 [out_addr+4], carry_out_1;
|
||||||
// set carry=1
|
st.s32 [out_addr+8], carry_out_2;
|
||||||
mad.lo.cc.u32 unused, 0, 0, 4294967295;
|
|
||||||
// overflow addition
|
|
||||||
madc.hi.cc.u32 result_1, 65536, 65536, 4294967295;
|
|
||||||
// write carry
|
|
||||||
madc.lo.u32 carry_out_1_1, 0, 0, 0;
|
|
||||||
// overflow is also detected by addc
|
|
||||||
addc.u32 carry_out_1_2, 0, 0;
|
|
||||||
// but not subc
|
|
||||||
subc.u32 carry_out_1_3, 2, 0;
|
|
||||||
|
|
||||||
st.s32 [out_addr+4], result_1;
|
|
||||||
st.s32 [out_addr+8], carry_out_1_1;
|
|
||||||
st.s32 [out_addr+12], carry_out_1_2;
|
|
||||||
st.s32 [out_addr+16], carry_out_1_3;
|
|
||||||
ret;
|
ret;
|
||||||
}
|
}
|
||||||
|
|
|
||||||
|
|
@ -290,7 +290,7 @@ test_ptx!(
|
||||||
[2147487519u32, 4294934539]
|
[2147487519u32, 4294934539]
|
||||||
);
|
);
|
||||||
test_ptx!(madc_cc2, [0xDEADu32], [0u32, 1, 1, 2]);
|
test_ptx!(madc_cc2, [0xDEADu32], [0u32, 1, 1, 2]);
|
||||||
test_ptx!(mad_hi_cc, [0x26223377u32, 0x70777766u32, 0x60666633u32], [0x71272866u32, 0u32, 1, 1, 2]); // Multi-tap :)
|
test_ptx!(mad_hi_cc, [0x26223377u32, 0x70777766u32, 0x60666633u32], [0x71272866u32, 0u32, 1u32]); // Multi-tap :)
|
||||||
test_ptx!(mov_vector_cast, [0x200000001u64], [2u32, 1u32]);
|
test_ptx!(mov_vector_cast, [0x200000001u64], [2u32, 1u32]);
|
||||||
test_ptx!(
|
test_ptx!(
|
||||||
cvt_clamp,
|
cvt_clamp,
|
||||||
|
|
|
||||||
Loading…
Add table
Add a link
Reference in a new issue