Improve compatibility with certain old code

8 files changed, 262 insertions, 36 deletions

diff --git a/ptx/src/ast.rs b/ptx/src/ast.rs
index 0281961..5e85d61 100644
--- a/ptx/src/ast.rs
+++ b/ptx/src/ast.rs
@@ -475,6 +475,7 @@ pub enum Instruction<P: ArgParams> {
     MatchAny(Arg3<P>),

     Red(AtomDetails, Arg2St<P>),

     Nanosleep(Arg1<P>),

+    Mul24(Mul24Details, Arg3<P>),

 }

 

 #[derive(Copy, Clone)]

@@ -495,6 +496,7 @@ pub enum ShflMode {
 #[derive(Copy, Clone)]

 pub struct VoteDetails {

     pub mode: VoteMode,

+    pub sync: bool,

     pub negate_pred: bool,

 }

 

@@ -833,6 +835,18 @@ pub enum MulIntControl {
     Wide,

 }

 

+#[derive(Copy, Clone)]

+pub struct Mul24Details {

+    pub type_: ScalarType,

+    pub control: Mul24Control,

+}

+

+#[derive(Copy, Clone, PartialEq, Eq)]

+pub enum Mul24Control {

+    Low,

+    High,

+}

+

 #[derive(PartialEq, Eq, Copy, Clone)]

 pub enum RoundingMode {

     NearestEven,

@@ -1364,11 +1378,19 @@ pub struct SurfaceDetails {
     pub geometry: TextureGeometry,

     pub vector: Option<u8>,

     pub type_: ScalarType,

+    pub clamp: SurfaceClamp,

     // direct = takes .texref, indirect = takes .u64

     pub direct: bool,

 }

 

 #[derive(Clone, Copy, PartialEq, Eq)]

+pub enum SurfaceClamp {

+    Trap,

+    Clamp,

+    Zero

+}

+

+#[derive(Clone, Copy, PartialEq, Eq)]

 pub enum TextureGeometry {

     OneD,

     TwoD,

diff --git a/ptx/src/emit.rs b/ptx/src/emit.rs
index 94cc973..44af810 100644
--- a/ptx/src/emit.rs
+++ b/ptx/src/emit.rs
@@ -1137,6 +1137,7 @@ fn emit_instruction(
         ast::Instruction::Vshr(arg) => emit_inst_vshr(ctx, arg)?,
         ast::Instruction::Set(details, arg) => emit_inst_set(ctx, details, arg)?,
         ast::Instruction::Red(details, arg) => emit_inst_red(ctx, details, arg)?,
+        ast::Instruction::Mul24(details, arg) => emit_inst_mul24(ctx, details, arg)?,
         // replaced by function calls or Statement variants
         ast::Instruction::Activemask { .. }
         | ast::Instruction::Bar(..)
@@ -1161,6 +1162,46 @@ fn emit_instruction(
     })
 }
 
+fn emit_inst_mul24(
+    ctx: &mut EmitContext,
+    details: &ast::Mul24Details,
+    arg: &ast::Arg3<ExpandedArgParams>,
+) -> Result<(), TranslateError> {
+    unsafe fn clear_top8_bits(
+        ctx: &mut EmitContext<'_>,
+        src: Id,
+        right_shift: unsafe extern "C" fn(
+            LLVMBuilderRef,
+            LLVMValueRef,
+            LLVMValueRef,
+            *const i8,
+        ) -> LLVMValueRef,
+    ) -> Result<LLVMValueRef, TranslateError> {
+        let u32_llvm_type = get_llvm_type(ctx, &ast::Type::Scalar(ast::ScalarType::U32))?;
+        let const_8 = LLVMConstInt(u32_llvm_type, 8, 0);
+        let src = ctx.names.value(src)?;
+        let src = LLVMBuildShl(ctx.builder.get(), src, const_8, LLVM_UNNAMED);
+        let src = right_shift(ctx.builder.get(), src, const_8, LLVM_UNNAMED);
+        Ok(src)
+    }
+    let right_shift = match details.type_ {
+        ast::ScalarType::U32 => LLVMBuildLShr,
+        ast::ScalarType::S32 => LLVMBuildAShr,
+        _ => return Err(TranslateError::unreachable()),
+    };
+    let src1 = unsafe { clear_top8_bits(ctx, arg.src1, right_shift)? };
+    let src2 = unsafe { clear_top8_bits(ctx, arg.src2, right_shift)? };
+    match details.control {
+        ast::Mul24Control::Low => {
+            emit_inst_mul_low_impl(ctx, Some(arg.dst), src1, src2, LLVMBuildMul)?
+        }
+        ast::Mul24Control::High => {
+            emit_inst_mul_hi_impl(ctx, details.type_, Some(arg.dst), src1, src2)?
+        }
+    };
+    Ok(())
+}
+
 fn emit_inst_red(
     ctx: &mut EmitContext,
     details: &ast::AtomDetails,
@@ -2751,13 +2792,13 @@ fn emit_inst_mul_hi(
 
 fn emit_inst_mul_hi_impl(
     ctx: &mut EmitContext,
-    type_: ast::ScalarType,
+    narrow_type: ast::ScalarType,
     dst: Option<Id>,
     src1: impl GetLLVMValue,
     src2: impl GetLLVMValue,
 ) -> Result<LLVMValueRef, TranslateError> {
-    let temp_dst = emit_inst_mul_wide(ctx, type_, None, src1, src2)?;
-    emit_get_high_bits(ctx, type_, temp_dst, dst)
+    let temp_dst = emit_inst_mul_wide(ctx, narrow_type, None, src1, src2)?;
+    emit_get_high_bits(ctx, narrow_type, temp_dst, dst)
 }
 
 fn emit_get_high_bits(
diff --git a/ptx/src/ptx.lalrpop b/ptx/src/ptx.lalrpop
index ae57575..3c8a3e5 100644
--- a/ptx/src/ptx.lalrpop
+++ b/ptx/src/ptx.lalrpop
@@ -213,6 +213,7 @@ match {
     "min",

     "mov",

     "mul",

+    "mul24",

     "nanosleep",

     "neg",

     "not",

@@ -294,6 +295,7 @@ ExtendedID : &'input str = {
     "min",

     "mov",

     "mul",

+    "mul24",

     "nanosleep",

     "neg",

     "not",

@@ -839,6 +841,7 @@ Instruction: ast::Instruction<ast::ParsedArgParams<'input>> = {
     InstMatch,

     InstRed,

     InstNanosleep,

+    InstMul24,

 };

 

 // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld

@@ -2099,17 +2102,17 @@ InstTex: ast::Instruction<ast::ParsedArgParams<'input>> = {
 

 // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#surface-instructions-suld

 InstSuld: ast::Instruction<ast::ParsedArgParams<'input>> = {

-    "suld" ".b" <geometry: UnlayeredTextureGeometry> <vector:VectorPrefix?> <type_:AnyBitType> ".trap" <dst:DstOperandVec> "," "[" <image:Operand> "," <coordinates:SrcOperandVec> "]"  => {

+    "suld" ".b" <geometry: UnlayeredTextureGeometry> <vector:VectorPrefix?> <type_:AnyBitType> <clamp:SurfaceClamp> <dst:DstOperandVec> "," "[" <image:Operand> "," <coordinates:SrcOperandVec> "]"  => {

         let args = ast::Arg4Tex {

             dst, 

             image,

             coordinates,

             layer: None,

         };

-        let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, };

+        let details = ast::SurfaceDetails { geometry, vector, type_, clamp, direct: false, };

         ast::Instruction::Suld(details, args)

     },

-    "suld" ".b" ".a1d" <vector:VectorPrefix?> <type_:AnyBitType> ".trap" <dst:DstOperandVec> "," "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "}" "]"  => {

+    "suld" ".b" ".a1d" <vector:VectorPrefix?> <type_:AnyBitType> <clamp:SurfaceClamp> <dst:DstOperandVec> "," "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "}" "]"  => {

         let geometry = ast::TextureGeometry::Array1D;

         let args = ast::Arg4Tex {

             dst,

@@ -2117,10 +2120,10 @@ InstSuld: ast::Instruction<ast::ParsedArgParams<'input>> = {
             coordinates: ast::Operand::VecPack(vec![x]),

             layer: Some(layer),

         };

-        let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, };

