diff --git a/ptx/src/ast.rs b/ptx/src/ast.rs index 7f2fc9a..653060b 100644 --- a/ptx/src/ast.rs +++ b/ptx/src/ast.rs @@ -113,6 +113,8 @@ sub_type! { VariableRegType { Scalar(ScalarType), Vector(SizedScalarType, u8), + // Array type is used when emiting SSA statements at the start of a method + Array(ScalarType, VecU32), // Pointer variant is used when passing around SLM pointer between // function calls for dynamic SLM Pointer(SizedScalarType, PointerStateSpace) @@ -267,7 +269,6 @@ pub enum MethodDecl<'a, ID> { Kernel { name: &'a str, in_args: Vec>, - uses_shared_mem: bool, }, } @@ -293,16 +294,51 @@ pub enum KernelArgumentType { Shared, } -impl From for Type { - fn from(t: FnArgumentType) -> Self { - match t { - FnArgumentType::Reg(x) => x.into(), - FnArgumentType::Param(x) => x.into(), +impl FnArgumentType { + pub fn to_type(&self, is_kernel: bool) -> Type { + if is_kernel { + self.to_kernel_type() + } else { + self.to_func_type() + } + } + + pub fn to_kernel_type(&self) -> Type { + match self { + FnArgumentType::Reg(x) => x.clone().into(), + FnArgumentType::Param(x) => x.clone().into(), FnArgumentType::Shared => { Type::Pointer(PointerType::Scalar(ScalarType::B8), LdStateSpace::Shared) } } } + + pub fn to_func_type(&self) -> Type { + match self { + FnArgumentType::Reg(x) => x.clone().into(), + FnArgumentType::Param(VariableParamType::Scalar(t)) => { + Type::Pointer(PointerType::Scalar((*t).into()), LdStateSpace::Param) + } + FnArgumentType::Param(VariableParamType::Array(t, dims)) => Type::Pointer( + PointerType::Array((*t).into(), dims.clone()), + LdStateSpace::Param, + ), + FnArgumentType::Param(VariableParamType::Pointer(t, space)) => Type::Pointer( + PointerType::Pointer((*t).into(), (*space).into()), + LdStateSpace::Param, + ), + FnArgumentType::Shared => { + Type::Pointer(PointerType::Scalar(ScalarType::B8), LdStateSpace::Shared) + } + } + } + + pub fn is_param(&self) -> bool { + match self { + FnArgumentType::Param(_) => true, + _ => false, + } + } } sub_enum!( @@ -323,11 +359,12 @@ pub enum Type { Pointer(PointerType, LdStateSpace), } -sub_type! { - PointerType { - Scalar(ScalarType), - Vector(ScalarType, u8), - } +#[derive(PartialEq, Eq, Clone)] +pub enum PointerType { + Scalar(ScalarType), + Vector(ScalarType, u8), + Array(ScalarType, VecU32), + Pointer(ScalarType, LdStateSpace), } impl From for PointerType { @@ -343,6 +380,8 @@ impl TryFrom for SizedScalarType { match value { PointerType::Scalar(t) => Ok(t.try_into()?), PointerType::Vector(_, _) => Err(()), + PointerType::Array(_, _) => Err(()), + PointerType::Pointer(_, _) => Err(()), } } } @@ -456,6 +495,7 @@ pub struct MultiVariable { pub count: Option, } +#[derive(Clone)] pub struct Variable { pub align: Option, pub v_type: T, @@ -543,6 +583,10 @@ pub enum Instruction { Sqrt(SqrtDetails, Arg2

), Rsqrt(RsqrtDetails, Arg2

), Neg(NegDetails, Arg2

), + Sin { flush_to_zero: bool, arg: Arg2

}, + Cos { flush_to_zero: bool, arg: Arg2

}, + Lg2 { flush_to_zero: bool, arg: Arg2

}, + Ex2 { flush_to_zero: bool, arg: Arg2

}, } #[derive(Copy, Clone)] @@ -744,6 +788,7 @@ pub enum MemScope { } #[derive(Copy, Clone, PartialEq, Eq, Debug)] +#[repr(u8)] pub enum LdStateSpace { Generic, Const, diff --git a/ptx/src/ptx.lalrpop b/ptx/src/ptx.lalrpop index 9d2adec..584ef84 100644 --- a/ptx/src/ptx.lalrpop +++ b/ptx/src/ptx.lalrpop @@ -144,12 +144,15 @@ match { "barrier", "bra", "call", + "cos", "cvt", "cvta", "debug", "div", + "ex2", "fma", "ld", + "lg2", "mad", "map_f64_to_f32", "max", @@ -166,6 +169,7 @@ match { "setp", "shl", "shr", + "sin", r"sm_[0-9]+" => ShaderModel, "sqrt", "st", @@ -187,12 +191,15 @@ ExtendedID : &'input str = { "barrier", "bra", "call", + "cos", "cvt", "cvta", "debug", "div", + "ex2", "fma", "ld", + "lg2", "mad", "map_f64_to_f32", "max", @@ -209,6 +216,7 @@ ExtendedID : &'input str = { "setp", "shl", "shr", + "sin", ShaderModel, "sqrt", "st", @@ -346,7 +354,7 @@ LinkingDirectives: ast::LinkingDirective = { MethodDecl: ast::MethodDecl<'input, &'input str> = { ".entry" => - ast::MethodDecl::Kernel{ name, in_args, uses_shared_mem: false }, + ast::MethodDecl::Kernel{ name, in_args }, ".func" => { ast::MethodDecl::Func(ret_vals.unwrap_or_else(|| Vec::new()), name, params) } @@ -687,6 +695,10 @@ Instruction: ast::Instruction> = { InstSqrt, InstRsqrt, InstNeg, + InstSin, + InstCos, + InstLg2, + InstEx2, }; // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld @@ -1600,6 +1612,34 @@ InstNeg: ast::Instruction> = { }, } +// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-sin +InstSin: ast::Instruction> = { + "sin" ".approx" ".f32" => { + ast::Instruction::Sin{ flush_to_zero: ftz.is_some(), arg } + }, +} + +// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-cos +InstCos: ast::Instruction> = { + "cos" ".approx" ".f32" => { + ast::Instruction::Cos{ flush_to_zero: ftz.is_some(), arg } + }, +} + +// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-lg2 +InstLg2: ast::Instruction> = { + "lg2" ".approx" ".f32" => { + ast::Instruction::Lg2{ flush_to_zero: ftz.is_some(), arg } + }, +} + +// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-ex2 +InstEx2: ast::Instruction> = { + "ex2" ".approx" ".f32" => { + ast::Instruction::Ex2{ flush_to_zero: ftz.is_some(), arg } + }, +} + NegTypeFtz: ast::ScalarType = { ".f16" => ast::ScalarType::F16, ".f16x2" => ast::ScalarType::F16x2, diff --git a/ptx/src/test/spirv_run/add.spvtxt b/ptx/src/test/spirv_run/add.spvtxt index 6810fec..d9a5b9e 100644 --- a/ptx/src/test/spirv_run/add.spvtxt +++ b/ptx/src/test/spirv_run/add.spvtxt @@ -2,21 +2,24 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %25 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %23 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "add" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %28 = OpTypeFunction %void %ulong %ulong + %26 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_1 = OpConstant %ulong 1 - %1 = OpFunction %void None %28 + %1 = OpFunction %void None %26 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %23 = OpLabel + %21 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -25,22 +28,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpIAdd %ulong %17 %ulong_1 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %22 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpIAdd %ulong %15 %ulong_1 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %20 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/and.spvtxt b/ptx/src/test/spirv_run/and.spvtxt index 8358c28..57c36e6 100644 --- a/ptx/src/test/spirv_run/and.spvtxt +++ b/ptx/src/test/spirv_run/and.spvtxt @@ -1,66 +1,58 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 41 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%33 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "and" -%34 = OpTypeVoid -%35 = OpTypeInt 64 0 -%36 = OpTypeFunction %34 %35 %35 -%37 = OpTypePointer Function %35 -%38 = OpTypeInt 32 0 -%39 = OpTypePointer Function %38 -%40 = OpTypePointer Generic %38 -%23 = OpConstant %35 4 -%1 = OpFunction %34 None %36 -%8 = OpFunctionParameter %35 -%9 = OpFunctionParameter %35 -%31 = OpLabel -%2 = OpVariable %37 Function -%3 = OpVariable %37 Function -%4 = OpVariable %37 Function -%5 = OpVariable %37 Function -%6 = OpVariable %39 Function -%7 = OpVariable %39 Function -OpStore %2 %8 -OpStore %3 %9 -%11 = OpLoad %35 %2 -%10 = OpCopyObject %35 %11 -OpStore %4 %10 -%13 = OpLoad %35 %3 -%12 = OpCopyObject %35 %13 -OpStore %5 %12 -%15 = OpLoad %35 %4 -%25 = OpConvertUToPtr %40 %15 -%14 = OpLoad %38 %25 -OpStore %6 %14 -%17 = OpLoad %35 %4 -%24 = OpIAdd %35 %17 %23 -%26 = OpConvertUToPtr %40 %24 -%16 = OpLoad %38 %26 -OpStore %7 %16 -%19 = OpLoad %38 %6 -%20 = OpLoad %38 %7 -%28 = OpCopyObject %38 %19 -%29 = OpCopyObject %38 %20 -%27 = OpBitwiseAnd %38 %28 %29 -%18 = OpCopyObject %38 %27 -OpStore %6 %18 -%21 = OpLoad %35 %5 -%22 = OpLoad %38 %6 -%30 = OpConvertUToPtr %40 %21 -OpStore %30 %22 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %31 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "and" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %34 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %ulong_4 = OpConstant %ulong 4 + %1 = OpFunction %void None %34 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %29 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_uint Function + %7 = OpVariable %_ptr_Function_uint Function + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %4 %10 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %23 + OpStore %6 %12 + %15 = OpLoad %ulong %4 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_uint %22 + %14 = OpLoad %uint %24 + OpStore %7 %14 + %17 = OpLoad %uint %6 + %18 = OpLoad %uint %7 + %26 = OpCopyObject %uint %17 + %27 = OpCopyObject %uint %18 + %25 = OpBitwiseAnd %uint %26 %27 + %16 = OpCopyObject %uint %25 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %6 + %28 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %28 %20 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/atom_add.spvtxt b/ptx/src/test/spirv_run/atom_add.spvtxt index 2c83fe9..6a977e0 100644 --- a/ptx/src/test/spirv_run/atom_add.spvtxt +++ b/ptx/src/test/spirv_run/atom_add.spvtxt @@ -1,84 +1,76 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 55 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%40 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "atom_add" %4 -OpDecorate %4 Alignment 4 -%41 = OpTypeVoid -%42 = OpTypeInt 32 0 -%43 = OpTypeInt 8 0 -%44 = OpConstant %42 1024 -%45 = OpTypeArray %43 %44 -%46 = OpTypePointer Workgroup %45 -%4 = OpVariable %46 Workgroup -%47 = OpTypeInt 64 0 -%48 = OpTypeFunction %41 %47 %47 -%49 = OpTypePointer Function %47 -%50 = OpTypePointer Function %42 -%51 = OpTypePointer Generic %42 -%27 = OpConstant %47 4 -%52 = OpTypePointer Workgroup %42 -%53 = OpConstant %42 1 -%54 = OpConstant %42 0 -%29 = OpConstant %47 4 -%1 = OpFunction %41 None %48 -%9 = OpFunctionParameter %47 -%10 = OpFunctionParameter %47 -%38 = OpLabel -%2 = OpVariable %49 Function -%3 = OpVariable %49 Function -%5 = OpVariable %49 Function -%6 = OpVariable %49 Function -%7 = OpVariable %50 Function -%8 = OpVariable %50 Function -OpStore %2 %9 -OpStore %3 %10 -%12 = OpLoad %47 %2 -%11 = OpCopyObject %47 %12 -OpStore %5 %11 -%14 = OpLoad %47 %3 -%13 = OpCopyObject %47 %14 -OpStore %6 %13 -%16 = OpLoad %47 %5 -%31 = OpConvertUToPtr %51 %16 -%15 = OpLoad %42 %31 -OpStore %7 %15 -%18 = OpLoad %47 %5 -%28 = OpIAdd %47 %18 %27 -%32 = OpConvertUToPtr %51 %28 -%17 = OpLoad %42 %32 -OpStore %8 %17 -%19 = OpLoad %42 %7 -%33 = OpBitcast %52 %4 -OpStore %33 %19 -%21 = OpLoad %42 %8 -%34 = OpBitcast %52 %4 -%20 = OpAtomicIAdd %42 %34 %53 %54 %21 -OpStore %7 %20 -%35 = OpBitcast %52 %4 -%22 = OpLoad %42 %35 -OpStore %8 %22 -%23 = OpLoad %47 %6 -%24 = OpLoad %42 %7 -%36 = OpConvertUToPtr %51 %23 -OpStore %36 %24 -%25 = OpLoad %47 %6 -%26 = OpLoad %42 %8 -%30 = OpIAdd %47 %25 %29 -%37 = OpConvertUToPtr %51 %30 -OpStore %37 %26 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %38 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "atom_add" %4 + OpDecorate %4 Alignment 4 + %void = OpTypeVoid + %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 + %uint_1024 = OpConstant %uint 1024 +%_arr_uchar_uint_1024 = OpTypeArray %uchar %uint_1024 +%_ptr_Workgroup__arr_uchar_uint_1024 = OpTypePointer Workgroup %_arr_uchar_uint_1024 + %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup + %ulong = OpTypeInt 64 0 + %46 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %ulong_4 = OpConstant %ulong 4 +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %uint_1 = OpConstant %uint 1 + %uint_0 = OpConstant %uint 0 + %ulong_4_0 = OpConstant %ulong 4 + %1 = OpFunction %void None %46 + %9 = OpFunctionParameter %ulong + %10 = OpFunctionParameter %ulong + %36 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_ulong Function + %7 = OpVariable %_ptr_Function_uint Function + %8 = OpVariable %_ptr_Function_uint Function + OpStore %2 %9 + OpStore %3 %10 + %11 = OpLoad %ulong %2 + OpStore %5 %11 + %12 = OpLoad %ulong %3 + OpStore %6 %12 + %14 = OpLoad %ulong %5 + %29 = OpConvertUToPtr %_ptr_Generic_uint %14 + %13 = OpLoad %uint %29 + OpStore %7 %13 + %16 = OpLoad %ulong %5 + %26 = OpIAdd %ulong %16 %ulong_4 + %30 = OpConvertUToPtr %_ptr_Generic_uint %26 + %15 = OpLoad %uint %30 + OpStore %8 %15 + %17 = OpLoad %uint %7 + %31 = OpBitcast %_ptr_Workgroup_uint %4 + OpStore %31 %17 + %19 = OpLoad %uint %8 + %32 = OpBitcast %_ptr_Workgroup_uint %4 + %18 = OpAtomicIAdd %uint %32 %uint_1 %uint_0 %19 + OpStore %7 %18 + %33 = OpBitcast %_ptr_Workgroup_uint %4 + %20 = OpLoad %uint %33 + OpStore %8 %20 + %21 = OpLoad %ulong %6 + %22 = OpLoad %uint %7 + %34 = OpConvertUToPtr %_ptr_Generic_uint %21 + OpStore %34 %22 + %23 = OpLoad %ulong %6 + %24 = OpLoad %uint %8 + %28 = OpIAdd %ulong %23 %ulong_4_0 + %35 = OpConvertUToPtr %_ptr_Generic_uint %28 + OpStore %35 %24 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/atom_cas.spvtxt b/ptx/src/test/spirv_run/atom_cas.spvtxt index c5fb922..b28d3bc 100644 --- a/ptx/src/test/spirv_run/atom_cas.spvtxt +++ b/ptx/src/test/spirv_run/atom_cas.spvtxt @@ -1,77 +1,69 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 51 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%41 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "atom_cas" -%42 = OpTypeVoid -%43 = OpTypeInt 64 0 -%44 = OpTypeFunction %42 %43 %43 -%45 = OpTypePointer Function %43 -%46 = OpTypeInt 32 0 -%47 = OpTypePointer Function %46 -%48 = OpTypePointer Generic %46 -%25 = OpConstant %43 4 -%27 = OpConstant %46 100 -%49 = OpConstant %46 1 -%50 = OpConstant %46 0 -%28 = OpConstant %43 4 -%30 = OpConstant %43 4 -%1 = OpFunction %42 None %44 -%8 = OpFunctionParameter %43 -%9 = OpFunctionParameter %43 -%39 = OpLabel -%2 = OpVariable %45 Function -%3 = OpVariable %45 Function -%4 = OpVariable %45 Function -%5 = OpVariable %45 Function -%6 = OpVariable %47 Function -%7 = OpVariable %47 Function -OpStore %2 %8 -OpStore %3 %9 -%11 = OpLoad %43 %2 -%10 = OpCopyObject %43 %11 -OpStore %4 %10 -%13 = OpLoad %43 %3 -%12 = OpCopyObject %43 %13 -OpStore %5 %12 -%15 = OpLoad %43 %4 -%32 = OpConvertUToPtr %48 %15 -%14 = OpLoad %46 %32 -OpStore %6 %14 -%17 = OpLoad %43 %4 -%18 = OpLoad %46 %6 -%26 = OpIAdd %43 %17 %25 -%34 = OpConvertUToPtr %48 %26 -%35 = OpCopyObject %46 %18 -%33 = OpAtomicCompareExchange %46 %34 %49 %50 %50 %27 %35 -%16 = OpCopyObject %46 %33 -OpStore %6 %16 -%20 = OpLoad %43 %4 -%29 = OpIAdd %43 %20 %28 -%36 = OpConvertUToPtr %48 %29 -%19 = OpLoad %46 %36 -OpStore %7 %19 -%21 = OpLoad %43 %5 -%22 = OpLoad %46 %6 -%37 = OpConvertUToPtr %48 %21 -OpStore %37 %22 -%23 = OpLoad %43 %5 -%24 = OpLoad %46 %7 -%31 = OpIAdd %43 %23 %30 -%38 = OpConvertUToPtr %48 %31 -OpStore %38 %24 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %39 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "atom_cas" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %42 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %ulong_4 = OpConstant %ulong 4 + %uint_100 = OpConstant %uint 100 + %uint_1 = OpConstant %uint 1 + %uint_0 = OpConstant %uint 0 + %ulong_4_0 = OpConstant %ulong 4 + %ulong_4_1 = OpConstant %ulong 4 + %1 = OpFunction %void None %42 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %37 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_uint Function + %7 = OpVariable %_ptr_Function_uint Function + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %4 %10 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %30 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %30 + OpStore %6 %12 + %15 = OpLoad %ulong %4 + %16 = OpLoad %uint %6 + %24 = OpIAdd %ulong %15 %ulong_4 + %32 = OpConvertUToPtr %_ptr_Generic_uint %24 + %33 = OpCopyObject %uint %16 + %31 = OpAtomicCompareExchange %uint %32 %uint_1 %uint_0 %uint_0 %uint_100 %33 + %14 = OpCopyObject %uint %31 + OpStore %6 %14 + %18 = OpLoad %ulong %4 + %27 = OpIAdd %ulong %18 %ulong_4_0 + %34 = OpConvertUToPtr %_ptr_Generic_uint %27 + %17 = OpLoad %uint %34 + OpStore %7 %17 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %6 + %35 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %35 %20 + %21 = OpLoad %ulong %5 + %22 = OpLoad %uint %7 + %29 = OpIAdd %ulong %21 %ulong_4_1 + %36 = OpConvertUToPtr %_ptr_Generic_uint %29 + OpStore %36 %22 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/b64tof64.spvtxt b/ptx/src/test/spirv_run/b64tof64.spvtxt index 9146c90..e8cfcf4 100644 --- a/ptx/src/test/spirv_run/b64tof64.spvtxt +++ b/ptx/src/test/spirv_run/b64tof64.spvtxt @@ -2,23 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %26 = OpExtInstImport "OpenCL.std" + %24 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "b64tof64" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %29 = OpTypeFunction %void %ulong %ulong + %27 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %double = OpTypeFloat 64 %_ptr_Function_double = OpTypePointer Function %double %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %29 + %1 = OpFunction %void None %27 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %24 = OpLabel + %22 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_double Function @@ -27,24 +29,22 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %20 = OpBitcast %double %11 - %10 = OpCopyObject %double %20 + %18 = OpBitcast %_ptr_Function_double %2 + %10 = OpLoad %double %18 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %6 %12 - %15 = OpLoad %double %4 - %21 = OpBitcast %ulong %15 - %14 = OpCopyObject %ulong %21 - OpStore %5 %14 - %17 = OpLoad %ulong %5 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %17 - %16 = OpLoad %ulong %22 - OpStore %7 %16 - %18 = OpLoad %ulong %6 - %19 = OpLoad %ulong %7 - %23 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %23 %19 + %11 = OpLoad %ulong %3 + OpStore %6 %11 + %13 = OpLoad %double %4 + %19 = OpBitcast %ulong %13 + %12 = OpCopyObject %ulong %19 + OpStore %5 %12 + %15 = OpLoad %ulong %5 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %15 + %14 = OpLoad %ulong %20 + OpStore %7 %14 + %16 = OpLoad %ulong %6 + %17 = OpLoad %ulong %7 + %21 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %21 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/block.spvtxt b/ptx/src/test/spirv_run/block.spvtxt index 534167d..fe7e63a 100644 --- a/ptx/src/test/spirv_run/block.spvtxt +++ b/ptx/src/test/spirv_run/block.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %29 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %27 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "block" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %32 = OpTypeFunction %void %ulong %ulong + %30 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_1 = OpConstant %ulong 1 %ulong_1_0 = OpConstant %ulong 1 - %1 = OpFunction %void None %32 + %1 = OpFunction %void None %30 %9 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong - %27 = OpLabel + %25 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -27,25 +30,23 @@ %8 = OpVariable %_ptr_Function_ulong Function OpStore %2 %9 OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %2 OpStore %4 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %5 %13 - %16 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_ulong %16 - %15 = OpLoad %ulong %25 - OpStore %6 %15 - %18 = OpLoad %ulong %6 - %17 = OpIAdd %ulong %18 %ulong_1 - OpStore %7 %17 - %20 = OpLoad %ulong %8 - %19 = OpIAdd %ulong %20 %ulong_1_0 - OpStore %8 %19 - %21 = OpLoad %ulong %5 - %22 = OpLoad %ulong %7 - %26 = OpConvertUToPtr %_ptr_Generic_ulong %21 - OpStore %26 %22 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_ulong %14 + %13 = OpLoad %ulong %23 + OpStore %6 %13 + %16 = OpLoad %ulong %6 + %15 = OpIAdd %ulong %16 %ulong_1 + OpStore %7 %15 + %18 = OpLoad %ulong %8 + %17 = OpIAdd %ulong %18 %ulong_1_0 + OpStore %8 %17 + %19 = OpLoad %ulong %5 + %20 = OpLoad %ulong %7 + %24 = OpConvertUToPtr %_ptr_Generic_ulong %19 + OpStore %24 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/bra.spvtxt b/ptx/src/test/spirv_run/bra.spvtxt index f59fda5..b20e61a 100644 --- a/ptx/src/test/spirv_run/bra.spvtxt +++ b/ptx/src/test/spirv_run/bra.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %31 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %29 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "bra" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %34 = OpTypeFunction %void %ulong %ulong + %32 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_1 = OpConstant %ulong 1 %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %34 + %1 = OpFunction %void None %32 %11 = OpFunctionParameter %ulong %12 = OpFunctionParameter %ulong - %29 = OpLabel + %27 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function @@ -26,31 +29,29 @@ %10 = OpVariable %_ptr_Function_ulong Function OpStore %2 %11 OpStore %3 %12 - %14 = OpLoad %ulong %2 - %13 = OpCopyObject %ulong %14 + %13 = OpLoad %ulong %2 OpStore %7 %13 - %16 = OpLoad %ulong %3 - %15 = OpCopyObject %ulong %16 - OpStore %8 %15 - %18 = OpLoad %ulong %7 - %27 = OpConvertUToPtr %_ptr_Generic_ulong %18 - %17 = OpLoad %ulong %27 - OpStore %9 %17 + %14 = OpLoad %ulong %3 + OpStore %8 %14 + %16 = OpLoad %ulong %7 + %25 = OpConvertUToPtr %_ptr_Generic_ulong %16 + %15 = OpLoad %ulong %25 + OpStore %9 %15 OpBranch %4 %4 = OpLabel + %18 = OpLoad %ulong %9 + %17 = OpIAdd %ulong %18 %ulong_1 + OpStore %10 %17 + OpBranch %6 + %35 = OpLabel %20 = OpLoad %ulong %9 - %19 = OpIAdd %ulong %20 %ulong_1 + %19 = OpIAdd %ulong %20 %ulong_2 OpStore %10 %19 OpBranch %6 - %37 = OpLabel - %22 = OpLoad %ulong %9 - %21 = OpIAdd %ulong %22 %ulong_2 - OpStore %10 %21 - OpBranch %6 %6 = OpLabel - %23 = OpLoad %ulong %8 - %24 = OpLoad %ulong %10 - %28 = OpConvertUToPtr %_ptr_Generic_ulong %23 - OpStore %28 %24 + %21 = OpLoad %ulong %8 + %22 = OpLoad %ulong %10 + %26 = OpConvertUToPtr %_ptr_Generic_ulong %21 + OpStore %26 %22 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/call.spvtxt b/ptx/src/test/spirv_run/call.spvtxt index ca4685a..31f5307 100644 --- a/ptx/src/test/spirv_run/call.spvtxt +++ b/ptx/src/test/spirv_run/call.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %47 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %37 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %4 "call" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %50 = OpTypeFunction %void %ulong %ulong + %40 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong - %53 = OpTypeFunction %ulong %ulong + %44 = OpTypeFunction %void %_ptr_Function_ulong %_ptr_Function_ulong %ulong_1 = OpConstant %ulong 1 - %4 = OpFunction %void None %50 + %4 = OpFunction %void None %40 %12 = OpFunctionParameter %ulong %13 = OpFunctionParameter %ulong - %32 = OpLabel + %26 = OpLabel %5 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function @@ -27,49 +30,38 @@ %11 = OpVariable %_ptr_Function_ulong Function OpStore %5 %12 OpStore %6 %13 - %15 = OpLoad %ulong %5 - %14 = OpCopyObject %ulong %15 + %14 = OpLoad %ulong %5 OpStore %7 %14 - %17 = OpLoad %ulong %6 - %16 = OpCopyObject %ulong %17 - OpStore %8 %16 - %19 = OpLoad %ulong %7 - %28 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %19 - %18 = OpLoad %ulong %28 - OpStore %9 %18 + %15 = OpLoad %ulong %6 + OpStore %8 %15 + %17 = OpLoad %ulong %7 + %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %17 + %16 = OpLoad %ulong %22 + OpStore %9 %16 + %18 = OpLoad %ulong %9 + %23 = OpBitcast %_ptr_Function_ulong %10 + %24 = OpCopyObject %ulong %18 + OpStore %23 %24 + %43 = OpFunctionCall %void %1 %11 %10 + %19 = OpLoad %ulong %11 + OpStore %9 %19 + %20 = OpLoad %ulong %8 %21 = OpLoad %ulong %9 - %29 = OpCopyObject %ulong %21 - %30 = OpCopyObject %ulong %29 - %20 = OpCopyObject %ulong %30 - OpStore %10 %20 - %23 = OpLoad %ulong %10 - %22 = OpFunctionCall %ulong %1 %23 - OpStore %11 %22 - %25 = OpLoad %ulong %11 - %24 = OpCopyObject %ulong %25 - OpStore %9 %24 - %26 = OpLoad %ulong %8 - %27 = OpLoad %ulong %9 - %31 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %26 - OpStore %31 %27 + %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %20 + OpStore %25 %21 OpReturn OpFunctionEnd - %1 = OpFunction %ulong None %53 - %36 = OpFunctionParameter %ulong - %45 = OpLabel - %34 = OpVariable %_ptr_Function_ulong Function - %33 = OpVariable %_ptr_Function_ulong Function - %35 = OpVariable %_ptr_Function_ulong Function - OpStore %34 %36 - %38 = OpLoad %ulong %34 - %37 = OpCopyObject %ulong %38 - OpStore %35 %37 - %40 = OpLoad %ulong %35 - %39 = OpIAdd %ulong %40 %ulong_1 - OpStore %35 %39 - %42 = OpLoad %ulong %35 - %41 = OpCopyObject %ulong %42 - OpStore %33 %41 - %43 = OpLoad %ulong %33 - OpReturnValue %43 + %1 = OpFunction %void None %44 + %27 = OpFunctionParameter %_ptr_Function_ulong + %28 = OpFunctionParameter %_ptr_Function_ulong + %35 = OpLabel + %29 = OpVariable %_ptr_Function_ulong Function + %30 = OpLoad %ulong %28 + OpStore %29 %30 + %32 = OpLoad %ulong %29 + %31 = OpIAdd %ulong %32 %ulong_1 + OpStore %29 %31 + %33 = OpLoad %ulong %29 + OpStore %27 %33 + OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/constant_f32.spvtxt b/ptx/src/test/spirv_run/constant_f32.spvtxt index 27c5f4e..46193a2 100644 --- a/ptx/src/test/spirv_run/constant_f32.spvtxt +++ b/ptx/src/test/spirv_run/constant_f32.spvtxt @@ -1,57 +1,48 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 32 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%24 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "constant_f32" -; OpDecorate %1 FunctionDenormModeINTEL 32 Preserve -%25 = OpTypeVoid -%26 = OpTypeInt 64 0 -%27 = OpTypeFunction %25 %26 %26 -%28 = OpTypePointer Function %26 -%29 = OpTypeFloat 32 -%30 = OpTypePointer Function %29 -%31 = OpTypePointer Generic %29 -%19 = OpConstant %29 0.5 -%1 = OpFunction %25 None %27 -%7 = OpFunctionParameter %26 -%8 = OpFunctionParameter %26 -%22 = OpLabel -%2 = OpVariable %28 Function -%3 = OpVariable %28 Function -%4 = OpVariable %28 Function -%5 = OpVariable %28 Function -%6 = OpVariable %30 Function -OpStore %2 %7 -OpStore %3 %8 -%10 = OpLoad %26 %2 -%9 = OpCopyObject %26 %10 -OpStore %4 %9 -%12 = OpLoad %26 %3 -%11 = OpCopyObject %26 %12 -OpStore %5 %11 -%14 = OpLoad %26 %4 -%20 = OpConvertUToPtr %31 %14 -%13 = OpLoad %29 %20 -OpStore %6 %13 -%16 = OpLoad %29 %6 -%15 = OpFMul %29 %16 %19 -OpStore %6 %15 -%17 = OpLoad %26 %5 -%18 = OpLoad %29 %6 -%21 = OpConvertUToPtr %31 %17 -OpStore %21 %18 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %22 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "constant_f32" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %25 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %float_0_5 = OpConstant %float 0.5 + %1 = OpFunction %void None %25 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %20 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %18 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpFMul %float %14 %float_0_5 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %19 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %19 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/constant_negative.spvtxt b/ptx/src/test/spirv_run/constant_negative.spvtxt index ec2ff72..5532e6e 100644 --- a/ptx/src/test/spirv_run/constant_negative.spvtxt +++ b/ptx/src/test/spirv_run/constant_negative.spvtxt @@ -1,56 +1,48 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 32 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%24 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "constant_negative" -%25 = OpTypeVoid -%26 = OpTypeInt 64 0 -%27 = OpTypeFunction %25 %26 %26 -%28 = OpTypePointer Function %26 -%29 = OpTypeInt 32 0 -%30 = OpTypePointer Function %29 -%31 = OpTypePointer Generic %29 -%19 = OpConstant %29 4294967295 -%1 = OpFunction %25 None %27 -%7 = OpFunctionParameter %26 -%8 = OpFunctionParameter %26 -%22 = OpLabel -%2 = OpVariable %28 Function -%3 = OpVariable %28 Function -%4 = OpVariable %28 Function -%5 = OpVariable %28 Function -%6 = OpVariable %30 Function -OpStore %2 %7 -OpStore %3 %8 -%10 = OpLoad %26 %2 -%9 = OpCopyObject %26 %10 -OpStore %4 %9 -%12 = OpLoad %26 %3 -%11 = OpCopyObject %26 %12 -OpStore %5 %11 -%14 = OpLoad %26 %4 -%20 = OpConvertUToPtr %31 %14 -%13 = OpLoad %29 %20 -OpStore %6 %13 -%16 = OpLoad %29 %6 -%15 = OpIMul %29 %16 %19 -OpStore %6 %15 -%17 = OpLoad %26 %5 -%18 = OpLoad %29 %6 -%21 = OpConvertUToPtr %31 %17 -OpStore %21 %18 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %22 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "constant_negative" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %25 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint +%uint_4294967295 = OpConstant %uint 4294967295 + %1 = OpFunction %void None %25 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %20 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_uint Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %18 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpIMul %uint %14 %uint_4294967295 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %uint %6 + %19 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %19 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/cos.ptx b/ptx/src/test/spirv_run/cos.ptx new file mode 100644 index 0000000..7983f20 --- /dev/null +++ b/ptx/src/test/spirv_run/cos.ptx @@ -0,0 +1,21 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.visible .entry cos( + .param .u64 input, + .param .u64 output +) +{ + .reg .u64 in_addr; + .reg .u64 out_addr; + .reg .f32 temp; + + ld.param.u64 in_addr, [input]; + ld.param.u64 out_addr, [output]; + + ld.f32 temp, [in_addr]; + cos.approx.f32 temp, temp; + st.f32 [out_addr], temp; + ret; +} diff --git a/ptx/src/test/spirv_run/cos.spvtxt b/ptx/src/test/spirv_run/cos.spvtxt new file mode 100644 index 0000000..6820142 --- /dev/null +++ b/ptx/src/test/spirv_run/cos.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "cos" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 cos %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/cvt_sat_s_u.spvtxt b/ptx/src/test/spirv_run/cvt_sat_s_u.spvtxt index 208c279..be321eb 100644 --- a/ptx/src/test/spirv_run/cvt_sat_s_u.spvtxt +++ b/ptx/src/test/spirv_run/cvt_sat_s_u.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %27 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %25 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "cvt_sat_s_u" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %30 = OpTypeFunction %void %ulong %ulong + %28 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint - %1 = OpFunction %void None %30 + %1 = OpFunction %void None %28 %9 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong - %25 = OpLabel + %23 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -27,25 +30,23 @@ %8 = OpVariable %_ptr_Function_uint Function OpStore %2 %9 OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %2 OpStore %4 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %5 %13 - %16 = OpLoad %ulong %4 - %23 = OpConvertUToPtr %_ptr_Generic_uint %16 - %15 = OpLoad %uint %23 - OpStore %6 %15 - %18 = OpLoad %uint %6 - %17 = OpSatConvertSToU %uint %18 - OpStore %7 %17 - %20 = OpLoad %uint %7 - %19 = OpBitcast %uint %20 - OpStore %8 %19 - %21 = OpLoad %ulong %5 - %22 = OpLoad %uint %8 - %24 = OpConvertUToPtr %_ptr_Generic_uint %21 - OpStore %24 %22 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %21 = OpConvertUToPtr %_ptr_Generic_uint %14 + %13 = OpLoad %uint %21 + OpStore %6 %13 + %16 = OpLoad %uint %6 + %15 = OpSatConvertSToU %uint %16 + OpStore %7 %15 + %18 = OpLoad %uint %7 + %17 = OpBitcast %uint %18 + OpStore %8 %17 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %8 + %22 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %22 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/cvta.spvtxt b/ptx/src/test/spirv_run/cvta.spvtxt index 84e7eac..cf6ff8b 100644 --- a/ptx/src/test/spirv_run/cvta.spvtxt +++ b/ptx/src/test/spirv_run/cvta.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %29 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %27 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "cvta" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %32 = OpTypeFunction %void %ulong %ulong + %30 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_Function_float = OpTypePointer Function %float %_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float - %1 = OpFunction %void None %32 + %1 = OpFunction %void None %30 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %27 = OpLabel + %25 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -25,29 +28,27 @@ %6 = OpVariable %_ptr_Function_float Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %20 = OpCopyObject %ulong %12 + %19 = OpCopyObject %ulong %20 + %11 = OpCopyObject %ulong %19 + OpStore %4 %11 + %14 = OpLoad %ulong %5 %22 = OpCopyObject %ulong %14 %21 = OpCopyObject %ulong %22 %13 = OpCopyObject %ulong %21 - OpStore %4 %13 - %16 = OpLoad %ulong %5 - %24 = OpCopyObject %ulong %16 - %23 = OpCopyObject %ulong %24 - %15 = OpCopyObject %ulong %23 - OpStore %5 %15 - %18 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %18 - %17 = OpLoad %float %25 - OpStore %6 %17 - %19 = OpLoad %ulong %5 - %20 = OpLoad %float %6 - %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %19 - OpStore %26 %20 + OpStore %5 %13 + %16 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %16 + %15 = OpLoad %float %23 + OpStore %6 %15 + %17 = OpLoad %ulong %5 + %18 = OpLoad %float %6 + %24 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %17 + OpStore %24 %18 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/div_approx.spvtxt b/ptx/src/test/spirv_run/div_approx.spvtxt index 40cc152..c62888c 100644 --- a/ptx/src/test/spirv_run/div_approx.spvtxt +++ b/ptx/src/test/spirv_run/div_approx.spvtxt @@ -1,65 +1,56 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 38 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%30 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "div_approx" -OpDecorate %1 FunctionDenormModeINTEL 32 Preserve -OpDecorate %18 FPFastMathMode AllowRecip -%31 = OpTypeVoid -%32 = OpTypeInt 64 0 -%33 = OpTypeFunction %31 %32 %32 -%34 = OpTypePointer Function %32 -%35 = OpTypeFloat 32 -%36 = OpTypePointer Function %35 -%37 = OpTypePointer Generic %35 -%23 = OpConstant %32 4 -%1 = OpFunction %31 None %33 -%8 = OpFunctionParameter %32 -%9 = OpFunctionParameter %32 -%28 = OpLabel -%2 = OpVariable %34 Function -%3 = OpVariable %34 Function -%4 = OpVariable %34 Function -%5 = OpVariable %34 Function -%6 = OpVariable %36 Function -%7 = OpVariable %36 Function -OpStore %2 %8 -OpStore %3 %9 -%11 = OpLoad %32 %2 -%10 = OpCopyObject %32 %11 -OpStore %4 %10 -%13 = OpLoad %32 %3 -%12 = OpCopyObject %32 %13 -OpStore %5 %12 -%15 = OpLoad %32 %4 -%25 = OpConvertUToPtr %37 %15 -%14 = OpLoad %35 %25 -OpStore %6 %14 -%17 = OpLoad %32 %4 -%24 = OpIAdd %32 %17 %23 -%26 = OpConvertUToPtr %37 %24 -%16 = OpLoad %35 %26 -OpStore %7 %16 -%19 = OpLoad %35 %6 -%20 = OpLoad %35 %7 -%18 = OpFDiv %35 %19 %20 -OpStore %6 %18 -%21 = OpLoad %32 %5 -%22 = OpLoad %35 %6 -%27 = OpConvertUToPtr %37 %21 -OpStore %27 %22 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %28 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "div_approx" + OpDecorate %16 FPFastMathMode AllowRecip + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %31 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %ulong_4 = OpConstant %ulong 4 + %1 = OpFunction %void None %31 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %26 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_float Function + %7 = OpVariable %_ptr_Function_float Function + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %4 %10 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_float %13 + %12 = OpLoad %float %23 + OpStore %6 %12 + %15 = OpLoad %ulong %4 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_float %22 + %14 = OpLoad %float %24 + OpStore %7 %14 + %17 = OpLoad %float %6 + %18 = OpLoad %float %7 + %16 = OpFDiv %float %17 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %float %6 + %25 = OpConvertUToPtr %_ptr_Generic_float %19 + OpStore %25 %20 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ex2.ptx b/ptx/src/test/spirv_run/ex2.ptx new file mode 100644 index 0000000..1edbcc6 --- /dev/null +++ b/ptx/src/test/spirv_run/ex2.ptx @@ -0,0 +1,21 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.visible .entry ex2( + .param .u64 input, + .param .u64 output +) +{ + .reg .u64 in_addr; + .reg .u64 out_addr; + .reg .f32 temp; + + ld.param.u64 in_addr, [input]; + ld.param.u64 out_addr, [output]; + + ld.f32 temp, [in_addr]; + ex2.approx.f32 temp, temp; + st.f32 [out_addr], temp; + ret; +} diff --git a/ptx/src/test/spirv_run/ex2.spvtxt b/ptx/src/test/spirv_run/ex2.spvtxt new file mode 100644 index 0000000..e9be3e3 --- /dev/null +++ b/ptx/src/test/spirv_run/ex2.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "ex2" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 exp2 %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/extern_shared.spvtxt b/ptx/src/test/spirv_run/extern_shared.spvtxt index b184980..fca8ee7 100644 --- a/ptx/src/test/spirv_run/extern_shared.spvtxt +++ b/ptx/src/test/spirv_run/extern_shared.spvtxt @@ -7,7 +7,7 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %32 = OpExtInstImport "OpenCL.std" + %30 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "extern_shared" %1 %void = OpTypeVoid @@ -18,51 +18,49 @@ %ulong = OpTypeInt 64 0 %uchar = OpTypeInt 8 0 %_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar - %40 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar + %38 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar %_ptr_Function__ptr_Workgroup_uchar = OpTypePointer Function %_ptr_Workgroup_uchar %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_ptr_Function__ptr_Workgroup_uint = OpTypePointer Function %_ptr_Workgroup_uint %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong - %2 = OpFunction %void None %40 + %2 = OpFunction %void None %38 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %28 = OpFunctionParameter %_ptr_Workgroup_uchar - %41 = OpLabel - %29 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function + %26 = OpFunctionParameter %_ptr_Workgroup_uchar + %39 = OpLabel + %27 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function - OpStore %29 %28 - OpBranch %26 - %26 = OpLabel + OpStore %27 %26 + OpBranch %24 + %24 = OpLabel OpStore %3 %8 OpStore %4 %9 - %11 = OpLoad %ulong %3 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %3 OpStore %5 %10 - %13 = OpLoad %ulong %4 - %12 = OpCopyObject %ulong %13 - OpStore %6 %12 - %15 = OpLoad %ulong %5 - %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %15 - %14 = OpLoad %ulong %22 - OpStore %7 %14 - %30 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %29 - %16 = OpLoad %_ptr_Workgroup_uint %30 - %17 = OpLoad %ulong %7 - %23 = OpBitcast %_ptr_Workgroup_ulong %16 - OpStore %23 %17 - %31 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %29 - %19 = OpLoad %_ptr_Workgroup_uint %31 - %24 = OpBitcast %_ptr_Workgroup_ulong %19 - %18 = OpLoad %ulong %24 - OpStore %7 %18 - %20 = OpLoad %ulong %6 - %21 = OpLoad %ulong %7 - %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %20 - OpStore %25 %21 + %11 = OpLoad %ulong %4 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %20 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %13 + %12 = OpLoad %ulong %20 + OpStore %7 %12 + %28 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %27 + %14 = OpLoad %_ptr_Workgroup_uint %28 + %15 = OpLoad %ulong %7 + %21 = OpBitcast %_ptr_Workgroup_ulong %14 + OpStore %21 %15 + %29 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %27 + %17 = OpLoad %_ptr_Workgroup_uint %29 + %22 = OpBitcast %_ptr_Workgroup_ulong %17 + %16 = OpLoad %ulong %22 + OpStore %7 %16 + %18 = OpLoad %ulong %6 + %19 = OpLoad %ulong %7 + %23 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %18 + OpStore %23 %19 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/fma.spvtxt b/ptx/src/test/spirv_run/fma.spvtxt index 4a90d09..9716198 100644 --- a/ptx/src/test/spirv_run/fma.spvtxt +++ b/ptx/src/test/spirv_run/fma.spvtxt @@ -1,72 +1,63 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 45 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%37 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "fma" -; OpDecorate %1 FunctionDenormModeINTEL 32 Preserve -%38 = OpTypeVoid -%39 = OpTypeInt 64 0 -%40 = OpTypeFunction %38 %39 %39 -%41 = OpTypePointer Function %39 -%42 = OpTypeFloat 32 -%43 = OpTypePointer Function %42 -%44 = OpTypePointer Generic %42 -%27 = OpConstant %39 4 -%29 = OpConstant %39 8 -%1 = OpFunction %38 None %40 -%9 = OpFunctionParameter %39 -%10 = OpFunctionParameter %39 -%35 = OpLabel -%2 = OpVariable %41 Function -%3 = OpVariable %41 Function -%4 = OpVariable %41 Function -%5 = OpVariable %41 Function -%6 = OpVariable %43 Function -%7 = OpVariable %43 Function -%8 = OpVariable %43 Function -OpStore %2 %9 -OpStore %3 %10 -%12 = OpLoad %39 %2 -%11 = OpCopyObject %39 %12 -OpStore %4 %11 -%14 = OpLoad %39 %3 -%13 = OpCopyObject %39 %14 -OpStore %5 %13 -%16 = OpLoad %39 %4 -%31 = OpConvertUToPtr %44 %16 -%15 = OpLoad %42 %31 -OpStore %6 %15 -%18 = OpLoad %39 %4 -%28 = OpIAdd %39 %18 %27 -%32 = OpConvertUToPtr %44 %28 -%17 = OpLoad %42 %32 -OpStore %7 %17 -%20 = OpLoad %39 %4 -%30 = OpIAdd %39 %20 %29 -%33 = OpConvertUToPtr %44 %30 -%19 = OpLoad %42 %33 -OpStore %8 %19 -%22 = OpLoad %42 %6 -%23 = OpLoad %42 %7 -%24 = OpLoad %42 %8 -%21 = OpExtInst %42 %37 mad %22 %23 %24 -OpStore %6 %21 -%25 = OpLoad %39 %5 -%26 = OpLoad %42 %6 -%34 = OpConvertUToPtr %44 %25 -OpStore %34 %26 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %35 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "fma" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %38 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %ulong_4 = OpConstant %ulong 4 + %ulong_8 = OpConstant %ulong 8 + %1 = OpFunction %void None %38 + %9 = OpFunctionParameter %ulong + %10 = OpFunctionParameter %ulong + %33 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_float Function + %7 = OpVariable %_ptr_Function_float Function + %8 = OpVariable %_ptr_Function_float Function + OpStore %2 %9 + OpStore %3 %10 + %11 = OpLoad %ulong %2 + OpStore %4 %11 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %29 = OpConvertUToPtr %_ptr_Generic_float %14 + %13 = OpLoad %float %29 + OpStore %6 %13 + %16 = OpLoad %ulong %4 + %26 = OpIAdd %ulong %16 %ulong_4 + %30 = OpConvertUToPtr %_ptr_Generic_float %26 + %15 = OpLoad %float %30 + OpStore %7 %15 + %18 = OpLoad %ulong %4 + %28 = OpIAdd %ulong %18 %ulong_8 + %31 = OpConvertUToPtr %_ptr_Generic_float %28 + %17 = OpLoad %float %31 + OpStore %8 %17 + %20 = OpLoad %float %6 + %21 = OpLoad %float %7 + %22 = OpLoad %float %8 + %19 = OpExtInst %float %35 mad %20 %21 %22 + OpStore %6 %19 + %23 = OpLoad %ulong %5 + %24 = OpLoad %float %6 + %32 = OpConvertUToPtr %_ptr_Generic_float %23 + OpStore %32 %24 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/global_array.spvtxt b/ptx/src/test/spirv_run/global_array.spvtxt index a4ed91d..25874ac 100644 --- a/ptx/src/test/spirv_run/global_array.spvtxt +++ b/ptx/src/test/spirv_run/global_array.spvtxt @@ -7,28 +7,28 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %22 = OpExtInstImport "OpenCL.std" + %21 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "global_array" %1 %void = OpTypeVoid %uint = OpTypeInt 32 0 %uint_4 = OpConstant %uint 4 %_arr_uint_uint_4 = OpTypeArray %uint %uint_4 -%_ptr_CrossWorkgroup__arr_uint_uint_4 = OpTypePointer CrossWorkgroup %_arr_uint_uint_4 - %uint_4_0 = OpConstant %uint 4 %uint_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 - %31 = OpConstantComposite %_arr_uint_uint_4 %uint_1 %uint_0 %uint_0 %uint_0 - %1 = OpVariable %_ptr_CrossWorkgroup__arr_uint_uint_4 CrossWorkgroup %31 + %28 = OpConstantComposite %_arr_uint_uint_4 %uint_1 %uint_0 %uint_0 %uint_0 + %uint_4_0 = OpConstant %uint 4 +%_ptr_CrossWorkgroup__arr_uint_uint_4 = OpTypePointer CrossWorkgroup %_arr_uint_uint_4 + %1 = OpVariable %_ptr_CrossWorkgroup__arr_uint_uint_4 CrossWorkgroup %28 %ulong = OpTypeInt 64 0 - %33 = OpTypeFunction %void %ulong %ulong + %32 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint - %2 = OpFunction %void None %33 + %2 = OpFunction %void None %32 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %20 = OpLabel + %19 = OpLabel %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function @@ -36,19 +36,18 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %3 %8 OpStore %4 %9 - %17 = OpConvertPtrToU %ulong %1 - %10 = OpCopyObject %ulong %17 + %16 = OpConvertPtrToU %ulong %1 + %10 = OpCopyObject %ulong %16 OpStore %5 %10 - %12 = OpLoad %ulong %4 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %4 OpStore %6 %11 - %14 = OpLoad %ulong %5 + %13 = OpLoad %ulong %5 + %17 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %13 + %12 = OpLoad %uint %17 + OpStore %7 %12 + %14 = OpLoad %ulong %6 + %15 = OpLoad %uint %7 %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %14 - %13 = OpLoad %uint %18 - OpStore %7 %13 - %15 = OpLoad %ulong %6 - %16 = OpLoad %uint %7 - %19 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %15 - OpStore %19 %16 + OpStore %18 %15 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/implicit_param.spvtxt b/ptx/src/test/spirv_run/implicit_param.spvtxt index c30788c..a78e53f 100644 --- a/ptx/src/test/spirv_run/implicit_param.spvtxt +++ b/ptx/src/test/spirv_run/implicit_param.spvtxt @@ -2,25 +2,27 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %28 = OpExtInstImport "OpenCL.std" + %24 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "implicit_param" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %31 = OpTypeFunction %void %ulong %ulong + %27 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_Function_float = OpTypePointer Function %float %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float - %1 = OpFunction %void None %31 + %1 = OpFunction %void None %27 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %26 = OpLabel + %22 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -29,27 +31,23 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %15 - %14 = OpLoad %float %22 - OpStore %6 %14 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %13 + %12 = OpLoad %float %18 + OpStore %6 %12 + %14 = OpLoad %float %6 + %19 = OpBitcast %_ptr_Function_float %7 + OpStore %19 %14 + %20 = OpBitcast %_ptr_Function_float %7 + %15 = OpLoad %float %20 + OpStore %6 %15 + %16 = OpLoad %ulong %5 %17 = OpLoad %float %6 - %23 = OpCopyObject %float %17 - %16 = OpBitcast %uint %23 - OpStore %7 %16 - %19 = OpLoad %uint %7 - %24 = OpBitcast %float %19 - %18 = OpCopyObject %float %24 - OpStore %6 %18 - %20 = OpLoad %ulong %5 - %21 = OpLoad %float %6 - %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %20 - OpStore %25 %21 + %21 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %16 + OpStore %21 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ld_st.spvtxt b/ptx/src/test/spirv_run/ld_st.spvtxt index d36db57..8e3f98d 100644 --- a/ptx/src/test/spirv_run/ld_st.spvtxt +++ b/ptx/src/test/spirv_run/ld_st.spvtxt @@ -2,20 +2,23 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %21 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %19 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "ld_st" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %24 = OpTypeFunction %void %ulong %ulong + %22 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %24 + %1 = OpFunction %void None %22 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %19 = OpLabel + %17 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -23,19 +26,17 @@ %6 = OpVariable %_ptr_Function_ulong Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 - %17 = OpConvertUToPtr %_ptr_Generic_ulong %14 - %13 = OpLoad %ulong %17 - OpStore %6 %13 - %15 = OpLoad %ulong %5 - %16 = OpLoad %ulong %6 - %18 = OpConvertUToPtr %_ptr_Generic_ulong %15 - OpStore %18 %16 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %15 = OpConvertUToPtr %_ptr_Generic_ulong %12 + %11 = OpLoad %ulong %15 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %14 = OpLoad %ulong %6 + %16 = OpConvertUToPtr %_ptr_Generic_ulong %13 + OpStore %16 %14 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ld_st_implicit.spvtxt b/ptx/src/test/spirv_run/ld_st_implicit.spvtxt index d4d9499..35f715b 100644 --- a/ptx/src/test/spirv_run/ld_st_implicit.spvtxt +++ b/ptx/src/test/spirv_run/ld_st_implicit.spvtxt @@ -7,20 +7,20 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %23 = OpExtInstImport "OpenCL.std" + %21 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "ld_st_implicit" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %26 = OpTypeFunction %void %ulong %ulong + %24 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float %uint = OpTypeInt 32 0 - %1 = OpFunction %void None %26 + %1 = OpFunction %void None %24 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %21 = OpLabel + %19 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -28,24 +28,22 @@ %6 = OpVariable %_ptr_Function_ulong Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 - %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %14 - %17 = OpLoad %float %18 - %31 = OpBitcast %uint %17 - %13 = OpUConvert %ulong %31 - OpStore %6 %13 - %15 = OpLoad %ulong %5 - %16 = OpLoad %ulong %6 - %19 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %15 - %32 = OpBitcast %ulong %16 - %33 = OpUConvert %uint %32 - %20 = OpBitcast %float %33 - OpStore %19 %20 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %16 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %12 + %15 = OpLoad %float %16 + %29 = OpBitcast %uint %15 + %11 = OpUConvert %ulong %29 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %14 = OpLoad %ulong %6 + %17 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %13 + %30 = OpBitcast %ulong %14 + %31 = OpUConvert %uint %30 + %18 = OpBitcast %float %31 + OpStore %17 %18 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ld_st_offset.spvtxt b/ptx/src/test/spirv_run/ld_st_offset.spvtxt index 208b53b..963d88a 100644 --- a/ptx/src/test/spirv_run/ld_st_offset.spvtxt +++ b/ptx/src/test/spirv_run/ld_st_offset.spvtxt @@ -2,24 +2,27 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %32 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %30 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "ld_st_offset" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %35 = OpTypeFunction %void %ulong %ulong + %33 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint %ulong_4 = OpConstant %ulong 4 %ulong_4_0 = OpConstant %ulong 4 - %1 = OpFunction %void None %35 + %1 = OpFunction %void None %33 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %30 = OpLabel + %28 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -28,29 +31,27 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %24 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %24 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %26 = OpConvertUToPtr %_ptr_Generic_uint %15 - %14 = OpLoad %uint %26 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %23 = OpIAdd %ulong %17 %ulong_4 - %27 = OpConvertUToPtr %_ptr_Generic_uint %23 - %16 = OpLoad %uint %27 - OpStore %7 %16 + %21 = OpIAdd %ulong %15 %ulong_4 + %25 = OpConvertUToPtr %_ptr_Generic_uint %21 + %14 = OpLoad %uint %25 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %uint %7 + %26 = OpConvertUToPtr %_ptr_Generic_uint %16 + OpStore %26 %17 %18 = OpLoad %ulong %5 - %19 = OpLoad %uint %7 - %28 = OpConvertUToPtr %_ptr_Generic_uint %18 - OpStore %28 %19 - %20 = OpLoad %ulong %5 - %21 = OpLoad %uint %6 - %25 = OpIAdd %ulong %20 %ulong_4_0 - %29 = OpConvertUToPtr %_ptr_Generic_uint %25 - OpStore %29 %21 + %19 = OpLoad %uint %6 + %23 = OpIAdd %ulong %18 %ulong_4_0 + %27 = OpConvertUToPtr %_ptr_Generic_uint %23 + OpStore %27 %19 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/lg2.ptx b/ptx/src/test/spirv_run/lg2.ptx new file mode 100644 index 0000000..c571955 --- /dev/null +++ b/ptx/src/test/spirv_run/lg2.ptx @@ -0,0 +1,21 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.visible .entry lg2( + .param .u64 input, + .param .u64 output +) +{ + .reg .u64 in_addr; + .reg .u64 out_addr; + .reg .f32 temp; + + ld.param.u64 in_addr, [input]; + ld.param.u64 out_addr, [output]; + + ld.f32 temp, [in_addr]; + lg2.approx.f32 temp, temp; + st.f32 [out_addr], temp; + ret; +} diff --git a/ptx/src/test/spirv_run/lg2.spvtxt b/ptx/src/test/spirv_run/lg2.spvtxt new file mode 100644 index 0000000..d30fe8a --- /dev/null +++ b/ptx/src/test/spirv_run/lg2.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "lg2" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 log2 %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/local_align.spvtxt b/ptx/src/test/spirv_run/local_align.spvtxt index 2482a75..915ac6f 100644 --- a/ptx/src/test/spirv_run/local_align.spvtxt +++ b/ptx/src/test/spirv_run/local_align.spvtxt @@ -2,26 +2,29 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %22 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %20 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "local_align" OpDecorate %4 Alignment 8 %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %25 = OpTypeFunction %void %ulong %ulong + %23 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong - %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 %uint_8 = OpConstant %uint 8 %_arr_uchar_uint_8 = OpTypeArray %uchar %uint_8 %_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8 %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %25 + %1 = OpFunction %void None %23 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %20 = OpLabel + %18 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function__arr_uchar_uint_8 Function @@ -30,19 +33,17 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %5 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %6 %12 - %15 = OpLoad %ulong %5 - %18 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %18 - OpStore %7 %14 - %16 = OpLoad %ulong %6 - %17 = OpLoad %ulong %7 - %19 = OpConvertUToPtr %_ptr_Generic_ulong %16 - OpStore %19 %17 + %11 = OpLoad %ulong %3 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %16 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %16 + OpStore %7 %12 + %14 = OpLoad %ulong %6 + %15 = OpLoad %ulong %7 + %17 = OpConvertUToPtr %_ptr_Generic_ulong %14 + OpStore %17 %15 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mad_s32.spvtxt b/ptx/src/test/spirv_run/mad_s32.spvtxt index 3a7153d..9150089 100644 --- a/ptx/src/test/spirv_run/mad_s32.spvtxt +++ b/ptx/src/test/spirv_run/mad_s32.spvtxt @@ -2,15 +2,17 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %48 = OpExtInstImport "OpenCL.std" + %46 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mad_s32" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %51 = OpTypeFunction %void %ulong %ulong + %49 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint @@ -19,10 +21,10 @@ %ulong_8 = OpConstant %ulong 8 %ulong_4_0 = OpConstant %ulong 4 %ulong_8_0 = OpConstant %ulong 8 - %1 = OpFunction %void None %51 + %1 = OpFunction %void None %49 %10 = OpFunctionParameter %ulong %11 = OpFunctionParameter %ulong - %46 = OpLabel + %44 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -33,45 +35,43 @@ %9 = OpVariable %_ptr_Function_uint Function OpStore %2 %10 OpStore %3 %11 - %13 = OpLoad %ulong %2 - %12 = OpCopyObject %ulong %13 + %12 = OpLoad %ulong %2 OpStore %4 %12 - %15 = OpLoad %ulong %3 - %14 = OpCopyObject %ulong %15 - OpStore %5 %14 + %13 = OpLoad %ulong %3 + OpStore %5 %13 + %15 = OpLoad %ulong %4 + %38 = OpConvertUToPtr %_ptr_Generic_uint %15 + %14 = OpLoad %uint %38 + OpStore %7 %14 %17 = OpLoad %ulong %4 - %40 = OpConvertUToPtr %_ptr_Generic_uint %17 - %16 = OpLoad %uint %40 - OpStore %7 %16 + %31 = OpIAdd %ulong %17 %ulong_4 + %39 = OpConvertUToPtr %_ptr_Generic_uint %31 + %16 = OpLoad %uint %39 + OpStore %8 %16 %19 = OpLoad %ulong %4 - %33 = OpIAdd %ulong %19 %ulong_4 - %41 = OpConvertUToPtr %_ptr_Generic_uint %33 - %18 = OpLoad %uint %41 - OpStore %8 %18 - %21 = OpLoad %ulong %4 - %35 = OpIAdd %ulong %21 %ulong_8 - %42 = OpConvertUToPtr %_ptr_Generic_uint %35 - %20 = OpLoad %uint %42 - OpStore %9 %20 - %23 = OpLoad %uint %7 - %24 = OpLoad %uint %8 - %25 = OpLoad %uint %9 - %56 = OpIMul %uint %23 %24 - %22 = OpIAdd %uint %25 %56 - OpStore %6 %22 + %33 = OpIAdd %ulong %19 %ulong_8 + %40 = OpConvertUToPtr %_ptr_Generic_uint %33 + %18 = OpLoad %uint %40 + OpStore %9 %18 + %21 = OpLoad %uint %7 + %22 = OpLoad %uint %8 + %23 = OpLoad %uint %9 + %54 = OpIMul %uint %21 %22 + %20 = OpIAdd %uint %23 %54 + OpStore %6 %20 + %24 = OpLoad %ulong %5 + %25 = OpLoad %uint %6 + %41 = OpConvertUToPtr %_ptr_Generic_uint %24 + OpStore %41 %25 %26 = OpLoad %ulong %5 %27 = OpLoad %uint %6 - %43 = OpConvertUToPtr %_ptr_Generic_uint %26 - OpStore %43 %27 + %35 = OpIAdd %ulong %26 %ulong_4_0 + %42 = OpConvertUToPtr %_ptr_Generic_uint %35 + OpStore %42 %27 %28 = OpLoad %ulong %5 %29 = OpLoad %uint %6 - %37 = OpIAdd %ulong %28 %ulong_4_0 - %44 = OpConvertUToPtr %_ptr_Generic_uint %37 - OpStore %44 %29 - %30 = OpLoad %ulong %5 - %31 = OpLoad %uint %6 - %39 = OpIAdd %ulong %30 %ulong_8_0 - %45 = OpConvertUToPtr %_ptr_Generic_uint %39 - OpStore %45 %31 + %37 = OpIAdd %ulong %28 %ulong_8_0 + %43 = OpConvertUToPtr %_ptr_Generic_uint %37 + OpStore %43 %29 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/max.spvtxt b/ptx/src/test/spirv_run/max.spvtxt index cab9a9a..05eb705 100644 --- a/ptx/src/test/spirv_run/max.spvtxt +++ b/ptx/src/test/spirv_run/max.spvtxt @@ -7,21 +7,21 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %30 = OpExtInstImport "OpenCL.std" + %28 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "max" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %33 = OpTypeFunction %void %ulong %ulong + %31 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint %ulong_4 = OpConstant %ulong 4 - %1 = OpFunction %void None %33 + %1 = OpFunction %void None %31 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %28 = OpLabel + %26 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -30,28 +30,26 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %23 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_uint %15 - %14 = OpLoad %uint %25 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %24 = OpIAdd %ulong %17 %ulong_4 - %26 = OpConvertUToPtr %_ptr_Generic_uint %24 - %16 = OpLoad %uint %26 - OpStore %7 %16 - %19 = OpLoad %uint %6 - %20 = OpLoad %uint %7 - %18 = OpExtInst %uint %30 s_max %19 %20 - OpStore %6 %18 - %21 = OpLoad %ulong %5 - %22 = OpLoad %uint %6 - %27 = OpConvertUToPtr %_ptr_Generic_uint %21 - OpStore %27 %22 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_uint %22 + %14 = OpLoad %uint %24 + OpStore %7 %14 + %17 = OpLoad %uint %6 + %18 = OpLoad %uint %7 + %16 = OpExtInst %uint %28 s_max %17 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %6 + %25 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %25 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/min.spvtxt b/ptx/src/test/spirv_run/min.spvtxt index 119cd15..d0d2b9a 100644 --- a/ptx/src/test/spirv_run/min.spvtxt +++ b/ptx/src/test/spirv_run/min.spvtxt @@ -7,21 +7,21 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %30 = OpExtInstImport "OpenCL.std" + %28 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "min" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %33 = OpTypeFunction %void %ulong %ulong + %31 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint %ulong_4 = OpConstant %ulong 4 - %1 = OpFunction %void None %33 + %1 = OpFunction %void None %31 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %28 = OpLabel + %26 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -30,28 +30,26 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %23 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_uint %15 - %14 = OpLoad %uint %25 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %24 = OpIAdd %ulong %17 %ulong_4 - %26 = OpConvertUToPtr %_ptr_Generic_uint %24 - %16 = OpLoad %uint %26 - OpStore %7 %16 - %19 = OpLoad %uint %6 - %20 = OpLoad %uint %7 - %18 = OpExtInst %uint %30 s_min %19 %20 - OpStore %6 %18 - %21 = OpLoad %ulong %5 - %22 = OpLoad %uint %6 - %27 = OpConvertUToPtr %_ptr_Generic_uint %21 - OpStore %27 %22 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_uint %22 + %14 = OpLoad %uint %24 + OpStore %7 %14 + %17 = OpLoad %uint %6 + %18 = OpLoad %uint %7 + %16 = OpExtInst %uint %28 s_min %17 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %6 + %25 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %25 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 7ba3c4d..3fa82ba 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -60,8 +60,7 @@ test_ptx!(call, [1u64], [2u64]); test_ptx!(vector, [1u32, 2u32], [3u32, 3u32]); test_ptx!(ld_st_offset, [1u32, 2u32], [2u32, 1u32]); test_ptx!(ntid, [3u32], [4u32]); -// TODO: enable test below -// test_ptx!(reg_local, [12u64], [13u64]); +test_ptx!(reg_local, [12u64], [13u64]); test_ptx!(mov_address, [0xDEADu64], [0u64]); test_ptx!(b64tof64, [111u64], [111u64]); test_ptx!(implicit_param, [34u32], [34u32]); @@ -105,6 +104,10 @@ test_ptx!(div_approx, [1f32, 2f32], [0.5f32]); test_ptx!(sqrt, [0.25f32], [0.5f32]); test_ptx!(rsqrt, [0.25f64], [2f64]); test_ptx!(neg, [181i32], [-181i32]); +test_ptx!(sin, [std::f32::consts::PI/2f32], [1f32]); +test_ptx!(cos, [std::f32::consts::PI], [-1f32]); +test_ptx!(lg2, [512f32], [9f32]); +test_ptx!(ex2, [10f32], [1024f32]); struct DisplayError { err: T, diff --git a/ptx/src/test/spirv_run/mov.spvtxt b/ptx/src/test/spirv_run/mov.spvtxt index d8a5029..15118aa 100644 --- a/ptx/src/test/spirv_run/mov.spvtxt +++ b/ptx/src/test/spirv_run/mov.spvtxt @@ -2,20 +2,23 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %24 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %22 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mov" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %27 = OpTypeFunction %void %ulong %ulong + %25 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %27 + %1 = OpFunction %void None %25 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %22 = OpLabel + %20 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -24,22 +27,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %20 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %20 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpCopyObject %ulong %17 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %21 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %18 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpCopyObject %ulong %15 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %19 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mul_hi.spvtxt b/ptx/src/test/spirv_run/mul_hi.spvtxt index bea23a9..8449183 100644 --- a/ptx/src/test/spirv_run/mul_hi.spvtxt +++ b/ptx/src/test/spirv_run/mul_hi.spvtxt @@ -2,21 +2,24 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %25 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %23 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mul_hi" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %28 = OpTypeFunction %void %ulong %ulong + %26 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %28 + %1 = OpFunction %void None %26 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %23 = OpLabel + %21 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -25,22 +28,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpExtInst %ulong %25 u_mul_hi %17 %ulong_2 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %22 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpExtInst %ulong %23 u_mul_hi %15 %ulong_2 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %20 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mul_lo.spvtxt b/ptx/src/test/spirv_run/mul_lo.spvtxt index e114374..d4b2566 100644 --- a/ptx/src/test/spirv_run/mul_lo.spvtxt +++ b/ptx/src/test/spirv_run/mul_lo.spvtxt @@ -2,21 +2,24 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %25 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %23 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mul_lo" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %28 = OpTypeFunction %void %ulong %ulong + %26 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %28 + %1 = OpFunction %void None %26 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %23 = OpLabel + %21 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -25,22 +28,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpIMul %ulong %17 %ulong_2 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %22 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpIMul %ulong %15 %ulong_2 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %20 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mul_non_ftz.spvtxt b/ptx/src/test/spirv_run/mul_non_ftz.spvtxt index 78153aa..cb20943 100644 --- a/ptx/src/test/spirv_run/mul_non_ftz.spvtxt +++ b/ptx/src/test/spirv_run/mul_non_ftz.spvtxt @@ -7,25 +7,21 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - OpCapability DenormFlushToZero - OpCapability DenormPreserve - OpExtension "SPV_KHR_float_controls" - %30 = OpExtInstImport "OpenCL.std" + %28 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mul_non_ftz" - OpExecutionMode %1 DenormPreserve 32 %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %33 = OpTypeFunction %void %ulong %ulong + %31 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_Function_float = OpTypePointer Function %float %_ptr_Generic_float = OpTypePointer Generic %float %ulong_4 = OpConstant %ulong 4 - %1 = OpFunction %void None %33 + %1 = OpFunction %void None %31 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %28 = OpLabel + %26 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -34,28 +30,26 @@ %7 = OpVariable %_ptr_Function_float Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_float %13 + %12 = OpLoad %float %23 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_float %15 - %14 = OpLoad %float %25 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %24 = OpIAdd %ulong %17 %ulong_4 - %26 = OpConvertUToPtr %_ptr_Generic_float %24 - %16 = OpLoad %float %26 - OpStore %7 %16 - %19 = OpLoad %float %6 - %20 = OpLoad %float %7 - %18 = OpFMul %float %19 %20 - OpStore %6 %18 - %21 = OpLoad %ulong %5 - %22 = OpLoad %float %6 - %27 = OpConvertUToPtr %_ptr_Generic_float %21 - OpStore %27 %22 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_float %22 + %14 = OpLoad %float %24 + OpStore %7 %14 + %17 = OpLoad %float %6 + %18 = OpLoad %float %7 + %16 = OpFMul %float %17 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %float %6 + %25 = OpConvertUToPtr %_ptr_Generic_float %19 + OpStore %25 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mul_wide.spvtxt b/ptx/src/test/spirv_run/mul_wide.spvtxt index 8ac0459..632fa3b 100644 --- a/ptx/src/test/spirv_run/mul_wide.spvtxt +++ b/ptx/src/test/spirv_run/mul_wide.spvtxt @@ -7,24 +7,24 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %32 = OpExtInstImport "OpenCL.std" + %30 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mul_wide" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %35 = OpTypeFunction %void %ulong %ulong + %33 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint %ulong_4 = OpConstant %ulong 4 - %_struct_40 = OpTypeStruct %uint %uint + %_struct_38 = OpTypeStruct %uint %uint %v2uint = OpTypeVector %uint 2 %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %35 + %1 = OpFunction %void None %33 %9 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong - %30 = OpLabel + %28 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -34,33 +34,31 @@ %8 = OpVariable %_ptr_Function_ulong Function OpStore %2 %9 OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %2 OpStore %4 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %5 %13 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %24 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %14 + %13 = OpLoad %uint %24 + OpStore %6 %13 %16 = OpLoad %ulong %4 - %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %16 - %15 = OpLoad %uint %26 - OpStore %6 %15 - %18 = OpLoad %ulong %4 - %25 = OpIAdd %ulong %18 %ulong_4 - %27 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %25 - %17 = OpLoad %uint %27 - OpStore %7 %17 - %20 = OpLoad %uint %6 - %21 = OpLoad %uint %7 - %41 = OpSMulExtended %_struct_40 %20 %21 - %42 = OpCompositeExtract %uint %41 0 - %43 = OpCompositeExtract %uint %41 1 - %45 = OpCompositeConstruct %v2uint %42 %43 - %19 = OpBitcast %ulong %45 - OpStore %8 %19 - %22 = OpLoad %ulong %5 - %23 = OpLoad %ulong %8 - %28 = OpConvertUToPtr %_ptr_Generic_ulong %22 - %29 = OpCopyObject %ulong %23 - OpStore %28 %29 + %23 = OpIAdd %ulong %16 %ulong_4 + %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %23 + %15 = OpLoad %uint %25 + OpStore %7 %15 + %18 = OpLoad %uint %6 + %19 = OpLoad %uint %7 + %39 = OpSMulExtended %_struct_38 %18 %19 + %40 = OpCompositeExtract %uint %39 0 + %41 = OpCompositeExtract %uint %39 1 + %43 = OpCompositeConstruct %v2uint %40 %41 + %17 = OpBitcast %ulong %43 + OpStore %8 %17 + %20 = OpLoad %ulong %5 + %21 = OpLoad %ulong %8 + %26 = OpConvertUToPtr %_ptr_Generic_ulong %20 + %27 = OpCopyObject %ulong %21 + OpStore %26 %27 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/neg.spvtxt b/ptx/src/test/spirv_run/neg.spvtxt index b358858..50726aa 100644 --- a/ptx/src/test/spirv_run/neg.spvtxt +++ b/ptx/src/test/spirv_run/neg.spvtxt @@ -2,46 +2,46 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %26 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL - OpEntryPoint Kernel %1 "not" + OpEntryPoint Kernel %1 "neg" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %29 = OpTypeFunction %void %ulong %ulong + %24 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong -%_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %29 + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %9 = OpFunctionParameter %ulong - %24 = OpLabel + %19 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function - %6 = OpVariable %_ptr_Function_ulong Function - %7 = OpVariable %_ptr_Function_ulong Function - OpStore %2 %8 - OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 - OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %20 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %20 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %22 = OpCopyObject %ulong %17 - %21 = OpNot %ulong %22 - %16 = OpCopyObject %ulong %21 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %23 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %23 %19 + %6 = OpVariable %_ptr_Function_uint Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %17 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpSNegate %uint %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %uint %6 + %18 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %18 %16 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/not.spvtxt b/ptx/src/test/spirv_run/not.spvtxt index b358858..d6bc389 100644 --- a/ptx/src/test/spirv_run/not.spvtxt +++ b/ptx/src/test/spirv_run/not.spvtxt @@ -2,20 +2,23 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %26 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %24 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "not" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %29 = OpTypeFunction %void %ulong %ulong + %27 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %29 + %1 = OpFunction %void None %27 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %24 = OpLabel + %22 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -24,24 +27,22 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %20 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %20 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %22 = OpCopyObject %ulong %17 - %21 = OpNot %ulong %22 - %16 = OpCopyObject %ulong %21 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %23 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %23 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %18 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %20 = OpCopyObject %ulong %15 + %19 = OpNot %ulong %20 + %14 = OpCopyObject %ulong %19 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %21 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %21 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ntid.spvtxt b/ptx/src/test/spirv_run/ntid.spvtxt index be16d2e..d1a3caa 100644 --- a/ptx/src/test/spirv_run/ntid.spvtxt +++ b/ptx/src/test/spirv_run/ntid.spvtxt @@ -2,10 +2,12 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %29 = OpExtInstImport "OpenCL.std" + %27 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "ntid" %gl_WorkGroupSize OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize @@ -15,14 +17,14 @@ %_ptr_UniformConstant_v4uint = OpTypePointer UniformConstant %v4uint %gl_WorkGroupSize = OpVariable %_ptr_UniformConstant_v4uint UniformConstant %ulong = OpTypeInt 64 0 - %35 = OpTypeFunction %void %ulong %ulong + %33 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint - %1 = OpFunction %void None %35 + %1 = OpFunction %void None %33 %9 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong - %27 = OpLabel + %25 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -31,27 +33,25 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %9 OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %2 OpStore %4 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %5 %13 - %16 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_uint %16 - %15 = OpLoad %uint %25 - OpStore %6 %15 - %18 = OpLoad %v4uint %gl_WorkGroupSize - %24 = OpCompositeExtract %uint %18 0 - %17 = OpCopyObject %uint %24 - OpStore %7 %17 - %20 = OpLoad %uint %6 - %21 = OpLoad %uint %7 - %19 = OpIAdd %uint %20 %21 - OpStore %6 %19 - %22 = OpLoad %ulong %5 - %23 = OpLoad %uint %6 - %26 = OpConvertUToPtr %_ptr_Generic_uint %22 - OpStore %26 %23 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_uint %14 + %13 = OpLoad %uint %23 + OpStore %6 %13 + %16 = OpLoad %v4uint %gl_WorkGroupSize + %22 = OpCompositeExtract %uint %16 0 + %15 = OpCopyObject %uint %22 + OpStore %7 %15 + %18 = OpLoad %uint %6 + %19 = OpLoad %uint %7 + %17 = OpIAdd %uint %18 %19 + OpStore %6 %17 + %20 = OpLoad %ulong %5 + %21 = OpLoad %uint %6 + %24 = OpConvertUToPtr %_ptr_Generic_uint %20 + OpStore %24 %21 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/or.spvtxt b/ptx/src/test/spirv_run/or.spvtxt index fbf80c5..312b1b3 100644 --- a/ptx/src/test/spirv_run/or.spvtxt +++ b/ptx/src/test/spirv_run/or.spvtxt @@ -7,19 +7,19 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %33 = OpExtInstImport "OpenCL.std" + %31 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "or" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %36 = OpTypeFunction %void %ulong %ulong + %34 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_8 = OpConstant %ulong 8 - %1 = OpFunction %void None %36 + %1 = OpFunction %void None %34 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %31 = OpLabel + %29 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -28,31 +28,29 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %23 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %25 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %24 = OpIAdd %ulong %17 %ulong_8 - %26 = OpConvertUToPtr %_ptr_Generic_ulong %24 - %16 = OpLoad %ulong %26 - OpStore %7 %16 - %19 = OpLoad %ulong %6 - %20 = OpLoad %ulong %7 - %28 = OpCopyObject %ulong %19 - %29 = OpCopyObject %ulong %20 - %27 = OpBitwiseOr %ulong %28 %29 - %18 = OpCopyObject %ulong %27 - OpStore %6 %18 - %21 = OpLoad %ulong %5 - %22 = OpLoad %ulong %6 - %30 = OpConvertUToPtr %_ptr_Generic_ulong %21 - OpStore %30 %22 + %22 = OpIAdd %ulong %15 %ulong_8 + %24 = OpConvertUToPtr %_ptr_Generic_ulong %22 + %14 = OpLoad %ulong %24 + OpStore %7 %14 + %17 = OpLoad %ulong %6 + %18 = OpLoad %ulong %7 + %26 = OpCopyObject %ulong %17 + %27 = OpCopyObject %ulong %18 + %25 = OpBitwiseOr %ulong %26 %27 + %16 = OpCopyObject %ulong %25 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %ulong %6 + %28 = OpConvertUToPtr %_ptr_Generic_ulong %19 + OpStore %28 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/pred_not.spvtxt b/ptx/src/test/spirv_run/pred_not.spvtxt index 410b1e4..178c98f 100644 --- a/ptx/src/test/spirv_run/pred_not.spvtxt +++ b/ptx/src/test/spirv_run/pred_not.spvtxt @@ -2,15 +2,17 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %44 = OpExtInstImport "OpenCL.std" + %42 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "pred_not" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %47 = OpTypeFunction %void %ulong %ulong + %45 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %bool = OpTypeBool %_ptr_Function_bool = OpTypePointer Function %bool @@ -20,10 +22,10 @@ %false = OpConstantFalse %bool %ulong_1 = OpConstant %ulong 1 %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %47 + %1 = OpFunction %void None %45 %14 = OpFunctionParameter %ulong %15 = OpFunctionParameter %ulong - %42 = OpLabel + %40 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -34,45 +36,43 @@ %9 = OpVariable %_ptr_Function_bool Function OpStore %2 %14 OpStore %3 %15 - %17 = OpLoad %ulong %2 - %16 = OpCopyObject %ulong %17 + %16 = OpLoad %ulong %2 OpStore %4 %16 - %19 = OpLoad %ulong %3 - %18 = OpCopyObject %ulong %19 - OpStore %5 %18 + %17 = OpLoad %ulong %3 + OpStore %5 %17 + %19 = OpLoad %ulong %4 + %37 = OpConvertUToPtr %_ptr_Generic_ulong %19 + %18 = OpLoad %ulong %37 + OpStore %6 %18 %21 = OpLoad %ulong %4 - %39 = OpConvertUToPtr %_ptr_Generic_ulong %21 - %20 = OpLoad %ulong %39 - OpStore %6 %20 - %23 = OpLoad %ulong %4 - %36 = OpIAdd %ulong %23 %ulong_8 - %40 = OpConvertUToPtr %_ptr_Generic_ulong %36 - %22 = OpLoad %ulong %40 - OpStore %7 %22 - %25 = OpLoad %ulong %6 - %26 = OpLoad %ulong %7 - %24 = OpULessThan %bool %25 %26 - OpStore %9 %24 - %28 = OpLoad %bool %9 - %27 = OpSelect %bool %28 %false %true - OpStore %9 %27 - %29 = OpLoad %bool %9 - OpBranchConditional %29 %10 %11 + %34 = OpIAdd %ulong %21 %ulong_8 + %38 = OpConvertUToPtr %_ptr_Generic_ulong %34 + %20 = OpLoad %ulong %38 + OpStore %7 %20 + %23 = OpLoad %ulong %6 + %24 = OpLoad %ulong %7 + %22 = OpULessThan %bool %23 %24 + OpStore %9 %22 + %26 = OpLoad %bool %9 + %25 = OpSelect %bool %26 %false %true + OpStore %9 %25 + %27 = OpLoad %bool %9 + OpBranchConditional %27 %10 %11 %10 = OpLabel - %30 = OpCopyObject %ulong %ulong_1 - OpStore %8 %30 + %28 = OpCopyObject %ulong %ulong_1 + OpStore %8 %28 OpBranch %11 %11 = OpLabel - %31 = OpLoad %bool %9 - OpBranchConditional %31 %13 %12 + %29 = OpLoad %bool %9 + OpBranchConditional %29 %13 %12 %12 = OpLabel - %32 = OpCopyObject %ulong %ulong_2 - OpStore %8 %32 + %30 = OpCopyObject %ulong %ulong_2 + OpStore %8 %30 OpBranch %13 %13 = OpLabel - %33 = OpLoad %ulong %5 - %34 = OpLoad %ulong %8 - %41 = OpConvertUToPtr %_ptr_Generic_ulong %33 - OpStore %41 %34 + %31 = OpLoad %ulong %5 + %32 = OpLoad %ulong %8 + %39 = OpConvertUToPtr %_ptr_Generic_ulong %31 + OpStore %39 %32 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/rcp.spvtxt b/ptx/src/test/spirv_run/rcp.spvtxt index fd10ff1..0ce2d75 100644 --- a/ptx/src/test/spirv_run/rcp.spvtxt +++ b/ptx/src/test/spirv_run/rcp.spvtxt @@ -7,24 +7,22 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - OpExtension "SPV_KHR_float_controls" - %23 = OpExtInstImport "OpenCL.std" + %21 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "rcp" - OpExecutionMode %1 DenormPreserve 32 - OpDecorate %15 FPFastMathMode AllowRecip + OpDecorate %13 FPFastMathMode AllowRecip %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %26 = OpTypeFunction %void %ulong %ulong + %24 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_Function_float = OpTypePointer Function %float %_ptr_Generic_float = OpTypePointer Generic %float %float_1 = OpConstant %float 1 - %1 = OpFunction %void None %26 + %1 = OpFunction %void None %24 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %21 = OpLabel + %19 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -32,22 +30,20 @@ %6 = OpVariable %_ptr_Function_float Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 - %19 = OpConvertUToPtr %_ptr_Generic_float %14 - %13 = OpLoad %float %19 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpFDiv %float %float_1 %14 OpStore %6 %13 + %15 = OpLoad %ulong %5 %16 = OpLoad %float %6 - %15 = OpFDiv %float %float_1 %16 - OpStore %6 %15 - %17 = OpLoad %ulong %5 - %18 = OpLoad %float %6 - %20 = OpConvertUToPtr %_ptr_Generic_float %17 - OpStore %20 %18 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/reg_local.ptx b/ptx/src/test/spirv_run/reg_local.ptx index f09b95a..5707cea 100644 --- a/ptx/src/test/spirv_run/reg_local.ptx +++ b/ptx/src/test/spirv_run/reg_local.ptx @@ -11,14 +11,13 @@ .reg .u64 in_addr; .reg .u64 out_addr; .reg .b64 temp; - .reg .s64 unused; ld.param.u64 in_addr, [input]; ld.param.u64 out_addr, [output]; ld.global.u64 temp, [in_addr]; st.u64 [local_x], temp + 1; - ld.u64 temp, [local_x]; - st.global.u64 [out_addr], temp; + ld.u64 temp, [local_x+0]; + st.global.u64 [out_addr+0], temp; ret; } \ No newline at end of file diff --git a/ptx/src/test/spirv_run/reg_local.spvtxt b/ptx/src/test/spirv_run/reg_local.spvtxt index 2d6bd08..596cedc 100644 --- a/ptx/src/test/spirv_run/reg_local.spvtxt +++ b/ptx/src/test/spirv_run/reg_local.spvtxt @@ -2,62 +2,66 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %35 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %34 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "reg_local" OpDecorate %4 Alignment 8 %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %38 = OpTypeFunction %void %ulong %ulong + %37 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong - %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 %uint_8 = OpConstant %uint 8 %_arr_uchar_uint_8 = OpTypeArray %uchar %uint_8 %_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8 %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %ulong_1 = OpConstant %ulong 1 - %1 = OpFunction %void None %38 + %ulong_0 = OpConstant %ulong 0 + %ulong_0_0 = OpConstant %ulong 0 + %1 = OpFunction %void None %37 + %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %10 = OpFunctionParameter %ulong - %33 = OpLabel + %32 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function__arr_uchar_uint_8 Function %5 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function - %8 = OpVariable %_ptr_Function_ulong Function - OpStore %2 %9 - OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %6 %13 - %16 = OpLoad %ulong %5 - %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %16 - %26 = OpLoad %ulong %25 - %15 = OpCopyObject %ulong %26 + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %5 %10 + %11 = OpLoad %ulong %3 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %13 + %24 = OpLoad %ulong %25 + %12 = OpCopyObject %ulong %24 + OpStore %7 %12 + %14 = OpLoad %ulong %7 + %26 = OpCopyObject %ulong %14 + %19 = OpIAdd %ulong %26 %ulong_1 + %27 = OpBitcast %_ptr_Function_ulong %4 + OpStore %27 %19 + %28 = OpBitcast %_ptr_Function_ulong %4 + %45 = OpBitcast %ulong %28 + %46 = OpIAdd %ulong %45 %ulong_0 + %21 = OpBitcast %_ptr_Function_ulong %46 + %29 = OpLoad %ulong %21 + %15 = OpCopyObject %ulong %29 OpStore %7 %15 - %18 = OpLoad %ulong %7 - %27 = OpCopyObject %ulong %18 - %24 = OpIAdd %ulong %27 %ulong_1 - %28 = OpCopyObject %ulong %24 - %17 = OpBitcast %ulong %28 - OpStore %4 %17 - %20 = OpLoad %_arr_uchar_uint_8 %4 - %29 = OpBitcast %ulong %20 - %30 = OpCopyObject %ulong %29 - %19 = OpCopyObject %ulong %30 - OpStore %7 %19 - %21 = OpLoad %ulong %6 - %22 = OpLoad %ulong %7 - %31 = OpCopyObject %ulong %22 - %32 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %21 - OpStore %32 %31 + %16 = OpLoad %ulong %6 + %17 = OpLoad %ulong %7 + %23 = OpIAdd %ulong %16 %ulong_0_0 + %30 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %23 + %31 = OpCopyObject %ulong %17 + OpStore %30 %31 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/rsqrt.spvtxt b/ptx/src/test/spirv_run/rsqrt.spvtxt index 5c3ba97..ed473bc 100644 --- a/ptx/src/test/spirv_run/rsqrt.spvtxt +++ b/ptx/src/test/spirv_run/rsqrt.spvtxt @@ -1,56 +1,47 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 31 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%23 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "rsqrt" -OpDecorate %1 FunctionDenormModeINTEL 64 Preserve -%24 = OpTypeVoid -%25 = OpTypeInt 64 0 -%26 = OpTypeFunction %24 %25 %25 -%27 = OpTypePointer Function %25 -%28 = OpTypeFloat 64 -%29 = OpTypePointer Function %28 -%30 = OpTypePointer Generic %28 -%1 = OpFunction %24 None %26 -%7 = OpFunctionParameter %25 -%8 = OpFunctionParameter %25 -%21 = OpLabel -%2 = OpVariable %27 Function -%3 = OpVariable %27 Function -%4 = OpVariable %27 Function -%5 = OpVariable %27 Function -%6 = OpVariable %29 Function -OpStore %2 %7 -OpStore %3 %8 -%10 = OpLoad %25 %2 -%9 = OpCopyObject %25 %10 -OpStore %4 %9 -%12 = OpLoad %25 %3 -%11 = OpCopyObject %25 %12 -OpStore %5 %11 -%14 = OpLoad %25 %4 -%19 = OpConvertUToPtr %30 %14 -%13 = OpLoad %28 %19 -OpStore %6 %13 -%16 = OpLoad %28 %6 -%15 = OpExtInst %28 %23 native_rsqrt %16 -OpStore %6 %15 -%17 = OpLoad %25 %5 -%18 = OpLoad %28 %6 -%20 = OpConvertUToPtr %30 %17 -OpStore %20 %18 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "rsqrt" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %double = OpTypeFloat 64 +%_ptr_Function_double = OpTypePointer Function %double +%_ptr_Generic_double = OpTypePointer Generic %double + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_double Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_double %12 + %11 = OpLoad %double %17 + OpStore %6 %11 + %14 = OpLoad %double %6 + %13 = OpExtInst %double %21 native_rsqrt %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %double %6 + %18 = OpConvertUToPtr %_ptr_Generic_double %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/selp.spvtxt b/ptx/src/test/spirv_run/selp.spvtxt index 6f73bc2..44e87e4 100644 --- a/ptx/src/test/spirv_run/selp.spvtxt +++ b/ptx/src/test/spirv_run/selp.spvtxt @@ -1,65 +1,57 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 40 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%31 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "selp" -%32 = OpTypeVoid -%33 = OpTypeInt 64 0 -%34 = OpTypeFunction %32 %33 %33 -%35 = OpTypePointer Function %33 -%36 = OpTypeInt 16 0 -%37 = OpTypePointer Function %36 -%38 = OpTypePointer Generic %36 -%23 = OpConstant %33 2 -%39 = OpTypeBool -%25 = OpConstantFalse %39 -%1 = OpFunction %32 None %34 -%8 = OpFunctionParameter %33 -%9 = OpFunctionParameter %33 -%29 = OpLabel -%2 = OpVariable %35 Function -%3 = OpVariable %35 Function -%4 = OpVariable %35 Function -%5 = OpVariable %35 Function -%6 = OpVariable %37 Function -%7 = OpVariable %37 Function -OpStore %2 %8 -OpStore %3 %9 -%11 = OpLoad %33 %2 -%10 = OpCopyObject %33 %11 -OpStore %4 %10 -%13 = OpLoad %33 %3 -%12 = OpCopyObject %33 %13 -OpStore %5 %12 -%15 = OpLoad %33 %4 -%26 = OpConvertUToPtr %38 %15 -%14 = OpLoad %36 %26 -OpStore %6 %14 -%17 = OpLoad %33 %4 -%24 = OpIAdd %33 %17 %23 -%27 = OpConvertUToPtr %38 %24 -%16 = OpLoad %36 %27 -OpStore %7 %16 -%19 = OpLoad %36 %6 -%20 = OpLoad %36 %7 -%18 = OpSelect %36 %25 %20 %20 -OpStore %6 %18 -%21 = OpLoad %33 %5 -%22 = OpLoad %36 %6 -%28 = OpConvertUToPtr %38 %21 -OpStore %28 %22 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %29 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "selp" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %32 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %ushort = OpTypeInt 16 0 +%_ptr_Function_ushort = OpTypePointer Function %ushort +%_ptr_Generic_ushort = OpTypePointer Generic %ushort + %ulong_2 = OpConstant %ulong 2 + %bool = OpTypeBool + %false = OpConstantFalse %bool + %1 = OpFunction %void None %32 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %27 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_ushort Function + %7 = OpVariable %_ptr_Function_ushort Function + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %4 %10 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %24 = OpConvertUToPtr %_ptr_Generic_ushort %13 + %12 = OpLoad %ushort %24 + OpStore %6 %12 + %15 = OpLoad %ulong %4 + %22 = OpIAdd %ulong %15 %ulong_2 + %25 = OpConvertUToPtr %_ptr_Generic_ushort %22 + %14 = OpLoad %ushort %25 + OpStore %7 %14 + %17 = OpLoad %ushort %6 + %18 = OpLoad %ushort %7 + %16 = OpSelect %ushort %false %18 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %ushort %6 + %26 = OpConvertUToPtr %_ptr_Generic_ushort %19 + OpStore %26 %20 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/setp.spvtxt b/ptx/src/test/spirv_run/setp.spvtxt index 5e18377..ec94573 100644 --- a/ptx/src/test/spirv_run/setp.spvtxt +++ b/ptx/src/test/spirv_run/setp.spvtxt @@ -2,14 +2,17 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %42 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %40 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "setp" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %45 = OpTypeFunction %void %ulong %ulong + %43 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %bool = OpTypeBool %_ptr_Function_bool = OpTypePointer Function %bool @@ -17,10 +20,10 @@ %ulong_8 = OpConstant %ulong 8 %ulong_1 = OpConstant %ulong 1 %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %45 + %1 = OpFunction %void None %43 %14 = OpFunctionParameter %ulong %15 = OpFunctionParameter %ulong - %40 = OpLabel + %38 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -31,42 +34,40 @@ %9 = OpVariable %_ptr_Function_bool Function OpStore %2 %14 OpStore %3 %15 - %17 = OpLoad %ulong %2 - %16 = OpCopyObject %ulong %17 + %16 = OpLoad %ulong %2 OpStore %4 %16 - %19 = OpLoad %ulong %3 - %18 = OpCopyObject %ulong %19 - OpStore %5 %18 + %17 = OpLoad %ulong %3 + OpStore %5 %17 + %19 = OpLoad %ulong %4 + %35 = OpConvertUToPtr %_ptr_Generic_ulong %19 + %18 = OpLoad %ulong %35 + OpStore %6 %18 %21 = OpLoad %ulong %4 - %37 = OpConvertUToPtr %_ptr_Generic_ulong %21 - %20 = OpLoad %ulong %37 - OpStore %6 %20 - %23 = OpLoad %ulong %4 - %34 = OpIAdd %ulong %23 %ulong_8 - %38 = OpConvertUToPtr %_ptr_Generic_ulong %34 - %22 = OpLoad %ulong %38 - OpStore %7 %22 - %25 = OpLoad %ulong %6 - %26 = OpLoad %ulong %7 - %24 = OpULessThan %bool %25 %26 - OpStore %9 %24 - %27 = OpLoad %bool %9 - OpBranchConditional %27 %10 %11 + %32 = OpIAdd %ulong %21 %ulong_8 + %36 = OpConvertUToPtr %_ptr_Generic_ulong %32 + %20 = OpLoad %ulong %36 + OpStore %7 %20 + %23 = OpLoad %ulong %6 + %24 = OpLoad %ulong %7 + %22 = OpULessThan %bool %23 %24 + OpStore %9 %22 + %25 = OpLoad %bool %9 + OpBranchConditional %25 %10 %11 %10 = OpLabel - %28 = OpCopyObject %ulong %ulong_1 - OpStore %8 %28 + %26 = OpCopyObject %ulong %ulong_1 + OpStore %8 %26 OpBranch %11 %11 = OpLabel - %29 = OpLoad %bool %9 - OpBranchConditional %29 %13 %12 + %27 = OpLoad %bool %9 + OpBranchConditional %27 %13 %12 %12 = OpLabel - %30 = OpCopyObject %ulong %ulong_2 - OpStore %8 %30 + %28 = OpCopyObject %ulong %ulong_2 + OpStore %8 %28 OpBranch %13 %13 = OpLabel - %31 = OpLoad %ulong %5 - %32 = OpLoad %ulong %8 - %39 = OpConvertUToPtr %_ptr_Generic_ulong %31 - OpStore %39 %32 + %29 = OpLoad %ulong %5 + %30 = OpLoad %ulong %8 + %37 = OpConvertUToPtr %_ptr_Generic_ulong %29 + OpStore %37 %30 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/shared_ptr_32.spvtxt b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt index 609cc0e..98e2501 100644 --- a/ptx/src/test/spirv_run/shared_ptr_32.spvtxt +++ b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt @@ -1,74 +1,66 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 47 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%34 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "shared_ptr_32" %4 -OpDecorate %4 Alignment 4 -%35 = OpTypeVoid -%36 = OpTypeInt 32 0 -%37 = OpTypeInt 8 0 -%38 = OpConstant %36 128 -%39 = OpTypeArray %37 %38 -%40 = OpTypePointer Workgroup %39 -%4 = OpVariable %40 Workgroup -%41 = OpTypeInt 64 0 -%42 = OpTypeFunction %35 %41 %41 -%43 = OpTypePointer Function %41 -%44 = OpTypePointer Function %36 -%45 = OpTypePointer CrossWorkgroup %41 -%46 = OpTypePointer Workgroup %41 -%25 = OpConstant %36 0 -%1 = OpFunction %35 None %42 -%10 = OpFunctionParameter %41 -%11 = OpFunctionParameter %41 -%32 = OpLabel -%2 = OpVariable %43 Function -%3 = OpVariable %43 Function -%5 = OpVariable %43 Function -%6 = OpVariable %43 Function -%7 = OpVariable %44 Function -%8 = OpVariable %43 Function -%9 = OpVariable %43 Function -OpStore %2 %10 -OpStore %3 %11 -%13 = OpLoad %41 %2 -%12 = OpCopyObject %41 %13 -OpStore %5 %12 -%15 = OpLoad %41 %3 -%14 = OpCopyObject %41 %15 -OpStore %6 %14 -%27 = OpConvertPtrToU %36 %4 -%16 = OpCopyObject %36 %27 -OpStore %7 %16 -%18 = OpLoad %41 %5 -%28 = OpConvertUToPtr %45 %18 -%17 = OpLoad %41 %28 -OpStore %8 %17 -%19 = OpLoad %36 %7 -%20 = OpLoad %41 %8 -%29 = OpConvertUToPtr %46 %19 -OpStore %29 %20 -%22 = OpLoad %36 %7 -%26 = OpIAdd %36 %22 %25 -%30 = OpConvertUToPtr %46 %26 -%21 = OpLoad %41 %30 -OpStore %9 %21 -%23 = OpLoad %41 %6 -%24 = OpLoad %41 %9 -%31 = OpConvertUToPtr %45 %23 -OpStore %31 %24 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %32 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "shared_ptr_32" %4 + OpDecorate %4 Alignment 4 + %void = OpTypeVoid + %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 + %uint_128 = OpConstant %uint 128 +%_arr_uchar_uint_128 = OpTypeArray %uchar %uint_128 +%_ptr_Workgroup__arr_uchar_uint_128 = OpTypePointer Workgroup %_arr_uchar_uint_128 + %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup + %ulong = OpTypeInt 64 0 + %40 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong +%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong + %uint_0 = OpConstant %uint 0 + %1 = OpFunction %void None %40 + %10 = OpFunctionParameter %ulong + %11 = OpFunctionParameter %ulong + %30 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_ulong Function + %7 = OpVariable %_ptr_Function_uint Function + %8 = OpVariable %_ptr_Function_ulong Function + %9 = OpVariable %_ptr_Function_ulong Function + OpStore %2 %10 + OpStore %3 %11 + %12 = OpLoad %ulong %2 + OpStore %5 %12 + %13 = OpLoad %ulong %3 + OpStore %6 %13 + %25 = OpConvertPtrToU %uint %4 + %14 = OpCopyObject %uint %25 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %16 + %15 = OpLoad %ulong %26 + OpStore %8 %15 + %17 = OpLoad %uint %7 + %18 = OpLoad %ulong %8 + %27 = OpConvertUToPtr %_ptr_Workgroup_ulong %17 + OpStore %27 %18 + %20 = OpLoad %uint %7 + %24 = OpIAdd %uint %20 %uint_0 + %28 = OpConvertUToPtr %_ptr_Workgroup_ulong %24 + %19 = OpLoad %ulong %28 + OpStore %9 %19 + %21 = OpLoad %ulong %6 + %22 = OpLoad %ulong %9 + %29 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %21 + OpStore %29 %22 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/shared_variable.spvtxt b/ptx/src/test/spirv_run/shared_variable.spvtxt index 1af2bd1..ffd6bd6 100644 --- a/ptx/src/test/spirv_run/shared_variable.spvtxt +++ b/ptx/src/test/spirv_run/shared_variable.spvtxt @@ -1,65 +1,57 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 39 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%27 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "shared_variable" %4 -OpDecorate %4 Alignment 4 -%28 = OpTypeVoid -%29 = OpTypeInt 32 0 -%30 = OpTypeInt 8 0 -%31 = OpConstant %29 128 -%32 = OpTypeArray %30 %31 -%33 = OpTypePointer Workgroup %32 -%4 = OpVariable %33 Workgroup -%34 = OpTypeInt 64 0 -%35 = OpTypeFunction %28 %34 %34 -%36 = OpTypePointer Function %34 -%37 = OpTypePointer CrossWorkgroup %34 -%38 = OpTypePointer Workgroup %34 -%1 = OpFunction %28 None %35 -%9 = OpFunctionParameter %34 -%10 = OpFunctionParameter %34 -%25 = OpLabel -%2 = OpVariable %36 Function -%3 = OpVariable %36 Function -%5 = OpVariable %36 Function -%6 = OpVariable %36 Function -%7 = OpVariable %36 Function -%8 = OpVariable %36 Function -OpStore %2 %9 -OpStore %3 %10 -%12 = OpLoad %34 %2 -%11 = OpCopyObject %34 %12 -OpStore %5 %11 -%14 = OpLoad %34 %3 -%13 = OpCopyObject %34 %14 -OpStore %6 %13 -%16 = OpLoad %34 %5 -%21 = OpConvertUToPtr %37 %16 -%15 = OpLoad %34 %21 -OpStore %7 %15 -%17 = OpLoad %34 %7 -%22 = OpBitcast %38 %4 -OpStore %22 %17 -%23 = OpBitcast %38 %4 -%18 = OpLoad %34 %23 -OpStore %8 %18 -%19 = OpLoad %34 %6 -%20 = OpLoad %34 %8 -%24 = OpConvertUToPtr %37 %19 -OpStore %24 %20 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %25 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "shared_variable" %4 + OpDecorate %4 Alignment 4 + %void = OpTypeVoid + %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 + %uint_128 = OpConstant %uint 128 +%_arr_uchar_uint_128 = OpTypeArray %uchar %uint_128 +%_ptr_Workgroup__arr_uchar_uint_128 = OpTypePointer Workgroup %_arr_uchar_uint_128 + %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup + %ulong = OpTypeInt 64 0 + %33 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong +%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong +%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong + %1 = OpFunction %void None %33 + %9 = OpFunctionParameter %ulong + %10 = OpFunctionParameter %ulong + %23 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_ulong Function + %7 = OpVariable %_ptr_Function_ulong Function + %8 = OpVariable %_ptr_Function_ulong Function + OpStore %2 %9 + OpStore %3 %10 + %11 = OpLoad %ulong %2 + OpStore %5 %11 + %12 = OpLoad %ulong %3 + OpStore %6 %12 + %14 = OpLoad %ulong %5 + %19 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %14 + %13 = OpLoad %ulong %19 + OpStore %7 %13 + %15 = OpLoad %ulong %7 + %20 = OpBitcast %_ptr_Workgroup_ulong %4 + OpStore %20 %15 + %21 = OpBitcast %_ptr_Workgroup_ulong %4 + %16 = OpLoad %ulong %21 + OpStore %8 %16 + %17 = OpLoad %ulong %6 + %18 = OpLoad %ulong %8 + %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %17 + OpStore %22 %18 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/shl.spvtxt b/ptx/src/test/spirv_run/shl.spvtxt index 4843a65..ce19fa5 100644 --- a/ptx/src/test/spirv_run/shl.spvtxt +++ b/ptx/src/test/spirv_run/shl.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %27 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %25 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "shl" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %30 = OpTypeFunction %void %ulong %ulong + %28 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %uint = OpTypeInt 32 0 %uint_2 = OpConstant %uint 2 - %1 = OpFunction %void None %30 + %1 = OpFunction %void None %28 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %25 = OpLabel + %23 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -26,24 +29,22 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %23 = OpCopyObject %ulong %17 - %22 = OpShiftLeftLogical %ulong %23 %uint_2 - %16 = OpCopyObject %ulong %22 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %24 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %24 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %21 = OpCopyObject %ulong %15 + %20 = OpShiftLeftLogical %ulong %21 %uint_2 + %14 = OpCopyObject %ulong %20 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %22 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %22 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/shr.spvtxt b/ptx/src/test/spirv_run/shr.spvtxt index 417839d..893dbf3 100644 --- a/ptx/src/test/spirv_run/shr.spvtxt +++ b/ptx/src/test/spirv_run/shr.spvtxt @@ -7,21 +7,21 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %24 = OpExtInstImport "OpenCL.std" + %22 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "shr" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %27 = OpTypeFunction %void %ulong %ulong + %25 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint %uint_1 = OpConstant %uint 1 - %1 = OpFunction %void None %27 + %1 = OpFunction %void None %25 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %22 = OpLabel + %20 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -29,22 +29,20 @@ %6 = OpVariable %_ptr_Function_uint Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 - %20 = OpConvertUToPtr %_ptr_Generic_uint %14 - %13 = OpLoad %uint %20 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %18 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpShiftRightArithmetic %uint %14 %uint_1 OpStore %6 %13 + %15 = OpLoad %ulong %5 %16 = OpLoad %uint %6 - %15 = OpShiftRightArithmetic %uint %16 %uint_1 - OpStore %6 %15 - %17 = OpLoad %ulong %5 - %18 = OpLoad %uint %6 - %21 = OpConvertUToPtr %_ptr_Generic_uint %17 - OpStore %21 %18 + %19 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %19 %16 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/sin.ptx b/ptx/src/test/spirv_run/sin.ptx new file mode 100644 index 0000000..fe94cac --- /dev/null +++ b/ptx/src/test/spirv_run/sin.ptx @@ -0,0 +1,21 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.visible .entry sin( + .param .u64 input, + .param .u64 output +) +{ + .reg .u64 in_addr; + .reg .u64 out_addr; + .reg .f32 temp; + + ld.param.u64 in_addr, [input]; + ld.param.u64 out_addr, [output]; + + ld.f32 temp, [in_addr]; + sin.approx.f32 temp, temp; + st.f32 [out_addr], temp; + ret; +} diff --git a/ptx/src/test/spirv_run/sin.spvtxt b/ptx/src/test/spirv_run/sin.spvtxt new file mode 100644 index 0000000..6656a43 --- /dev/null +++ b/ptx/src/test/spirv_run/sin.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "sin" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 sin %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/sqrt.spvtxt b/ptx/src/test/spirv_run/sqrt.spvtxt index d2c5b20..6d1cfd2 100644 --- a/ptx/src/test/spirv_run/sqrt.spvtxt +++ b/ptx/src/test/spirv_run/sqrt.spvtxt @@ -1,56 +1,47 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 31 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%23 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "sqrt" -OpDecorate %1 FunctionDenormModeINTEL 32 Preserve -%24 = OpTypeVoid -%25 = OpTypeInt 64 0 -%26 = OpTypeFunction %24 %25 %25 -%27 = OpTypePointer Function %25 -%28 = OpTypeFloat 32 -%29 = OpTypePointer Function %28 -%30 = OpTypePointer Generic %28 -%1 = OpFunction %24 None %26 -%7 = OpFunctionParameter %25 -%8 = OpFunctionParameter %25 -%21 = OpLabel -%2 = OpVariable %27 Function -%3 = OpVariable %27 Function -%4 = OpVariable %27 Function -%5 = OpVariable %27 Function -%6 = OpVariable %29 Function -OpStore %2 %7 -OpStore %3 %8 -%10 = OpLoad %25 %2 -%9 = OpCopyObject %25 %10 -OpStore %4 %9 -%12 = OpLoad %25 %3 -%11 = OpCopyObject %25 %12 -OpStore %5 %11 -%14 = OpLoad %25 %4 -%19 = OpConvertUToPtr %30 %14 -%13 = OpLoad %28 %19 -OpStore %6 %13 -%16 = OpLoad %28 %6 -%15 = OpExtInst %28 %23 native_sqrt %16 -OpStore %6 %15 -%17 = OpLoad %25 %5 -%18 = OpLoad %28 %6 -%20 = OpConvertUToPtr %30 %17 -OpStore %20 %18 -OpReturn -OpFunctionEnd \ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "sqrt" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %4 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 native_sqrt %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/sub.spvtxt b/ptx/src/test/spirv_run/sub.spvtxt index 8520168..88017ae 100644 --- a/ptx/src/test/spirv_run/sub.spvtxt +++ b/ptx/src/test/spirv_run/sub.spvtxt @@ -7,19 +7,19 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %25 = OpExtInstImport "OpenCL.std" + %23 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "sub" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %28 = OpTypeFunction %void %ulong %ulong + %26 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_1 = OpConstant %ulong 1 - %1 = OpFunction %void None %28 + %1 = OpFunction %void None %26 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %23 = OpLabel + %21 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -28,22 +28,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpISub %ulong %17 %ulong_1 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %22 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpISub %ulong %15 %ulong_1 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %20 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/vector.spvtxt b/ptx/src/test/spirv_run/vector.spvtxt index ff0ee97..535e480 100644 --- a/ptx/src/test/spirv_run/vector.spvtxt +++ b/ptx/src/test/spirv_run/vector.spvtxt @@ -2,26 +2,29 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %60 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %57 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %31 "vector" %void = OpTypeVoid %uint = OpTypeInt 32 0 %v2uint = OpTypeVector %uint 2 - %64 = OpTypeFunction %v2uint %v2uint + %61 = OpTypeFunction %v2uint %v2uint %_ptr_Function_v2uint = OpTypePointer Function %v2uint %_ptr_Function_uint = OpTypePointer Function %uint %ulong = OpTypeInt 64 0 - %68 = OpTypeFunction %void %ulong %ulong + %65 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_v2uint = OpTypePointer Generic %v2uint - %1 = OpFunction %v2uint None %64 + %1 = OpFunction %v2uint None %61 %7 = OpFunctionParameter %v2uint %30 = OpLabel - %3 = OpVariable %_ptr_Function_v2uint Function %2 = OpVariable %_ptr_Function_v2uint Function + %3 = OpVariable %_ptr_Function_v2uint Function %4 = OpVariable %_ptr_Function_v2uint Function %5 = OpVariable %_ptr_Function_uint Function %6 = OpVariable %_ptr_Function_uint Function @@ -57,10 +60,10 @@ %26 = OpLoad %v2uint %2 OpReturnValue %26 OpFunctionEnd - %31 = OpFunction %void None %68 + %31 = OpFunction %void None %65 %40 = OpFunctionParameter %ulong %41 = OpFunctionParameter %ulong - %58 = OpLabel + %55 = OpLabel %32 = OpVariable %_ptr_Function_ulong Function %33 = OpVariable %_ptr_Function_ulong Function %34 = OpVariable %_ptr_Function_ulong Function @@ -71,27 +74,24 @@ %39 = OpVariable %_ptr_Function_ulong Function OpStore %32 %40 OpStore %33 %41 - %43 = OpLoad %ulong %32 - %42 = OpCopyObject %ulong %43 + %42 = OpLoad %ulong %32 OpStore %34 %42 - %45 = OpLoad %ulong %33 - %44 = OpCopyObject %ulong %45 - OpStore %35 %44 - %47 = OpLoad %ulong %34 - %54 = OpConvertUToPtr %_ptr_Generic_v2uint %47 - %46 = OpLoad %v2uint %54 + %43 = OpLoad %ulong %33 + OpStore %35 %43 + %45 = OpLoad %ulong %34 + %52 = OpConvertUToPtr %_ptr_Generic_v2uint %45 + %44 = OpLoad %v2uint %52 + OpStore %36 %44 + %47 = OpLoad %v2uint %36 + %46 = OpFunctionCall %v2uint %1 %47 OpStore %36 %46 %49 = OpLoad %v2uint %36 - %48 = OpFunctionCall %v2uint %1 %49 - OpStore %36 %48 + %53 = OpBitcast %ulong %49 + %48 = OpCopyObject %ulong %53 + OpStore %39 %48 + %50 = OpLoad %ulong %35 %51 = OpLoad %v2uint %36 - %55 = OpBitcast %ulong %51 - %56 = OpCopyObject %ulong %55 - %50 = OpCopyObject %ulong %56 - OpStore %39 %50 - %52 = OpLoad %ulong %35 - %53 = OpLoad %v2uint %36 - %57 = OpConvertUToPtr %_ptr_Generic_v2uint %52 - OpStore %57 %53 + %54 = OpConvertUToPtr %_ptr_Generic_v2uint %50 + OpStore %54 %51 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/vector_extract.spvtxt b/ptx/src/test/spirv_run/vector_extract.spvtxt index 45df3a8..4943189 100644 --- a/ptx/src/test/spirv_run/vector_extract.spvtxt +++ b/ptx/src/test/spirv_run/vector_extract.spvtxt @@ -7,12 +7,12 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %75 = OpExtInstImport "OpenCL.std" + %73 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "vector_extract" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %78 = OpTypeFunction %void %ulong %ulong + %76 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %ushort = OpTypeInt 16 0 %_ptr_Function_ushort = OpTypePointer Function %ushort @@ -21,10 +21,10 @@ %uchar = OpTypeInt 8 0 %v4uchar = OpTypeVector %uchar 4 %_ptr_CrossWorkgroup_v4uchar = OpTypePointer CrossWorkgroup %v4uchar - %1 = OpFunction %void None %78 + %1 = OpFunction %void None %76 %11 = OpFunctionParameter %ulong %12 = OpFunctionParameter %ulong - %73 = OpLabel + %71 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -36,89 +36,87 @@ %10 = OpVariable %_ptr_Function_v4ushort Function OpStore %2 %11 OpStore %3 %12 - %14 = OpLoad %ulong %2 - %13 = OpCopyObject %ulong %14 + %13 = OpLoad %ulong %2 OpStore %4 %13 - %16 = OpLoad %ulong %3 - %15 = OpCopyObject %ulong %16 - OpStore %5 %15 - %21 = OpLoad %ulong %4 - %63 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %21 - %45 = OpLoad %v4uchar %63 - %64 = OpCompositeExtract %uchar %45 0 + %14 = OpLoad %ulong %3 + OpStore %5 %14 + %19 = OpLoad %ulong %4 + %61 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %19 + %43 = OpLoad %v4uchar %61 + %62 = OpCompositeExtract %uchar %43 0 + %85 = OpBitcast %uchar %62 + %15 = OpUConvert %ushort %85 + %63 = OpCompositeExtract %uchar %43 1 + %86 = OpBitcast %uchar %63 + %16 = OpUConvert %ushort %86 + %64 = OpCompositeExtract %uchar %43 2 %87 = OpBitcast %uchar %64 %17 = OpUConvert %ushort %87 - %65 = OpCompositeExtract %uchar %45 1 + %65 = OpCompositeExtract %uchar %43 3 %88 = OpBitcast %uchar %65 %18 = OpUConvert %ushort %88 - %66 = OpCompositeExtract %uchar %45 2 - %89 = OpBitcast %uchar %66 - %19 = OpUConvert %ushort %89 - %67 = OpCompositeExtract %uchar %45 3 - %90 = OpBitcast %uchar %67 - %20 = OpUConvert %ushort %90 - OpStore %6 %17 - OpStore %7 %18 - OpStore %8 %19 - OpStore %9 %20 - %23 = OpLoad %ushort %7 - %24 = OpLoad %ushort %8 - %25 = OpLoad %ushort %9 - %26 = OpLoad %ushort %6 - %46 = OpUndef %v4ushort - %47 = OpCompositeInsert %v4ushort %23 %46 0 - %48 = OpCompositeInsert %v4ushort %24 %47 1 - %49 = OpCompositeInsert %v4ushort %25 %48 2 - %50 = OpCompositeInsert %v4ushort %26 %49 3 - %22 = OpCopyObject %v4ushort %50 - OpStore %10 %22 - %31 = OpLoad %v4ushort %10 - %51 = OpCopyObject %v4ushort %31 - %27 = OpCompositeExtract %ushort %51 0 - %28 = OpCompositeExtract %ushort %51 1 - %29 = OpCompositeExtract %ushort %51 2 - %30 = OpCompositeExtract %ushort %51 3 - OpStore %8 %27 - OpStore %9 %28 - OpStore %6 %29 - OpStore %7 %30 - %36 = OpLoad %ushort %8 - %37 = OpLoad %ushort %9 - %38 = OpLoad %ushort %6 - %39 = OpLoad %ushort %7 - %53 = OpUndef %v4ushort - %54 = OpCompositeInsert %v4ushort %36 %53 0 - %55 = OpCompositeInsert %v4ushort %37 %54 1 - %56 = OpCompositeInsert %v4ushort %38 %55 2 - %57 = OpCompositeInsert %v4ushort %39 %56 3 - %52 = OpCopyObject %v4ushort %57 - %32 = OpCompositeExtract %ushort %52 0 - %33 = OpCompositeExtract %ushort %52 1 - %34 = OpCompositeExtract %ushort %52 2 - %35 = OpCompositeExtract %ushort %52 3 - OpStore %9 %32 - OpStore %6 %33 - OpStore %7 %34 - OpStore %8 %35 - %40 = OpLoad %ulong %5 - %41 = OpLoad %ushort %6 - %42 = OpLoad %ushort %7 - %43 = OpLoad %ushort %8 - %44 = OpLoad %ushort %9 - %58 = OpUndef %v4uchar + OpStore %6 %15 + OpStore %7 %16 + OpStore %8 %17 + OpStore %9 %18 + %21 = OpLoad %ushort %7 + %22 = OpLoad %ushort %8 + %23 = OpLoad %ushort %9 + %24 = OpLoad %ushort %6 + %44 = OpUndef %v4ushort + %45 = OpCompositeInsert %v4ushort %21 %44 0 + %46 = OpCompositeInsert %v4ushort %22 %45 1 + %47 = OpCompositeInsert %v4ushort %23 %46 2 + %48 = OpCompositeInsert %v4ushort %24 %47 3 + %20 = OpCopyObject %v4ushort %48 + OpStore %10 %20 + %29 = OpLoad %v4ushort %10 + %49 = OpCopyObject %v4ushort %29 + %25 = OpCompositeExtract %ushort %49 0 + %26 = OpCompositeExtract %ushort %49 1 + %27 = OpCompositeExtract %ushort %49 2 + %28 = OpCompositeExtract %ushort %49 3 + OpStore %8 %25 + OpStore %9 %26 + OpStore %6 %27 + OpStore %7 %28 + %34 = OpLoad %ushort %8 + %35 = OpLoad %ushort %9 + %36 = OpLoad %ushort %6 + %37 = OpLoad %ushort %7 + %51 = OpUndef %v4ushort + %52 = OpCompositeInsert %v4ushort %34 %51 0 + %53 = OpCompositeInsert %v4ushort %35 %52 1 + %54 = OpCompositeInsert %v4ushort %36 %53 2 + %55 = OpCompositeInsert %v4ushort %37 %54 3 + %50 = OpCopyObject %v4ushort %55 + %30 = OpCompositeExtract %ushort %50 0 + %31 = OpCompositeExtract %ushort %50 1 + %32 = OpCompositeExtract %ushort %50 2 + %33 = OpCompositeExtract %ushort %50 3 + OpStore %9 %30 + OpStore %6 %31 + OpStore %7 %32 + OpStore %8 %33 + %38 = OpLoad %ulong %5 + %39 = OpLoad %ushort %6 + %40 = OpLoad %ushort %7 + %41 = OpLoad %ushort %8 + %42 = OpLoad %ushort %9 + %56 = OpUndef %v4uchar + %89 = OpBitcast %ushort %39 + %66 = OpUConvert %uchar %89 + %57 = OpCompositeInsert %v4uchar %66 %56 0 + %90 = OpBitcast %ushort %40 + %67 = OpUConvert %uchar %90 + %58 = OpCompositeInsert %v4uchar %67 %57 1 %91 = OpBitcast %ushort %41 %68 = OpUConvert %uchar %91 - %59 = OpCompositeInsert %v4uchar %68 %58 0 + %59 = OpCompositeInsert %v4uchar %68 %58 2 %92 = OpBitcast %ushort %42 %69 = OpUConvert %uchar %92 - %60 = OpCompositeInsert %v4uchar %69 %59 1 - %93 = OpBitcast %ushort %43 - %70 = OpUConvert %uchar %93 - %61 = OpCompositeInsert %v4uchar %70 %60 2 - %94 = OpBitcast %ushort %44 - %71 = OpUConvert %uchar %94 - %62 = OpCompositeInsert %v4uchar %71 %61 3 - %72 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %40 - OpStore %72 %62 + %60 = OpCompositeInsert %v4uchar %69 %59 3 + %70 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %38 + OpStore %70 %60 OpReturn OpFunctionEnd diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs index 36e15f9..7a0dd08 100644 --- a/ptx/src/translate.rs +++ b/ptx/src/translate.rs @@ -1,7 +1,7 @@ use crate::ast; use half::f16; use rspirv::{binary::Disassemble, dr}; -use std::{borrow::Cow, hash::Hash, iter, mem}; +use std::{borrow::Cow, convert::TryFrom, hash::Hash, iter, mem}; use std::{ collections::{hash_map, HashMap, HashSet}, convert::TryInto, @@ -50,20 +50,45 @@ impl From for SpirvType { ast::Type::Scalar(t) => SpirvType::Base(t.into()), ast::Type::Vector(typ, len) => SpirvType::Vector(typ.into(), len), ast::Type::Array(t, len) => SpirvType::Array(t.into(), len), - ast::Type::Pointer(ast::PointerType::Scalar(typ), state_space) => SpirvType::Pointer( - Box::new(SpirvType::Base(typ.into())), + ast::Type::Pointer(pointer_t, state_space) => SpirvType::Pointer( + Box::new(SpirvType::from(ast::Type::from(pointer_t))), state_space.to_spirv(), ), - ast::Type::Pointer(ast::PointerType::Vector(typ, len), state_space) => { - SpirvType::Pointer( - Box::new(SpirvType::Vector(typ.into(), len)), - state_space.to_spirv(), - ) + } + } +} + +impl From for ast::Type { + fn from(t: ast::PointerType) -> Self { + match t { + ast::PointerType::Scalar(t) => ast::Type::Scalar(t), + ast::PointerType::Vector(t, len) => ast::Type::Vector(t, len), + ast::PointerType::Array(t, dims) => ast::Type::Array(t, dims), + ast::PointerType::Pointer(t, space) => { + ast::Type::Pointer(ast::PointerType::Scalar(t), space) } } } } +impl ast::Type { + fn pointer_to(self, space: ast::LdStateSpace) -> Result { + Ok(match self { + ast::Type::Scalar(t) => ast::Type::Pointer(ast::PointerType::Scalar(t), space), + ast::Type::Vector(t, len) => { + ast::Type::Pointer(ast::PointerType::Vector(t, len), space) + } + ast::Type::Array(t, dims) => { + ast::Type::Pointer(ast::PointerType::Array(t, dims), space) + } + ast::Type::Pointer(ast::PointerType::Scalar(t), space) => { + ast::Type::Pointer(ast::PointerType::Pointer(t, space), space) + } + ast::Type::Pointer(_, _) => return Err(TranslateError::Unreachable), + }) + } +} + impl Into for ast::PointerStateSpace { fn into(self) -> spirv::StorageClass { match self { @@ -229,8 +254,8 @@ impl TypeWordMap { fn get_or_add_fn( &mut self, b: &mut dr::Builder, - mut out_params: impl ExactSizeIterator, in_params: impl ExactSizeIterator, + mut out_params: impl ExactSizeIterator, ) -> (spirv::Word, spirv::Word) { let (out_args, out_spirv_type) = if out_params.len() == 0 { (None, self.void()) @@ -486,7 +511,7 @@ fn emit_directives<'input>( map: &mut TypeWordMap, id_defs: &GlobalStringIdResolver<'input>, opencl_id: spirv::Word, - denorm_information: &HashMap, HashMap>, + denorm_information: &HashMap, HashMap>, call_map: &HashMap<&'input str, HashSet>, directives: Vec, kernel_info: &mut HashMap, @@ -516,7 +541,7 @@ fn emit_directives<'input>( map, &id_defs, &f.globals, - &f.func_decl, + &f.spirv_decl, &denorm_information, call_map, &directives, @@ -553,7 +578,7 @@ fn get_call_map<'input>( body: Some(statements), .. }) => { - let call_key = CallgraphKey::new(&func_decl); + let call_key = MethodName::new(&func_decl); for statement in statements { match statement { Statement::Call(call) => { @@ -569,28 +594,28 @@ fn get_call_map<'input>( let mut result = HashMap::new(); for (method_key, children) in directly_called_by.iter() { match method_key { - CallgraphKey::Kernel(name) => { + MethodName::Kernel(name) => { let mut visited = HashSet::new(); for child in children { add_call_map_single(&directly_called_by, &mut visited, *child); } result.insert(*name, visited); } - CallgraphKey::Func(_) => {} + MethodName::Func(_) => {} } } result } fn add_call_map_single<'input>( - directly_called_by: &MultiHashMap, spirv::Word>, + directly_called_by: &MultiHashMap, spirv::Word>, visited: &mut HashSet, current: spirv::Word, ) { if !visited.insert(current) { return; } - if let Some(children) = directly_called_by.get(&CallgraphKey::Func(current)) { + if let Some(children) = directly_called_by.get(&MethodName::Func(current)) { for child in children { add_call_map_single(directly_called_by, visited, *child); } @@ -645,8 +670,9 @@ fn convert_dynamic_shared_memory_usage<'input>( globals, body: Some(statements), import_as, + spirv_decl, }) => { - let call_key = CallgraphKey::new(&func_decl); + let call_key = MethodName::new(&func_decl); let statements = statements .into_iter() .map(|statement| match statement { @@ -667,6 +693,7 @@ fn convert_dynamic_shared_memory_usage<'input>( globals, body: Some(statements), import_as, + spirv_decl, }) } directive => directive, @@ -680,44 +707,34 @@ fn convert_dynamic_shared_memory_usage<'input>( .into_iter() .map(|directive| match directive { Directive::Method(Function { - mut func_decl, + func_decl, globals, body: Some(statements), import_as, + mut spirv_decl, }) => { - let call_key = CallgraphKey::new(&func_decl); - if !methods_using_extern_shared.contains(&call_key) { + if !methods_using_extern_shared.contains(&spirv_decl.name) { return Directive::Method(Function { func_decl, globals, body: Some(statements), import_as, + spirv_decl, }); } let shared_id_param = new_id(); - match &mut func_decl { - ast::MethodDecl::Func(_, _, input_args) => { - input_args.push(ast::Variable { - align: None, - v_type: ast::FnArgumentType::Shared, - array_init: Vec::new(), - name: shared_id_param, - }); + spirv_decl.input.push({ + ast::Variable { + align: None, + v_type: ast::Type::Pointer( + ast::PointerType::Scalar(ast::ScalarType::U8), + ast::LdStateSpace::Shared, + ), + array_init: Vec::new(), + name: shared_id_param, } - ast::MethodDecl::Kernel { - in_args, - uses_shared_mem, - .. - } => { - *uses_shared_mem = true; - in_args.push(ast::Variable { - align: None, - v_type: ast::KernelArgumentType::Shared, - array_init: Vec::new(), - name: shared_id_param, - }); - } - } + }); + spirv_decl.uses_shared_mem = true; let shared_var_id = new_id(); let shared_var = ExpandedStatement::Variable(ast::Variable { align: None, @@ -750,6 +767,7 @@ fn convert_dynamic_shared_memory_usage<'input>( globals, body: Some(new_statements), import_as, + spirv_decl, }) } directive => directive, @@ -761,7 +779,7 @@ fn replace_uses_of_shared_memory<'a>( result: &mut Vec, new_id: &mut impl FnMut() -> spirv::Word, extern_shared_decls: &HashMap, - methods_using_extern_shared: &mut HashSet>, + methods_using_extern_shared: &mut HashSet>, shared_id_param: spirv::Word, shared_var_id: spirv::Word, statements: Vec, @@ -772,7 +790,7 @@ fn replace_uses_of_shared_memory<'a>( // We can safely skip checking call arguments, // because there's simply no way to pass shared ptr // without converting it to .b64 first - if methods_using_extern_shared.contains(&CallgraphKey::Func(call.func)) { + if methods_using_extern_shared.contains(&MethodName::Func(call.func)) { call.param_list .push((shared_id_param, ast::FnArgumentType::Shared)); } @@ -809,13 +827,13 @@ fn replace_uses_of_shared_memory<'a>( } fn get_callers_of_extern_shared<'a>( - methods_using_extern_shared: &mut HashSet>, - directly_called_by: &MultiHashMap>, + methods_using_extern_shared: &mut HashSet>, + directly_called_by: &MultiHashMap>, ) { let direct_uses_of_extern_shared = methods_using_extern_shared .iter() .filter_map(|method| { - if let CallgraphKey::Func(f_id) = method { + if let MethodName::Func(f_id) = method { Some(*f_id) } else { None @@ -828,14 +846,14 @@ fn get_callers_of_extern_shared<'a>( } fn get_callers_of_extern_shared_single<'a>( - methods_using_extern_shared: &mut HashSet>, - directly_called_by: &MultiHashMap>, + methods_using_extern_shared: &mut HashSet>, + directly_called_by: &MultiHashMap>, fn_id: spirv::Word, ) { if let Some(callers) = directly_called_by.get(&fn_id) { for caller in callers { if methods_using_extern_shared.insert(*caller) { - if let CallgraphKey::Func(caller_fn) = caller { + if let MethodName::Func(caller_fn) = caller { get_callers_of_extern_shared_single( methods_using_extern_shared, directly_called_by, @@ -877,7 +895,7 @@ fn denorm_count_map_update_impl( // and emit suitable execution mode fn compute_denorm_information<'input>( module: &[Directive<'input>], -) -> HashMap, HashMap> { +) -> HashMap, HashMap> { let mut denorm_methods = HashMap::new(); for directive in module { match directive { @@ -888,7 +906,7 @@ fn compute_denorm_information<'input>( .. }) => { let mut flush_counter = DenormCountMap::new(); - let method_key = CallgraphKey::new(func_decl); + let method_key = MethodName::new(func_decl); for statement in statements { match statement { Statement::Instruction(inst) => { @@ -907,6 +925,7 @@ fn compute_denorm_information<'input>( Statement::Undef(_, _) => {} Statement::Label(_) => {} Statement::Variable(_) => {} + Statement::PtrAdd { .. } => {} } } denorm_methods.insert(method_key, flush_counter); @@ -933,16 +952,16 @@ fn compute_denorm_information<'input>( } #[derive(Hash, PartialEq, Eq, Copy, Clone)] -enum CallgraphKey<'input> { +enum MethodName<'input> { Kernel(&'input str), Func(spirv::Word), } -impl<'input> CallgraphKey<'input> { +impl<'input> MethodName<'input> { fn new(decl: &ast::MethodDecl<'input, spirv::Word>) -> Self { match decl { - ast::MethodDecl::Kernel { name, .. } => CallgraphKey::Kernel(name), - ast::MethodDecl::Func(_, id, _) => CallgraphKey::Func(*id), + ast::MethodDecl::Kernel { name, .. } => MethodName::Kernel(name), + ast::MethodDecl::Func(_, id, _) => MethodName::Func(*id), } } } @@ -979,30 +998,30 @@ fn emit_function_header<'a>( map: &mut TypeWordMap, defined_globals: &GlobalStringIdResolver<'a>, synthetic_globals: &[ast::Variable], - func_directive: &ast::MethodDecl, - denorm_information: &HashMap, HashMap>, + func_decl: &SpirvMethodDecl<'a>, + _denorm_information: &HashMap, HashMap>, call_map: &HashMap<&'a str, HashSet>, direcitves: &[Directive], kernel_info: &mut HashMap, ) -> Result<(), TranslateError> { - if let ast::MethodDecl::Kernel { - name, - in_args, - uses_shared_mem, - } = &func_directive - { - let args_lens = in_args.iter().map(|param| param.v_type.width()).collect(); + if let MethodName::Kernel(name) = func_decl.name { + let args_lens = func_decl + .input + .iter() + .map(|param| param.v_type.size_of()) + .collect(); kernel_info.insert( name.to_string(), KernelInfo { arguments_sizes: args_lens, - uses_shared_mem: *uses_shared_mem, + uses_shared_mem: func_decl.uses_shared_mem, }, ); } - let (ret_type, func_type) = get_function_type(builder, map, &func_directive); - let fn_id = match func_directive { - ast::MethodDecl::Kernel { name, .. } => { + let (ret_type, func_type) = + get_function_type(builder, map, &func_decl.input, &func_decl.output); + let fn_id = match func_decl.name { + MethodName::Kernel(name) => { let fn_id = defined_globals.get_id(name)?; let mut global_variables = defined_globals .variables_type_check @@ -1035,17 +1054,11 @@ fn emit_function_header<'a>( _ => {} } } - global_variables.append(&mut interface); - builder.entry_point( - spirv::ExecutionModel::Kernel, - fn_id, - *name, - global_variables, - ); + builder.entry_point(spirv::ExecutionModel::Kernel, fn_id, name, global_variables); fn_id } - ast::MethodDecl::Func(_, name, _) => *name, + MethodName::Func(name) => name, }; builder.begin_function( ret_type, @@ -1053,7 +1066,9 @@ fn emit_function_header<'a>( spirv::FunctionControl::NONE, func_type, )?; - if let Some(denorm_modes) = denorm_information.get(&CallgraphKey::new(&func_directive)) { + // TODO: re-enable when Intel float control extension works + /* + if let Some(denorm_modes) = denorm_information.get(&func_decl.name) { for (size_of, denorm_mode) in denorm_modes { builder.decorate( fn_id, @@ -1065,16 +1080,17 @@ fn emit_function_header<'a>( ) } } - func_directive.visit_args(&mut |arg| { - let result_type = map.get_or_add(builder, ast::Type::from(arg.v_type.clone()).into()); + */ + for input in &func_decl.input { + let result_type = map.get_or_add(builder, SpirvType::from(input.v_type.clone())); let inst = dr::Instruction::new( spirv::Op::FunctionParameter, Some(result_type), - Some(arg.name), + Some(input.name), Vec::new(), ); builder.function.as_mut().unwrap().parameters.push(inst); - }); + } Ok(()) } @@ -1103,12 +1119,14 @@ fn emit_capabilities(builder: &mut dr::Builder) { builder.capability(spirv::Capability::Int64); builder.capability(spirv::Capability::Float16); builder.capability(spirv::Capability::Float64); - builder.capability(spirv::Capability::FunctionFloatControlINTEL); + // TODO: re-enable when Intel float control extension works + //builder.capability(spirv::Capability::FunctionFloatControlINTEL); } // http://htmlpreview.github.io/?https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/KHR/SPV_KHR_float_controls.html -fn emit_extensions(builder: &mut dr::Builder) { - builder.extension("SPV_INTEL_float_controls2"); +fn emit_extensions(_builder: &mut dr::Builder) { + // TODO: re-enable when Intel float control extension works + //builder.extension("SPV_INTEL_float_controls2"); } fn emit_opencl_import(builder: &mut dr::Builder) -> spirv::Word { @@ -1153,42 +1171,48 @@ fn translate_function<'a>( ptx_impl_imports: &mut HashMap, f: ast::ParsedFunction<'a>, ) -> Result, TranslateError> { - let (str_resolver, fn_resolver, fn_decl) = id_defs.start_fn(&f.func_directive); + let (str_resolver, fn_resolver, fn_decl) = id_defs.start_fn(&f.func_directive)?; to_ssa(ptx_impl_imports, str_resolver, fn_resolver, fn_decl, f.body) } fn expand_kernel_params<'a, 'b>( fn_resolver: &mut FnStringIdResolver<'a, 'b>, args: impl Iterator>, -) -> Vec> { - args.map(|a| ast::KernelArgument { - name: fn_resolver.add_def( - a.name, - Some((StateSpace::Param, ast::Type::from(a.v_type.clone()))), - ), - v_type: a.v_type.clone(), - align: a.align, - array_init: Vec::new(), +) -> Result>, TranslateError> { + args.map(|a| { + Ok(ast::KernelArgument { + name: fn_resolver.add_def( + a.name, + Some(( + StateSpace::Param, + ast::Type::from(a.v_type.clone()).pointer_to(ast::LdStateSpace::Param)?, + )), + ), + v_type: a.v_type.clone(), + align: a.align, + array_init: Vec::new(), + }) }) - .collect() + .collect::>() } fn expand_fn_params<'a, 'b>( fn_resolver: &mut FnStringIdResolver<'a, 'b>, args: impl Iterator>, -) -> Vec> { +) -> Result>, TranslateError> { args.map(|a| { + let var_type = a.v_type.to_func_type(); let ss = match a.v_type { ast::FnArgumentType::Reg(_) => StateSpace::Reg, ast::FnArgumentType::Param(_) => StateSpace::Param, ast::FnArgumentType::Shared => StateSpace::Shared, }; - ast::FnArgument { - name: fn_resolver.add_def(a.name, Some((ss, ast::Type::from(a.v_type.clone())))), + Ok(ast::FnArgument { + name: fn_resolver.add_def(a.name, Some((ss, var_type))), v_type: a.v_type.clone(), align: a.align, array_init: Vec::new(), - } + }) }) .collect() } @@ -1200,6 +1224,7 @@ fn to_ssa<'input, 'b>( f_args: ast::MethodDecl<'input, spirv::Word>, f_body: Option>>>, ) -> Result, TranslateError> { + let mut spirv_decl = SpirvMethodDecl::new(&f_args); let f_body = match f_body { Some(vec) => vec, None => { @@ -1208,6 +1233,7 @@ fn to_ssa<'input, 'b>( body: None, globals: Vec::new(), import_as: None, + spirv_decl, }) } }; @@ -1217,8 +1243,8 @@ fn to_ssa<'input, 'b>( let typed_statements = convert_to_typed_statements(unadorned_statements, &fn_defs, &numeric_id_defs)?; let mut numeric_id_defs = numeric_id_defs.finish(); - let (f_args, ssa_statements) = - insert_mem_ssa_statements(typed_statements, &mut numeric_id_defs, f_args)?; + let ssa_statements = + insert_mem_ssa_statements(typed_statements, &mut numeric_id_defs, &mut spirv_decl)?; let expanded_statements = expand_arguments(ssa_statements, &mut numeric_id_defs)?; let expanded_statements = insert_implicit_conversions(expanded_statements, &mut numeric_id_defs)?; @@ -1231,6 +1257,7 @@ fn to_ssa<'input, 'b>( globals: globals, body: Some(f_body), import_as: None, + spirv_decl, }) } @@ -1338,13 +1365,21 @@ fn convert_to_typed_statements( ast::Instruction::Call(call) => { // TODO: error out if lengths don't match let fn_def = fn_defs.get_fn_decl(call.func)?; - let ret_params = to_resolved_fn_args(call.ret_params, &*fn_def.ret_vals); - let param_list = to_resolved_fn_args(call.param_list, &*fn_def.params); + let out_args = to_resolved_fn_args(call.ret_params, &*fn_def.ret_vals); + let in_args = to_resolved_fn_args(call.param_list, &*fn_def.params); + let (out_params, out_non_params): (Vec<_>, Vec<_>) = out_args + .into_iter() + .partition(|(_, arg_type)| arg_type.is_param()); + let normalized_input_args = out_params + .into_iter() + .map(|(id, typ)| (ast::CallOperand::Reg(id), typ)) + .chain(in_args.into_iter()) + .collect(); let resolved_call = ResolvedCall { uniform: call.uniform, - ret_params, + ret_params: out_non_params, func: call.func, - param_list, + param_list: normalized_input_args, }; result.push(Statement::Call(resolved_call)); } @@ -1514,6 +1549,30 @@ fn convert_to_typed_statements( ast::Instruction::Neg(d, a) => { result.push(Statement::Instruction(ast::Instruction::Neg(d, a.cast()))) } + ast::Instruction::Sin { flush_to_zero, arg } => { + result.push(Statement::Instruction(ast::Instruction::Sin { + flush_to_zero, + arg: arg.cast(), + })) + } + ast::Instruction::Cos { flush_to_zero, arg } => { + result.push(Statement::Instruction(ast::Instruction::Cos { + flush_to_zero, + arg: arg.cast(), + })) + } + ast::Instruction::Lg2 { flush_to_zero, arg } => { + result.push(Statement::Instruction(ast::Instruction::Lg2 { + flush_to_zero, + arg: arg.cast(), + })) + } + ast::Instruction::Ex2 { flush_to_zero, arg } => { + result.push(Statement::Instruction(ast::Instruction::Ex2 { + flush_to_zero, + arg: arg.cast(), + })) + } }, Statement::Label(i) => result.push(Statement::Label(i)), Statement::Variable(v) => result.push(Statement::Variable(v)), @@ -1525,7 +1584,9 @@ fn convert_to_typed_statements( Statement::Conversion(c) => result.push(Statement::Conversion(c)), Statement::Constant(c) => result.push(Statement::Constant(c)), Statement::RetValue(d, id) => result.push(Statement::RetValue(d, id)), - Statement::Undef(_, _) => return Err(TranslateError::Unreachable), + Statement::Undef(_, _) | Statement::PtrAdd { .. } => { + return Err(TranslateError::Unreachable) + } } } Ok(result) @@ -1584,11 +1645,13 @@ fn to_ptx_impl_atomic_call( }, ], ); + let spirv_decl = SpirvMethodDecl::new(&func_decl); let func = Function { func_decl, globals: Vec::new(), body: None, import_as: Some(entry.key().clone()), + spirv_decl, }; entry.insert(Directive::Method(func)); fn_id @@ -1660,7 +1723,8 @@ fn normalize_labels( | Statement::Conversion(_) | Statement::Constant(_) | Statement::Label(_) - | Statement::Undef(_, _) => (), + | Statement::Undef(_, _) + | Statement::PtrAdd { .. } => {} } } iter::once(Statement::Label(id_def.new_id(None))) @@ -1716,66 +1780,45 @@ fn normalize_predicates( fn insert_mem_ssa_statements<'a, 'b>( func: Vec, id_def: &mut MutableNumericIdResolver, - mut f_args: ast::MethodDecl<'a, spirv::Word>, -) -> Result<(ast::MethodDecl<'a, spirv::Word>, Vec), TranslateError> { + fn_decl: &mut SpirvMethodDecl, +) -> Result, TranslateError> { let mut result = Vec::with_capacity(func.len()); - let out_param = match &mut f_args { - ast::MethodDecl::Kernel { in_args, .. } => { - for p in in_args.iter_mut() { - let typ = ast::Type::from(p.v_type.clone()); + for arg in fn_decl.output.iter() { + match type_to_variable_type(&arg.v_type)? { + Some(var_type) => { + result.push(Statement::Variable(ast::Variable { + align: arg.align, + v_type: var_type, + name: arg.name, + array_init: arg.array_init.clone(), + })); + } + None => return Err(TranslateError::Unreachable), + } + } + for arg in fn_decl.input.iter_mut() { + match type_to_variable_type(&arg.v_type)? { + Some(var_type) => { + let typ = arg.v_type.clone(); let new_id = id_def.new_id(typ.clone()); result.push(Statement::Variable(ast::Variable { - align: p.align, - v_type: ast::VariableType::Param(p.v_type.clone().to_param()), - name: p.name, - array_init: p.array_init.clone(), + align: arg.align, + v_type: var_type, + name: arg.name, + array_init: arg.array_init.clone(), })); result.push(Statement::StoreVar( ast::Arg2St { - src1: p.name, + src1: arg.name, src2: new_id, }, typ, )); - p.name = new_id; + arg.name = new_id; } - None + None => {} } - ast::MethodDecl::Func(out_params, _, in_params) => { - for p in in_params.iter_mut() { - let typ = ast::Type::from(p.v_type.clone()); - let new_id = id_def.new_id(typ.clone()); - let var_typ = ast::VariableType::from(p.v_type.clone()); - result.push(Statement::Variable(ast::Variable { - align: p.align, - v_type: var_typ, - name: p.name, - array_init: p.array_init.clone(), - })); - result.push(Statement::StoreVar( - ast::Arg2St { - src1: p.name, - src2: new_id, - }, - typ.clone(), - )); - p.name = new_id; - } - match &mut **out_params { - [p] => { - result.push(Statement::Variable(ast::Variable { - align: p.align, - v_type: ast::VariableType::from(p.v_type.clone()), - name: p.name, - array_init: p.array_init.clone(), - })); - Some(p.name) - } - [] => None, - _ => todo!(), - } - } - }; + } for s in func { match s { Statement::Call(call) => { @@ -1783,13 +1826,14 @@ fn insert_mem_ssa_statements<'a, 'b>( } Statement::Instruction(inst) => match inst { ast::Instruction::Ret(d) => { - if let Some(out_param) = out_param { - let typ = id_def.get_typed(out_param)?; + // TODO: handle multiple output args + if let &[out_param] = &fn_decl.output.as_slice() { + let typ = id_def.get_typed(out_param.name)?; let new_id = id_def.new_id(typ.clone()); result.push(Statement::LoadVar( ast::Arg2 { dst: new_id, - src: out_param, + src: out_param.name, }, typ.clone(), )); @@ -1818,11 +1862,31 @@ fn insert_mem_ssa_statements<'a, 'b>( | Statement::Conversion(_) | Statement::RetValue(_, _) | Statement::Constant(_) - | Statement::Undef(_, _) => {} + | Statement::Undef(_, _) + | Statement::PtrAdd { .. } => {} Statement::Composite(_) => todo!(), } } - Ok((f_args, result)) + Ok(result) +} + +fn type_to_variable_type(t: &ast::Type) -> Result, TranslateError> { + Ok(match t { + ast::Type::Scalar(typ) => Some(ast::VariableType::Reg(ast::VariableRegType::Scalar(*typ))), + ast::Type::Vector(typ, len) => Some(ast::VariableType::Reg(ast::VariableRegType::Vector( + (*typ) + .try_into() + .map_err(|_| TranslateError::MismatchedType)?, + *len, + ))), + ast::Type::Array(typ, len) => Some(ast::VariableType::Reg(ast::VariableRegType::Array( + (*typ) + .try_into() + .map_err(|_| TranslateError::MismatchedType)?, + len.clone(), + ))), + ast::Type::Pointer(_, _) => None, + }) } trait VisitVariable: Sized { @@ -1849,6 +1913,28 @@ trait VisitVariableExpanded { ) -> Result; } +struct VisitArgumentDescriptor<'a, Ctor: FnOnce(spirv::Word) -> ExpandedStatement> { + desc: ArgumentDescriptor, + typ: &'a ast::Type, + stmt_ctor: Ctor, +} + +impl<'a, Ctor: FnOnce(spirv::Word) -> ExpandedStatement> VisitVariableExpanded + for VisitArgumentDescriptor<'a, Ctor> +{ + fn visit_variable_extended< + F: FnMut( + ArgumentDescriptor, + Option<&ast::Type>, + ) -> Result, + >( + self, + f: &mut F, + ) -> Result { + f(self.desc, Some(self.typ)).map(self.stmt_ctor) + } +} + fn insert_mem_ssa_statement_default<'a, F: VisitVariable>( id_def: &mut MutableNumericIdResolver, result: &mut Vec, @@ -1857,7 +1943,7 @@ fn insert_mem_ssa_statement_default<'a, F: VisitVariable>( let mut post_statements = Vec::new(); let new_statement = stmt.visit_variable(&mut |desc: ArgumentDescriptor, instr_type| { - if instr_type.is_none() { + if instr_type.is_none() || desc.sema == ArgumentSemantics::RegisterPointer { return Ok(desc.op); } let id_type = match (id_def.get_typed(desc.op)?, desc.sema) { @@ -1925,6 +2011,56 @@ fn expand_arguments<'a, 'b>( name, array_init, })), + Statement::PtrAdd { + underlying_type, + state_space, + dst, + ptr_src, + constant_src, + } => { + let mut visitor = FlattenArguments::new(&mut result, id_def); + let sema = match state_space { + ast::LdStateSpace::Const + | ast::LdStateSpace::Global + | ast::LdStateSpace::Shared + | ast::LdStateSpace::Generic => ArgumentSemantics::PhysicalPointer, + ast::LdStateSpace::Local | ast::LdStateSpace::Param => { + ArgumentSemantics::RegisterPointer + } + }; + let ptr_type = ast::Type::Pointer(underlying_type.clone(), state_space); + let new_dst = visitor.id( + ArgumentDescriptor { + op: dst, + is_dst: true, + sema, + }, + Some(&ptr_type), + )?; + let new_ptr_src = visitor.id( + ArgumentDescriptor { + op: ptr_src, + is_dst: false, + sema, + }, + Some(&ptr_type), + )?; + let new_constant_src = visitor.id( + ArgumentDescriptor { + op: constant_src, + is_dst: false, + sema: ArgumentSemantics::Default, + }, + Some(&ast::Type::Scalar(ast::ScalarType::S64)), + )?; + result.push(Statement::PtrAdd { + underlying_type, + state_space, + dst: new_dst, + ptr_src: new_ptr_src, + constant_src: new_constant_src, + }) + } Statement::Label(id) => result.push(Statement::Label(id)), Statement::Conditional(bra) => result.push(Statement::Conditional(bra)), Statement::LoadVar(arg, typ) => result.push(Statement::LoadVar(arg, typ)), @@ -1990,83 +2126,94 @@ impl<'a, 'b> FlattenArguments<'a, 'b> { desc: ArgumentDescriptor<(spirv::Word, i32)>, typ: &ast::Type, ) -> Result { - let mut typ = typ.clone(); let (reg, offset) = desc.op; - match desc.sema { - ArgumentSemantics::Default - | ArgumentSemantics::DefaultRelaxed - | ArgumentSemantics::PhysicalPointer => { - if desc.sema == ArgumentSemantics::PhysicalPointer { - typ = self.id_def.get_typed(reg)?; - } - let (width, kind) = match typ { - ast::Type::Scalar(scalar_t) => { - let kind = match scalar_t.kind() { - kind @ ScalarKind::Bit - | kind @ ScalarKind::Unsigned - | kind @ ScalarKind::Signed => kind, - ScalarKind::Float => return Err(TranslateError::MismatchedType), - ScalarKind::Float2 => return Err(TranslateError::MismatchedType), - ScalarKind::Pred => return Err(TranslateError::MismatchedType), - }; - (scalar_t.size_of(), kind) - } - _ => return Err(TranslateError::MismatchedType), - }; - let arith_detail = if kind == ScalarKind::Signed { - ast::ArithDetails::Signed(ast::ArithSInt { - typ: ast::SIntType::from_size(width), - saturate: false, - }) - } else { - ast::ArithDetails::Unsigned(ast::UIntType::from_size(width)) - }; - let id_constant_stmt = self.id_def.new_id(typ.clone()); - let result_id = self.id_def.new_id(typ); - // TODO: check for edge cases around min value/max value/wrapping - if offset < 0 && kind != ScalarKind::Signed { + let add_type; + match typ { + ast::Type::Pointer(underlying_type, state_space) => { + let reg_typ = self.id_def.get_typed(reg)?; + if let ast::Type::Pointer(_, _) = reg_typ { + let id_constant_stmt = self.id_def.new_id(typ.clone()); self.func.push(Statement::Constant(ConstantDefinition { dst: id_constant_stmt, - typ: ast::ScalarType::from_parts(width, kind), - value: ast::ImmediateValue::U64(-(offset as i64) as u64), - })); - self.func.push(Statement::Instruction( - ast::Instruction::::Sub( - arith_detail, - ast::Arg3 { - dst: result_id, - src1: reg, - src2: id_constant_stmt, - }, - ), - )); - } else { - self.func.push(Statement::Constant(ConstantDefinition { - dst: id_constant_stmt, - typ: ast::ScalarType::from_parts(width, kind), + typ: ast::ScalarType::S64, value: ast::ImmediateValue::S64(offset as i64), })); - self.func.push(Statement::Instruction( - ast::Instruction::::Add( - arith_detail, - ast::Arg3 { - dst: result_id, - src1: reg, - src2: id_constant_stmt, - }, - ), - )); + let dst = self.id_def.new_id(typ.clone()); + self.func.push(Statement::PtrAdd { + underlying_type: underlying_type.clone(), + state_space: *state_space, + dst, + ptr_src: reg, + constant_src: id_constant_stmt, + }); + return Ok(dst); + } else { + add_type = self.id_def.get_typed(reg)?; } - Ok(result_id) } - ArgumentSemantics::RegisterPointer => { - if offset == 0 { - return Ok(reg); - } - todo!() + _ => { + add_type = typ.clone(); } - ArgumentSemantics::Address => todo!(), + }; + let (width, kind) = match add_type { + ast::Type::Scalar(scalar_t) => { + let kind = match scalar_t.kind() { + kind @ ScalarKind::Bit + | kind @ ScalarKind::Unsigned + | kind @ ScalarKind::Signed => kind, + ScalarKind::Float => return Err(TranslateError::MismatchedType), + ScalarKind::Float2 => return Err(TranslateError::MismatchedType), + ScalarKind::Pred => return Err(TranslateError::MismatchedType), + }; + (scalar_t.size_of(), kind) + } + _ => return Err(TranslateError::MismatchedType), + }; + let arith_detail = if kind == ScalarKind::Signed { + ast::ArithDetails::Signed(ast::ArithSInt { + typ: ast::SIntType::from_size(width), + saturate: false, + }) + } else { + ast::ArithDetails::Unsigned(ast::UIntType::from_size(width)) + }; + let id_constant_stmt = self.id_def.new_id(add_type.clone()); + let result_id = self.id_def.new_id(add_type); + // TODO: check for edge cases around min value/max value/wrapping + if offset < 0 && kind != ScalarKind::Signed { + self.func.push(Statement::Constant(ConstantDefinition { + dst: id_constant_stmt, + typ: ast::ScalarType::from_parts(width, kind), + value: ast::ImmediateValue::U64(-(offset as i64) as u64), + })); + self.func.push(Statement::Instruction( + ast::Instruction::::Sub( + arith_detail, + ast::Arg3 { + dst: result_id, + src1: reg, + src2: id_constant_stmt, + }, + ), + )); + } else { + self.func.push(Statement::Constant(ConstantDefinition { + dst: id_constant_stmt, + typ: ast::ScalarType::from_parts(width, kind), + value: ast::ImmediateValue::S64(offset as i64), + })); + self.func.push(Statement::Instruction( + ast::Instruction::::Add( + arith_detail, + ast::Arg3 { + dst: result_id, + src1: reg, + src2: id_constant_stmt, + }, + ), + )); } + Ok(result_id) } fn immediate( @@ -2281,6 +2428,36 @@ fn insert_implicit_conversions( should_bitcast_wrapper, None, )?, + Statement::PtrAdd { + underlying_type, + state_space, + dst, + ptr_src, + constant_src, + } => { + let visit_desc = VisitArgumentDescriptor { + desc: ArgumentDescriptor { + op: ptr_src, + is_dst: false, + sema: ArgumentSemantics::PhysicalPointer, + }, + typ: &ast::Type::Pointer(underlying_type.clone(), state_space), + stmt_ctor: |new_ptr_src| Statement::PtrAdd { + underlying_type, + state_space, + dst, + ptr_src: new_ptr_src, + constant_src, + }, + }; + insert_implicit_conversions_impl( + &mut result, + id_def, + visit_desc, + bitcast_physical_pointer, + Some(state_space), + )?; + } s @ Statement::Conditional(_) | s @ Statement::Label(_) | s @ Statement::Constant(_) @@ -2327,7 +2504,7 @@ fn insert_implicit_conversions_impl( conversion_fn = bitcast_physical_pointer; } ArgumentSemantics::RegisterPointer => { - conversion_fn = bitcast_logical_pointer; + conversion_fn = bitcast_register_pointer; } ArgumentSemantics::Address => { conversion_fn = force_bitcast_ptr_to_bit; @@ -2369,26 +2546,18 @@ fn insert_implicit_conversions_impl( fn get_function_type( builder: &mut dr::Builder, map: &mut TypeWordMap, - method_decl: &ast::MethodDecl, + spirv_input: &[ast::Variable], + spirv_output: &[ast::Variable], ) -> (spirv::Word, spirv::Word) { - match method_decl { - ast::MethodDecl::Func(out_params, _, in_params) => map.get_or_add_fn( - builder, - out_params - .iter() - .map(|p| SpirvType::from(ast::Type::from(p.v_type.clone()))), - in_params - .iter() - .map(|p| SpirvType::from(ast::Type::from(p.v_type.clone()))), - ), - ast::MethodDecl::Kernel { in_args, .. } => map.get_or_add_fn( - builder, - iter::empty(), - in_args - .iter() - .map(|p| SpirvType::from(ast::Type::from(p.v_type.clone()))), - ), - } + map.get_or_add_fn( + builder, + spirv_input + .iter() + .map(|var| SpirvType::from(var.v_type.clone())), + spirv_output + .iter() + .map(|var| SpirvType::from(var.v_type.clone())), + ) } fn emit_function_body_ops( @@ -2416,7 +2585,7 @@ fn emit_function_body_ops( Statement::Call(call) => { let (result_type, result_id) = match &*call.ret_params { [(id, typ)] => ( - map.get_or_add(builder, SpirvType::from(ast::Type::from(typ.clone()))), + map.get_or_add(builder, SpirvType::from(typ.to_func_type())), Some(*id), ), [] => (map.void(), None), @@ -2545,42 +2714,13 @@ fn emit_function_body_ops( } let result_type = map.get_or_add(builder, SpirvType::from(ast::Type::from(data.typ.clone()))); - match data.state_space { - ast::LdStateSpace::Generic - | ast::LdStateSpace::Global - | ast::LdStateSpace::Shared => { - builder.load(result_type, Some(arg.dst), arg.src, None, [])?; - } - ast::LdStateSpace::Param | ast::LdStateSpace::Local => { - let result_type = map.get_or_add( - builder, - SpirvType::from(ast::Type::from(data.typ.clone())), - ); - builder.copy_object(result_type, Some(arg.dst), arg.src)?; - } - _ => todo!(), - } + builder.load(result_type, Some(arg.dst), arg.src, None, [])?; } ast::Instruction::St(data, arg) => { if data.qualifier != ast::LdStQualifier::Weak { todo!() } - if data.state_space == ast::StStateSpace::Param - || data.state_space == ast::StStateSpace::Local - { - let result_type = map.get_or_add( - builder, - SpirvType::from(ast::Type::from(data.typ.clone())), - ); - builder.copy_object(result_type, Some(arg.src1), arg.src2)?; - } else if data.state_space == ast::StStateSpace::Generic - || data.state_space == ast::StStateSpace::Global - || data.state_space == ast::StStateSpace::Shared - { - builder.store(arg.src1, arg.src2, None, &[])?; - } else { - todo!() - } + builder.store(arg.src1, arg.src2, None, &[])?; } // SPIR-V does not support ret as guaranteed-converged ast::Instruction::Ret(_) => builder.ret()?, @@ -2817,6 +2957,46 @@ fn emit_function_body_ops( }; negate_func(builder, result_type, Some(arg.dst), arg.src)?; } + ast::Instruction::Sin { arg, .. } => { + let result_type = map.get_or_add_scalar(builder, ast::ScalarType::F32); + builder.ext_inst( + result_type, + Some(arg.dst), + opencl, + spirv::CLOp::sin as u32, + [arg.src], + )?; + } + ast::Instruction::Cos { arg, .. } => { + let result_type = map.get_or_add_scalar(builder, ast::ScalarType::F32); + builder.ext_inst( + result_type, + Some(arg.dst), + opencl, + spirv::CLOp::cos as u32, + [arg.src], + )?; + } + ast::Instruction::Lg2 { arg, .. } => { + let result_type = map.get_or_add_scalar(builder, ast::ScalarType::F32); + builder.ext_inst( + result_type, + Some(arg.dst), + opencl, + spirv::CLOp::log2 as u32, + [arg.src], + )?; + } + ast::Instruction::Ex2 { arg, .. } => { + let result_type = map.get_or_add_scalar(builder, ast::ScalarType::F32); + builder.ext_inst( + result_type, + Some(arg.dst), + opencl, + spirv::CLOp::exp2 as u32, + [arg.src], + )?; + } }, Statement::LoadVar(arg, typ) => { let type_id = map.get_or_add(builder, SpirvType::from(typ.clone())); @@ -2842,6 +3022,22 @@ fn emit_function_body_ops( let result_type = map.get_or_add(builder, SpirvType::from(t.clone())); builder.undef(result_type, Some(*id)); } + Statement::PtrAdd { + underlying_type, + state_space, + dst, + ptr_src, + constant_src, + } => { + let s64_type = map.get_or_add_scalar(builder, ast::ScalarType::S64); + let ptr_as_s64 = builder.bitcast(s64_type, None, *ptr_src)?; + let added_ptr = builder.i_add(s64_type, None, ptr_as_s64, *constant_src)?; + let result_type = map.get_or_add( + builder, + SpirvType::from(ast::Type::Pointer(underlying_type.clone(), *state_space)), + ); + builder.bitcast(result_type, Some(*dst), added_ptr)?; + } } } Ok(()) @@ -3722,16 +3918,21 @@ fn expand_map_variables<'a, 'b>( ast::Statement::Variable(var) => { let ss = match var.var.v_type { ast::VariableType::Reg(_) => StateSpace::Reg, - ast::VariableType::Local(_) => StateSpace::Local, - ast::VariableType::Param(_) => StateSpace::ParamReg, ast::VariableType::Global(_) => StateSpace::Global, ast::VariableType::Shared(_) => StateSpace::Shared, + ast::VariableType::Param(_) => StateSpace::ParamReg, + ast::VariableType::Local(_) => StateSpace::Local, + }; + let mut var_type = ast::Type::from(var.var.v_type.clone()); + var_type = match var.var.v_type { + ast::VariableType::Reg(_) | ast::VariableType::Shared(_) => var_type, + ast::VariableType::Global(_) => var_type.pointer_to(ast::LdStateSpace::Global)?, + ast::VariableType::Param(_) => var_type.pointer_to(ast::LdStateSpace::Param)?, + ast::VariableType::Local(_) => var_type.pointer_to(ast::LdStateSpace::Local)?, }; match var.count { Some(count) => { - for new_id in - id_defs.add_defs(var.var.name, count, ss, var.var.v_type.clone().into()) - { + for new_id in id_defs.add_defs(var.var.name, count, ss, var_type) { result.push(Statement::Variable(ast::Variable { align: var.var.align, v_type: var.var.v_type.clone(), @@ -3741,8 +3942,7 @@ fn expand_map_variables<'a, 'b>( } } None => { - let new_id = - id_defs.add_def(var.var.name, Some((ss, var.var.v_type.clone().into()))); + let new_id = id_defs.add_def(var.var.name, Some((ss, var_type))); result.push(Statement::Variable(ast::Variable { align: var.var.align, v_type: var.var.v_type.clone(), @@ -3858,11 +4058,14 @@ impl<'a> GlobalStringIdResolver<'a> { fn start_fn<'b>( &'b mut self, header: &'b ast::MethodDecl<'a, &'a str>, - ) -> ( - FnStringIdResolver<'a, 'b>, - GlobalFnDeclResolver<'a, 'b>, - ast::MethodDecl<'a, spirv::Word>, - ) { + ) -> Result< + ( + FnStringIdResolver<'a, 'b>, + GlobalFnDeclResolver<'a, 'b>, + ast::MethodDecl<'a, spirv::Word>, + ), + TranslateError, + > { // In case a function decl was inserted earlier we want to use its id let name_id = self.get_or_add_def(header.name()); let mut fn_resolver = FnStringIdResolver { @@ -3874,18 +4077,13 @@ impl<'a> GlobalStringIdResolver<'a> { type_check: HashMap::new(), }; let new_fn_decl = match header { - ast::MethodDecl::Kernel { + ast::MethodDecl::Kernel { name, in_args } => ast::MethodDecl::Kernel { name, - in_args, - uses_shared_mem, - } => ast::MethodDecl::Kernel { - name, - in_args: expand_kernel_params(&mut fn_resolver, in_args.iter()), - uses_shared_mem: *uses_shared_mem, + in_args: expand_kernel_params(&mut fn_resolver, in_args.iter())?, }, ast::MethodDecl::Func(ret_params, _, params) => { - let ret_params_ids = expand_fn_params(&mut fn_resolver, ret_params.iter()); - let params_ids = expand_fn_params(&mut fn_resolver, params.iter()); + let ret_params_ids = expand_fn_params(&mut fn_resolver, ret_params.iter())?; + let params_ids = expand_fn_params(&mut fn_resolver, params.iter())?; self.fns.insert( name_id, FnDecl { @@ -3896,14 +4094,14 @@ impl<'a> GlobalStringIdResolver<'a> { ast::MethodDecl::Func(ret_params_ids, name_id, params_ids) } }; - ( + Ok(( fn_resolver, GlobalFnDeclResolver { variables: &self.variables, fns: &self.fns, }, new_fn_decl, - ) + )) } } @@ -4080,6 +4278,13 @@ enum Statement { Constant(ConstantDefinition), RetValue(ast::RetData, spirv::Word), Undef(ast::Type, spirv::Word), + PtrAdd { + underlying_type: ast::PointerType, + state_space: ast::LdStateSpace, + dst: spirv::Word, + ptr_src: spirv::Word, + constant_src: spirv::Word, + }, } impl ExpandedStatement { @@ -4141,6 +4346,24 @@ impl ExpandedStatement { let id = f(id); Statement::Undef(typ, id) } + Statement::PtrAdd { + underlying_type, + state_space, + dst, + ptr_src, + constant_src, + } => { + let dst = f(dst); + let ptr_src = f(ptr_src); + let constant_src = f(constant_src); + Statement::PtrAdd { + underlying_type, + state_space, + dst, + ptr_src, + constant_src, + } + } } } } @@ -4175,10 +4398,10 @@ impl> ResolvedCall { let new_id = visitor.id( ArgumentDescriptor { op: id, - is_dst: true, - sema: ArgumentSemantics::Default, + is_dst: !typ.is_param(), + sema: typ.semantics(), }, - Some(&typ.clone().into()), + Some(&typ.to_func_type()), )?; Ok((new_id, typ)) }) @@ -4199,9 +4422,9 @@ impl> ResolvedCall { ArgumentDescriptor { op: id, is_dst: false, - sema: ArgumentSemantics::Default, + sema: typ.semantics(), }, - &typ.clone().into(), + &typ.to_func_type(), )?; Ok((new_id, typ)) }) @@ -4364,6 +4587,7 @@ enum Directive<'input> { struct Function<'input> { pub func_decl: ast::MethodDecl<'input, spirv::Word>, + pub spirv_decl: SpirvMethodDecl<'input>, pub globals: Vec>, pub body: Option>, import_as: Option, @@ -4574,9 +4798,7 @@ impl ast::Instruction { // Call instruction is converted to a call statement early on ast::Instruction::Call(_) => return Err(TranslateError::Unreachable), ast::Instruction::Ld(d, a) => { - let is_param = d.state_space == ast::LdStateSpace::Param - || d.state_space == ast::LdStateSpace::Local; - let new_args = a.map(visitor, &d, is_param)?; + let new_args = a.map(visitor, &d)?; ast::Instruction::Ld(d, new_args) } ast::Instruction::Mov(d, a) => { @@ -4629,9 +4851,7 @@ impl ast::Instruction { ast::Instruction::Shr(t, a.map_shift(visitor, &ast::Type::Scalar(t.into()))?) } ast::Instruction::St(d, a) => { - let is_param = d.state_space == ast::StStateSpace::Param - || d.state_space == ast::StStateSpace::Local; - let new_args = a.map(visitor, &d, is_param)?; + let new_args = a.map(visitor, &d)?; ast::Instruction::St(d, new_args) } ast::Instruction::Bra(d, a) => ast::Instruction::Bra(d, a.map(visitor, None)?), @@ -4693,6 +4913,34 @@ impl ast::Instruction { ast::Instruction::Neg(d, a) => { ast::Instruction::Neg(d, a.map(visitor, &ast::Type::Scalar(d.typ))?) } + ast::Instruction::Sin { flush_to_zero, arg } => { + let typ = ast::Type::Scalar(ast::ScalarType::F32); + ast::Instruction::Sin { + flush_to_zero, + arg: arg.map(visitor, &typ)?, + } + } + ast::Instruction::Cos { flush_to_zero, arg } => { + let typ = ast::Type::Scalar(ast::ScalarType::F32); + ast::Instruction::Cos { + flush_to_zero, + arg: arg.map(visitor, &typ)?, + } + } + ast::Instruction::Lg2 { flush_to_zero, arg } => { + let typ = ast::Type::Scalar(ast::ScalarType::F32); + ast::Instruction::Lg2 { + flush_to_zero, + arg: arg.map(visitor, &typ)?, + } + } + ast::Instruction::Ex2 { flush_to_zero, arg } => { + let typ = ast::Type::Scalar(ast::ScalarType::F32); + ast::Instruction::Ex2 { + flush_to_zero, + arg: arg.map(visitor, &typ)?, + } + } }) } } @@ -4864,6 +5112,24 @@ impl ast::Type { components: vec![*len as u32], state_space: *state_space, }, + ast::Type::Pointer(ast::PointerType::Array(scalar, components), state_space) => { + TypeParts { + kind: TypeKind::PointerArray, + scalar_kind: scalar.kind(), + width: scalar.size_of(), + components: components.clone(), + state_space: *state_space, + } + } + ast::Type::Pointer(ast::PointerType::Pointer(scalar, inner_space), state_space) => { + TypeParts { + kind: TypeKind::PointerPointer, + scalar_kind: scalar.kind(), + width: scalar.size_of(), + components: vec![*inner_space as u32], + state_space: *state_space, + } + } } } @@ -4891,6 +5157,31 @@ impl ast::Type { ), t.state_space, ), + TypeKind::PointerArray => ast::Type::Pointer( + ast::PointerType::Array( + ast::ScalarType::from_parts(t.width, t.scalar_kind), + t.components, + ), + t.state_space, + ), + TypeKind::PointerPointer => ast::Type::Pointer( + ast::PointerType::Pointer( + ast::ScalarType::from_parts(t.width, t.scalar_kind), + unsafe { mem::transmute::<_, ast::LdStateSpace>(t.components[0] as u8) }, + ), + t.state_space, + ), + } + } + + fn size_of(&self) -> usize { + match self { + ast::Type::Scalar(typ) => typ.size_of() as usize, + ast::Type::Vector(typ, len) => (typ.size_of() as usize) * (*len as usize), + ast::Type::Array(typ, len) => len + .iter() + .fold(typ.size_of() as usize, |x, y| (x as usize) * (*y as usize)), + ast::Type::Pointer(_, _) => mem::size_of::(), } } } @@ -4911,6 +5202,8 @@ enum TypeKind { Array, PointerScalar, PointerVector, + PointerArray, + PointerPointer, } impl ast::Instruction { @@ -5002,6 +5295,12 @@ impl ast::Instruction { ast::Instruction::Neg(details, _) => details .flush_to_zero .map(|ftz| (ftz, details.typ.size_of())), + ast::Instruction::Sin { flush_to_zero, .. } + | ast::Instruction::Cos { flush_to_zero, .. } + | ast::Instruction::Lg2 { flush_to_zero, .. } + | ast::Instruction::Ex2 { flush_to_zero, .. } => { + Some((*flush_to_zero, mem::size_of::() as u8)) + } } } } @@ -5139,28 +5438,6 @@ impl<'a> ast::Instruction> { } } -impl ast::VariableParamType { - fn width(&self) -> usize { - match self { - ast::VariableParamType::Scalar(t) => ast::ScalarType::from(*t).size_of() as usize, - ast::VariableParamType::Array(t, len) => { - (ast::ScalarType::from(*t).size_of() as usize) - * (len.iter().fold(1, |x, y| x * (*y)) as usize) - } - ast::VariableParamType::Pointer(_, _) => mem::size_of::(), - } - } -} - -impl ast::KernelArgumentType { - fn width(&self) -> usize { - match self { - ast::KernelArgumentType::Normal(t) => t.width(), - ast::KernelArgumentType::Shared => mem::size_of::(), - } - } -} - impl From for ast::Type { fn from(this: ast::KernelArgumentType) -> Self { match this { @@ -5173,18 +5450,6 @@ impl From for ast::Type { } } -impl ast::KernelArgumentType { - fn to_param(self) -> ast::VariableParamType { - match self { - ast::KernelArgumentType::Normal(p) => p, - ast::KernelArgumentType::Shared => ast::VariableParamType::Pointer( - ast::SizedScalarType::B8, - ast::PointerStateSpace::Shared, - ), - } - } -} - impl ast::Arg1 { fn cast>(self) -> ast::Arg1 { ast::Arg1 { src: self.src } @@ -5303,7 +5568,6 @@ impl ast::Arg2Ld { self, visitor: &mut V, details: &ast::LdDetails, - is_param: bool, ) -> Result, TranslateError> { let dst = visitor.id_or_vector( ArgumentDescriptor { @@ -5313,24 +5577,22 @@ impl ast::Arg2Ld { }, &ast::Type::from(details.typ.clone()), )?; + let is_logical_ptr = details.state_space == ast::LdStateSpace::Param + || details.state_space == ast::LdStateSpace::Local; let src = visitor.operand( ArgumentDescriptor { op: self.src, is_dst: false, - sema: if is_param { + sema: if is_logical_ptr { ArgumentSemantics::RegisterPointer } else { ArgumentSemantics::PhysicalPointer }, }, - &(if is_param { - ast::Type::from(details.typ.clone()) - } else { - ast::Type::Pointer( - ast::PointerType::from(details.typ.clone()), - details.state_space, - ) - }), + &ast::Type::Pointer( + ast::PointerType::from(details.typ.clone()), + details.state_space, + ), )?; Ok(ast::Arg2Ld { dst, src }) } @@ -5350,26 +5612,23 @@ impl ast::Arg2St { self, visitor: &mut V, details: &ast::StData, - is_param: bool, ) -> Result, TranslateError> { + let is_logical_ptr = details.state_space == ast::StStateSpace::Param + || details.state_space == ast::StStateSpace::Local; let src1 = visitor.operand( ArgumentDescriptor { op: self.src1, - is_dst: is_param, - sema: if is_param { + is_dst: false, + sema: if is_logical_ptr { ArgumentSemantics::RegisterPointer } else { ArgumentSemantics::PhysicalPointer }, }, - &(if is_param { - details.typ.clone().into() - } else { - ast::Type::Pointer( - ast::PointerType::from(details.typ.clone()), - details.state_space.to_ld_ss(), - ) - }), + &ast::Type::Pointer( + ast::PointerType::from(details.typ.clone()), + details.state_space.to_ld_ss(), + ), )?; let src2 = visitor.operand_or_vector( ArgumentDescriptor { @@ -6190,7 +6449,7 @@ impl ast::LdStateSpace { ast::LdStateSpace::Global => spirv::StorageClass::CrossWorkgroup, ast::LdStateSpace::Local => spirv::StorageClass::Function, ast::LdStateSpace::Shared => spirv::StorageClass::Workgroup, - ast::LdStateSpace::Param => unreachable!(), + ast::LdStateSpace::Param => spirv::StorageClass::Function, } } } @@ -6264,18 +6523,24 @@ impl ast::AtomSemantics { } } -fn bitcast_logical_pointer( - operand: &ast::Type, - instr: &ast::Type, - _: Option, -) -> Result, TranslateError> { - if instr != operand { - Ok(Some(ConversionKind::Default)) - } else { - Ok(None) +impl ast::FnArgumentType { + fn semantics(&self) -> ArgumentSemantics { + match self { + ast::FnArgumentType::Reg(_) => ArgumentSemantics::Default, + ast::FnArgumentType::Param(_) => ArgumentSemantics::RegisterPointer, + ast::FnArgumentType::Shared => ArgumentSemantics::PhysicalPointer, + } } } +fn bitcast_register_pointer( + operand_type: &ast::Type, + instr_type: &ast::Type, + ss: Option, +) -> Result, TranslateError> { + bitcast_physical_pointer(operand_type, instr_type, ss) +} + fn bitcast_physical_pointer( operand_type: &ast::Type, instr_type: &ast::Type, @@ -6315,13 +6580,15 @@ fn bitcast_physical_pointer( } ast::Type::Scalar(ast::ScalarType::B32) | ast::Type::Scalar(ast::ScalarType::U32) - | ast::Type::Scalar(ast::ScalarType::S32) => { - if let Some(ast::LdStateSpace::Shared) = ss { + | ast::Type::Scalar(ast::ScalarType::S32) => match ss { + Some(ast::LdStateSpace::Shared) + | Some(ast::LdStateSpace::Generic) + | Some(ast::LdStateSpace::Param) + | Some(ast::LdStateSpace::Local) => { Ok(Some(ConversionKind::BitToPtr(ast::LdStateSpace::Shared))) - } else { - Err(TranslateError::MismatchedType) } - } + _ => Err(TranslateError::MismatchedType), + }, ast::Type::Pointer(op_scalar_t, op_space) => { if let ast::Type::Pointer(instr_scalar_t, instr_space) = instr_type { if op_space == instr_space { @@ -6566,21 +6833,68 @@ impl<'a> ast::MethodDecl<'a, &'a str> { } } -impl<'a> ast::MethodDecl<'a, spirv::Word> { - fn visit_args(&self, f: &mut impl FnMut(&ast::FnArgument)) { - match self { - ast::MethodDecl::Func(_, _, params) => params.iter().for_each(f), - ast::MethodDecl::Kernel { in_args, .. } => in_args.iter().for_each(|arg| { - f(&ast::FnArgument { - align: arg.align, - name: arg.name, - v_type: match arg.v_type.clone() { - ast::KernelArgumentType::Normal(typ) => ast::FnArgumentType::Param(typ), - ast::KernelArgumentType::Shared => ast::FnArgumentType::Shared, - }, - array_init: arg.array_init.clone(), - }) - }), +struct SpirvMethodDecl<'input> { + input: Vec>, + output: Vec>, + name: MethodName<'input>, + uses_shared_mem: bool, +} + +impl<'input> SpirvMethodDecl<'input> { + fn new(ast_decl: &ast::MethodDecl<'input, spirv::Word>) -> Self { + let (input, output) = match ast_decl { + ast::MethodDecl::Kernel { in_args, .. } => { + let spirv_input = in_args + .iter() + .map(|var| { + let v_type = match &var.v_type { + ast::KernelArgumentType::Normal(t) => { + ast::FnArgumentType::Param(t.clone()) + } + ast::KernelArgumentType::Shared => ast::FnArgumentType::Shared, + }; + ast::Variable { + name: var.name, + align: var.align, + v_type: v_type.to_kernel_type(), + array_init: var.array_init.clone(), + } + }) + .collect(); + (spirv_input, Vec::new()) + } + ast::MethodDecl::Func(out_args, _, in_args) => { + let (param_output, non_param_output): (Vec<_>, Vec<_>) = + out_args.iter().partition(|var| var.v_type.is_param()); + let spirv_output = non_param_output + .into_iter() + .cloned() + .map(|var| ast::Variable { + name: var.name, + align: var.align, + v_type: var.v_type.to_func_type(), + array_init: var.array_init.clone(), + }) + .collect(); + let spirv_input = param_output + .into_iter() + .cloned() + .chain(in_args.iter().cloned()) + .map(|var| ast::Variable { + name: var.name, + align: var.align, + v_type: var.v_type.to_func_type(), + array_init: var.array_init.clone(), + }) + .collect(); + (spirv_input, spirv_output) + } + }; + SpirvMethodDecl { + input, + output, + name: MethodName::new(ast_decl), + uses_shared_mem: false, } } }