+        let details = ast::SurfaceDetails { geometry, vector, type_, clamp, direct: false, };

         ast::Instruction::Suld(details, args)

     },

-    "suld" ".b" ".a2d" <vector:VectorPrefix?> <type_:AnyBitType> ".trap" <dst:DstOperandVec> "," "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "," <y:RegOrImmediate> "," RegOrImmediate "}" "]"  => {

+    "suld" ".b" ".a2d" <vector:VectorPrefix?> <type_:AnyBitType> <clamp:SurfaceClamp> <dst:DstOperandVec> "," "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "," <y:RegOrImmediate> "," RegOrImmediate "}" "]"  => {

         let geometry = ast::TextureGeometry::Array2D;

         let args = ast::Arg4Tex {

             dst,

@@ -2128,24 +2131,30 @@ InstSuld: ast::Instruction<ast::ParsedArgParams<'input>> = {
             coordinates: ast::Operand::VecPack(vec![x, y]),

             layer: Some(layer),

         };

-        let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, };

+        let details = ast::SurfaceDetails { geometry, vector, type_, clamp, direct: false, };

         ast::Instruction::Suld(details, args)

     }

 }

 

+SurfaceClamp: ast::SurfaceClamp = {

+    ".trap" => ast::SurfaceClamp::Trap,

+    ".clamp" => ast::SurfaceClamp::Clamp,

+    ".zero" => ast::SurfaceClamp::Zero,

+}

+

 // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#surface-instructions-sust

 InstSust: ast::Instruction<ast::ParsedArgParams<'input>> = {

-    "sust" ".b" <geometry: UnlayeredTextureGeometry> <vector:VectorPrefix?> <type_:AnyBitType> SustClamp "[" <image:Operand> "," <coordinates:SrcOperandVec> "]" "," <value:SrcOperandVec> => {

+    "sust" ".b" <geometry: UnlayeredTextureGeometry> <vector:VectorPrefix?> <type_:AnyBitType> <clamp:SurfaceClamp> "[" <image:Operand> "," <coordinates:SrcOperandVec> "]" "," <value:SrcOperandVec> => {

         let args = ast::Arg4Sust {

             image,

             coordinates,

             layer: None,

             value,

         };

-        let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, };

+        let details = ast::SurfaceDetails { geometry, vector, type_, clamp, direct: false, };

         ast::Instruction::Sust(details, args)

     },

-    "sust" ".b" ".a1d" <vector:VectorPrefix?> <type_:AnyBitType> SustClamp "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "}" "]" "," <value:SrcOperandVec> => {

+    "sust" ".b" ".a1d" <vector:VectorPrefix?> <type_:AnyBitType> <clamp:SurfaceClamp> "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "}" "]" "," <value:SrcOperandVec> => {

         let geometry = ast::TextureGeometry::Array1D;

         let args = ast::Arg4Sust {

             image,

@@ -2153,10 +2162,10 @@ InstSust: ast::Instruction<ast::ParsedArgParams<'input>> = {
             layer: Some(layer),

             value,

         };

-        let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, };

+        let details = ast::SurfaceDetails { geometry, vector, type_, clamp, direct: false, };

         ast::Instruction::Sust(details, args)

     },

-    "sust" ".b" ".a2d" <vector:VectorPrefix?> <type_:AnyBitType> SustClamp "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "," <y:RegOrImmediate> "," RegOrImmediate "}" "]" "," <value:SrcOperandVec> => {

+    "sust" ".b" ".a2d" <vector:VectorPrefix?> <type_:AnyBitType> <clamp:SurfaceClamp> "[" <image:Operand> "," "{" <layer:SrcOperand> "," <x:RegOrImmediate> "," <y:RegOrImmediate> "," RegOrImmediate "}" "]" "," <value:SrcOperandVec> => {

         let geometry = ast::TextureGeometry::Array2D;

         let args = ast::Arg4Sust {

             image,

@@ -2164,17 +2173,11 @@ InstSust: ast::Instruction<ast::ParsedArgParams<'input>> = {
             layer: Some(layer),

             value,

         };

-        let details = ast::SurfaceDetails { geometry, vector, type_, direct: false, };

+        let details = ast::SurfaceDetails { geometry, vector, type_, clamp, direct: false, };

         ast::Instruction::Sust(details, args)

     }

 }

 

-SustClamp = {

-    ".trap",

-    ".clamp",

-    ".zero"

-}

-

 UnlayeredTextureGeometry: ast::TextureGeometry = {

     ".1d" => ast::TextureGeometry::OneD,

     ".2d" => ast::TextureGeometry::TwoD,

@@ -2244,27 +2247,31 @@ ShiftNormalization: ast::ShiftNormalization = {
 InstVote: ast::Instruction<ast::ParsedArgParams<'input>> = {

     "vote" ".ballot" ".sync" ".b32" <dst:DstOperand> "," <negate:"!"?> <src1:Operand> "," <src2:Operand> => {

         let mode = ast::VoteMode::Ballot;

+        let sync = true;

         let negate_pred = negate.is_some();

         let args = ast::Arg3 {dst, src1, src2};

-        ast::Instruction::Vote(ast::VoteDetails{mode, negate_pred}, args)

+        ast::Instruction::Vote(ast::VoteDetails{mode, sync, negate_pred}, args)

     },

-    "vote" ".sync" ".ballot" ".b32" <dst:DstOperand> "," <negate:"!"?> <src1:Operand> "," <src2:Operand> => {

+    "vote" <sync:".sync"?> ".ballot" ".b32" <dst:DstOperand> "," <negate:"!"?> <src1:Operand> "," <src2:Operand> => {

         let mode = ast::VoteMode::Ballot;

+        let sync = sync.is_some();

         let negate_pred = negate.is_some();

         let args = ast::Arg3 {dst, src1, src2};

-        ast::Instruction::Vote(ast::VoteDetails{mode, negate_pred}, args)

+        ast::Instruction::Vote(ast::VoteDetails{mode, sync, negate_pred}, args)

     },

-    "vote" ".sync" ".any" ".pred" <dst:DstOperand> "," <negate:"!"?> <src1:Operand> "," <src2:Operand> => {

+    "vote" <sync:".sync"?> ".any" ".pred" <dst:DstOperand> "," <negate:"!"?> <src1:Operand> "," <src2:Operand> => {

         let mode = ast::VoteMode::Any;

+        let sync = sync.is_some();

         let negate_pred = negate.is_some();

         let args = ast::Arg3 {dst, src1, src2};

-        ast::Instruction::Vote(ast::VoteDetails{mode, negate_pred}, args)

+        ast::Instruction::Vote(ast::VoteDetails{mode, sync, negate_pred}, args)

     },

-    "vote" ".sync" ".all" ".pred" <dst:DstOperand> "," <negate:"!"?> <src1:Operand> "," <src2:Operand> => {

+    "vote" <sync:".sync"?> ".all" ".pred" <dst:DstOperand> "," <negate:"!"?> <src1:Operand> "," <src2:Operand> => {

         let mode = ast::VoteMode::All;

+        let sync = sync.is_some();

         let negate_pred = negate.is_some();

         let args = ast::Arg3 {dst, src1, src2};

-        ast::Instruction::Vote(ast::VoteDetails{mode, negate_pred}, args)   

+        ast::Instruction::Vote(ast::VoteDetails{mode, sync, negate_pred}, args)   

     }

 }

 

@@ -2372,6 +2379,29 @@ InstNanosleep: ast::Instruction<ast::ParsedArgParams<'input>> = {
     }

 }

 

+// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-mul24

+InstMul24: ast::Instruction<ast::ParsedArgParams<'input>> = {

+    "mul24" <control:Mul24Control> ".u32" <a:Arg3> => {

+        let details = ast::Mul24Details {

+            type_: ast::ScalarType::U32,

+            control

+        };

+        ast::Instruction::Mul24(details, a)

+    },

+    "mul24" <control:Mul24Control> ".s32" <a:Arg3> => {

+        let details = ast::Mul24Details {

+            type_: ast::ScalarType::S32,

+            control

+        };

+        ast::Instruction::Mul24(details, a)

+    }

+}

+

+Mul24Control: ast::Mul24Control = {

+    ".lo" => ast::Mul24Control::Low,

+    ".hi" => ast::Mul24Control::High,

+}

+

 NegTypeFtz: ast::ScalarType = {

     ".f16" => ast::ScalarType::F16,

     ".f16x2" => ast::ScalarType::F16x2,

diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index a65240c..9f7af1a 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -271,7 +271,11 @@ test_ptx!(const, [0u16], [10u16, 20, 30, 40]);
 test_ptx!(cvt_s16_s8, [0x139231C2u32], [0xFFFFFFC2u32]);

 test_ptx!(cvt_f64_f32, [0.125f32], [0.125f64]);

 test_ptx!(cvt_f32_f16, [0xa1u16], [0x37210000u32]);

-test_ptx!(prmt, [0x70c507d6u32, 0x6fbd4b5cu32], [0x6fbdd65cu32, 0x6FFFD600]);

+test_ptx!(

+    prmt,

+    [0x70c507d6u32, 0x6fbd4b5cu32],

+    [0x6fbdd65cu32, 0x6FFFD600]

+);

 test_ptx!(

     prmt_non_immediate,

     [0x70c507d6u32, 0x6fbd4b5cu32],

@@ -336,7 +340,11 @@ test_ptx!(
     [f16::from_f32(2.0), f16::from_f32(3.0)],

     [f16::from_f32(2.0), f16::from_f32(5.0)]

 );

-test_ptx!(set_f16x2, [0xc1690e6eu32, 0x13739444u32, 0x424834CC, 0x4248B4CC], [0xffffu32, 0x3C000000]);

+test_ptx!(

+    set_f16x2,

+    [0xc1690e6eu32, 0x13739444u32, 0x424834CC, 0x4248B4CC],

+    [0xffffu32, 0x3C000000]

+);

 test_ptx!(

     dp4a,

     [0xde3032f5u32, 0x2474fe15, 0xf51d8d6c],

@@ -350,6 +358,7 @@ test_ptx!(
     [1923569713u64, 1923569712],

     [1923569713u64, 1923569712]

 );

+test_ptx!(mul_24_lo, [0xeffa4964u32, 0x46e7e28c], [3336989360u32, 3068553904]);

 

 test_ptx_warp!(

     shfl,

diff --git a/ptx/src/test/spirv_run/mul_24_lo.ll b/ptx/src/test/spirv_run/mul_24_lo.ll
new file mode 100644
index 0000000..12ffae4
--- /dev/null
+++ b/ptx/src/test/spirv_run/mul_24_lo.ll
@@ -0,0 +1,57 @@
+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"
+
+define protected amdgpu_kernel void @mul_24_lo(ptr addrspace(4) byref(i64) %"32", ptr addrspace(4) byref(i64) %"33") #0 {
+"44":
+  %"10" = alloca i1, align 1, addrspace(5)
+  store i1 false, ptr addrspace(5) %"10", align 1
+  %"11" = alloca i1, align 1, addrspace(5)
+  store i1 false, ptr addrspace(5) %"11", align 1
+  %"4" = alloca i64, align 8, addrspace(5)
+  %"5" = alloca i64, align 8, addrspace(5)
+  %"6" = alloca i32, align 4, addrspace(5)
+  %"7" = alloca i32, align 4, addrspace(5)
+  %"8" = alloca i32, align 4, addrspace(5)
+  %"9" = alloca i32, align 4, addrspace(5)
+  %"12" = load i64, ptr addrspace(4) %"32", align 8
+  store i64 %"12", ptr addrspace(5) %"4", align 8
+  %"13" = load i64, ptr addrspace(4) %"33", align 8
+  store i64 %"13", ptr addrspace(5) %"5", align 8
+  %"15" = load i64, ptr addrspace(5) %"4", align 8
+  %"34" = inttoptr i64 %"15" to ptr
+  %"14" = load i32, ptr %"34", align 4
+  store i32 %"14", ptr addrspace(5) %"6", align 4
+  %"17" = load i64, ptr addrspace(5) %"4", align 8
+  %"35" = inttoptr i64 %"17" to ptr
+  %"46" = getelementptr inbounds i8, ptr %"35", i64 4
+  %"16" = load i32, ptr %"46", align 4
+  store i32 %"16", ptr addrspace(5) %"7", align 4
+  %"19" = load i32, ptr addrspace(5) %"6", align 4
+  %"20" = load i32, ptr addrspace(5) %"7", align 4
+  %0 = shl i32 %"19", 8
+  %1 = ashr i32 %0, 8
+  %2 = shl i32 %"20", 8
+  %3 = ashr i32 %2, 8
+  %"36" = mul i32 %1, %3
+  store i32 %"36", ptr addrspace(5) %"8", align 4
+  %"22" = load i32, ptr addrspace(5) %"6", align 4
+  %"23" = load i32, ptr addrspace(5) %"7", align 4
+  %4 = shl i32 %"22", 8
+  %5 = lshr i32 %4, 8
+  %6 = shl i32 %"23", 8
+  %7 = lshr i32 %6, 8
+  %"39" = mul i32 %5, %7
+  store i32 %"39", ptr addrspace(5) %"9", align 4
+  %"24" = load i64, ptr addrspace(5) %"5", align 8
+  %"25" = load i32, ptr addrspace(5) %"8", align 4
+  %"42" = inttoptr i64 %"24" to ptr
+  store i32 %"25", ptr %"42", align 4
+  %"26" = load i64, ptr addrspace(5) %"5", align 8
+  %"27" = load i32, ptr addrspace(5) %"9", align 4
+  %"43" = inttoptr i64 %"26" to ptr
+  %"48" = getelementptr inbounds i8, ptr %"43", i64 4
+  store i32 %"27", ptr %"48", align 4
+  ret void
+}
+
+attributes #0 = { "amdgpu-unsafe-fp-atomics"="true" "denormal-fp-math"="ieee,ieee" "denormal-fp-math-f32"="ieee,ieee" "no-trapping-math"="true" "uniform-work-group-size"="true" }
diff --git a/ptx/src/test/spirv_run/mul_24_lo.ptx b/ptx/src/test/spirv_run/mul_24_lo.ptx
new file mode 100644
index 0000000..9dcbafc
--- /dev/null
+++ b/ptx/src/test/spirv_run/mul_24_lo.ptx
@@ -0,0 +1,27 @@
+.version 6.5

+.target sm_30

+.address_size 64

+

+.visible .entry mul_24_lo(

+	.param .u64 input,

+	.param .u64 output

+)

+{

+	.reg .u64 	    in_addr;

+    .reg .u64 	    out_addr;

+    .reg .b32 	    temp1;

+    .reg .b32 	    temp2;

+    .reg .b32 	    result1;

+    .reg .b32 	    result2;

+

+	ld.param.u64 	in_addr, [input];

+    ld.param.u64 	out_addr, [output];

+

+    ld.b32          temp1, [in_addr];

+    ld.b32          temp2, [in_addr+4];

+	mul24.lo.s32	result1, temp1, temp2;

+	mul24.lo.u32	result2, temp1, temp2;

+    st.b32          [out_addr], result1;

+    st.b32          [out_addr+4], result2;

+	ret;

+}

diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs
index 041c690..3b3de1f 100644
--- a/ptx/src/translate.rs
+++ b/ptx/src/translate.rs
@@ -2653,6 +2653,7 @@ fn replace_instructions_with_builtins_impl<'input>(
             Statement::Instruction(ast::Instruction::Vote(

                 ast::VoteDetails {

                     mode: ast::VoteMode::Any,

+                    sync,

                     negate_pred,

                 },

                 arg,

@@ -2672,6 +2673,7 @@ fn replace_instructions_with_builtins_impl<'input>(
                     ast::Instruction::Vote(

                         ast::VoteDetails {

                             mode: ast::VoteMode::Any,

+                            sync,

                             negate_pred,

                         },

                         arg,

@@ -2682,6 +2684,7 @@ fn replace_instructions_with_builtins_impl<'input>(
             Statement::Instruction(ast::Instruction::Vote(

                 ast::VoteDetails {

                     mode: ast::VoteMode::All,

+                    sync,

                     negate_pred,

                 },

                 arg,

@@ -2701,6 +2704,7 @@ fn replace_instructions_with_builtins_impl<'input>(
                     ast::Instruction::Vote(

                         ast::VoteDetails {

                             mode: ast::VoteMode::All,

+                            sync,

                             negate_pred,

                         },

                         arg,

@@ -2711,6 +2715,7 @@ fn replace_instructions_with_builtins_impl<'input>(
             Statement::Instruction(ast::Instruction::Vote(

                 ast::VoteDetails {

                     mode: ast::VoteMode::Ballot,

+                    sync,

                     negate_pred,

                 },

                 arg,

@@ -2730,6 +2735,7 @@ fn replace_instructions_with_builtins_impl<'input>(
                     ast::Instruction::Vote(

                         ast::VoteDetails {

                             mode: ast::VoteMode::Ballot,

+                            sync,

                             negate_pred,

                         },

                         arg,

@@ -6594,6 +6600,10 @@ impl<T: ArgParamsEx> ast::Instruction<T> {
                     ast::StateSpace::Reg,

                 )),

             )?),

+            ast::Instruction::Mul24(details, args) => ast::Instruction::Mul24(

+                details,

+                args.map_generic(visitor, &ast::Type::Scalar(details.type_), false)?,

+            ),

         })

     }

 }

@@ -6915,6 +6925,7 @@ impl<T: ast::ArgParams> ast::Instruction<T> {
             ast::Instruction::Vshr { .. } => None,

             ast::Instruction::Dp4a { .. } => None,

             ast::Instruction::MatchAny { .. } => None,

+            ast::Instruction::Mul24 { .. } => None,

             ast::Instruction::Sub(ast::ArithDetails::Signed(_), _) => None,

             ast::Instruction::Sub(ast::ArithDetails::Unsigned(_), _) => None,

             ast::Instruction::Add(ast::ArithDetails::Signed(_), _) => None,

diff --git a/zluda_dump/src/dark_api.rs b/zluda_dump/src/dark_api.rs
index a619bca..72e4198 100644
--- a/zluda_dump/src/dark_api.rs
+++ b/zluda_dump/src/dark_api.rs
@@ -432,13 +432,42 @@ impl CudaDarkApiDump for CudaDarkApiDumpFns {
     }

 

     unsafe fn ctx_create_v2_bypass_impl(

-        _guid: &[u8; 16],

-        _idx: usize,

-        _pctx: *mut cuda_types::CUcontext,

-        _flags: c_uint,

-        _dev: cuda_types::CUdevice,

+        guid: &[u8; 16],

+        idx: usize,

+        pctx: *mut cuda_types::CUcontext,

+        flags: c_uint,

+        dev: cuda_types::CUdevice,

     ) -> CUresult {

-        todo!()

+        let arguments_writer = Box::new(move |writer: &mut dyn std::io::Write| {

+            writer.write_all(b"(pctx: ")?;

+            format::CudaDisplay::write(&pctx, "", 0, writer)?;

+            writer.write_all(b", flags: ")?;

+            format::CudaDisplay::write(&flags, "", 0, writer)?;

+            writer.write_all(b", dev: ")?;

+            format::CudaDisplay::write(&dev, "", 0, writer)?;

+            writer.write_all(b")")

+        });

+        let global_state = &mut *super::GLOBAL_STATE.lock().unwrap();

+        let mut fn_logger = global_state.log_factory.get_logger_dark_api(

+            CUuuid {

+                bytes: guid.clone(),

+            },

+            idx,

+            Some(arguments_writer),

+        );

+        let cuda_state = &mut global_state.delayed_state.unwrap_mut().cuda_state;

+        let original_ptr = cuda_state.dark_api.overrides[guid].1.add(idx);

+        let original_fn = mem::transmute::<

+            _,

+            unsafe extern "system" fn(

+                *mut cuda_types::CUcontext,

+                c_uint,

+                cuda_types::CUdevice,

+            ) -> CUresult,

+        >(*original_ptr);

+        let original_result = original_fn(pctx, flags, dev);

+        fn_logger.result = Some(original_result);

+        original_result

     }

 

     unsafe fn dlss_feature_evaluate_init_impl(