diff options
| author | Andrzej Janik <[email protected]> | 2020-11-05 00:27:46 +0100 |
|---|---|---|
| committer | Andrzej Janik <[email protected]> | 2020-11-05 00:27:46 +0100 |
| commit | 96702d86c96ef2d14795a71af43015a8eacd0a94 (patch) | |
| tree | 7da0d15a396f31393668fe6ad2a8e999b6742fcc | |
| parent | e5a53ed5d30fad3d8ebae6d72ead1564d2b97275 (diff) | |
| download | ZLUDA-96702d86c96ef2d14795a71af43015a8eacd0a94.tar.gz ZLUDA-96702d86c96ef2d14795a71af43015a8eacd0a94.zip | |
Fix issues with .param/.local and implement sin, cos, ex2, lg2
60 files changed, 2721 insertions, 2170 deletions
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<KernelArgument<ID>>, - uses_shared_mem: bool, }, } @@ -293,16 +294,51 @@ pub enum KernelArgumentType { Shared, } -impl From<FnArgumentType> 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<SizedScalarType> for PointerType { @@ -343,6 +380,8 @@ impl TryFrom<PointerType> 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<ID> { pub count: Option<u32>, } +#[derive(Clone)] pub struct Variable<T, ID> { pub align: Option<u32>, pub v_type: T, @@ -543,6 +583,10 @@ pub enum Instruction<P: ArgParams> { Sqrt(SqrtDetails, Arg2<P>), Rsqrt(RsqrtDetails, Arg2<P>), Neg(NegDetails, Arg2<P>), + Sin { flush_to_zero: bool, arg: Arg2<P> }, + Cos { flush_to_zero: bool, arg: Arg2<P> }, + Lg2 { flush_to_zero: bool, arg: Arg2<P> }, + Ex2 { flush_to_zero: bool, arg: Arg2<P> }, } #[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" <name:ExtendedID> <in_args:KernelArguments> => - ast::MethodDecl::Kernel{ name, in_args, uses_shared_mem: false }, + ast::MethodDecl::Kernel{ name, in_args }, ".func" <ret_vals:FnArguments?> <name:ExtendedID> <params:FnArguments> => { ast::MethodDecl::Func(ret_vals.unwrap_or_else(|| Vec::new()), name, params) } @@ -687,6 +695,10 @@ Instruction: ast::Instruction<ast::ParsedArgParams<'input>> = { 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<ast::ParsedArgParams<'input>> = { }, } +// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-sin +InstSin: ast::Instruction<ast::ParsedArgParams<'input>> = { + "sin" ".approx" <ftz:".ftz"?> ".f32" <arg:Arg2> => { + 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<ast::ParsedArgParams<'input>> = { + "cos" ".approx" <ftz:".ftz"?> ".f32" <arg:Arg2> => { + 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<ast::ParsedArgParams<'input>> = { + "lg2" ".approx" <ftz:".ftz"?> ".f32" <arg:Arg2> => { + 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<ast::ParsedArgParams<'input>> = { + "ex2" ".approx" <ftz:".ftz"?> ".f32" <arg:Arg2> => { + 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<T: Debug> {
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<ast::Type> 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<ast::PointerType> 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<Self, TranslateError> {
+ 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<spirv::StorageClass> 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<Item = SpirvType>,
in_params: impl ExactSizeIterator<Item = SpirvType>,
+ mut out_params: impl ExactSizeIterator<Item = SpirvType>,
) -> (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<CallgraphKey<'input>, HashMap<u8, spirv::FPDenormMode>>,
+ denorm_information: &HashMap<MethodName<'input>, HashMap<u8, spirv::FPDenormMode>>,
call_map: &HashMap<&'input str, HashSet<spirv::Word>>,
directives: Vec<Directive>,
kernel_info: &mut HashMap<String, KernelInfo>,
@@ -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<CallgraphKey<'input>, spirv::Word>,
+ directly_called_by: &MultiHashMap<MethodName<'input>, spirv::Word>,
visited: &mut HashSet<spirv::Word>,
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,
- });
- }
- 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.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,
}
- }
+ });
+ 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<ExpandedStatement>,
new_id: &mut impl FnMut() -> spirv::Word,
extern_shared_decls: &HashMap<spirv::Word, ast::SizedScalarType>,
- methods_using_extern_shared: &mut HashSet<CallgraphKey<'a>>,
+ methods_using_extern_shared: &mut HashSet<MethodName<'a>>,
shared_id_param: spirv::Word,
shared_var_id: spirv::Word,
statements: Vec<ExpandedStatement>,
@@ -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<CallgraphKey<'a>>,
- directly_called_by: &MultiHashMap<spirv::Word, CallgraphKey<'a>>,
+ methods_using_extern_shared: &mut HashSet<MethodName<'a>>,
+ directly_called_by: &MultiHashMap<spirv::Word, MethodName<'a>>,
) {
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<CallgraphKey<'a>>,
- directly_called_by: &MultiHashMap<spirv::Word, CallgraphKey<'a>>,
+ methods_using_extern_shared: &mut HashSet<MethodName<'a>>,
+ directly_called_by: &MultiHashMap<spirv::Word, MethodName<'a>>,
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<T: Eq + Hash>( // and emit suitable execution mode
fn compute_denorm_information<'input>(
module: &[Directive<'input>],
-) -> HashMap<CallgraphKey<'input>, HashMap<u8, spirv::FPDenormMode>> {
+) -> HashMap<MethodName<'input>, HashMap<u8, spirv::FPDenormMode>> {
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<ast::VariableType, spirv::Word>],
- func_directive: &ast::MethodDecl<spirv::Word>,
- denorm_information: &HashMap<CallgraphKey<'a>, HashMap<u8, spirv::FPDenormMode>>,
+ func_decl: &SpirvMethodDecl<'a>,
+ _denorm_information: &HashMap<MethodName<'a>, HashMap<u8, spirv::FPDenormMode>>,
call_map: &HashMap<&'a str, HashSet<spirv::Word>>,
direcitves: &[Directive],
kernel_info: &mut HashMap<String, KernelInfo>,
) -> 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<String, Directive>,
f: ast::ParsedFunction<'a>,
) -> Result<Function<'a>, 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<Item = &'b ast::KernelArgument<&'a str>>,
-) -> Vec<ast::KernelArgument<spirv::Word>> {
- 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<Vec<ast::KernelArgument<spirv::Word>>, 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::<Result<_, _>>()
}
fn expand_fn_params<'a, 'b>(
fn_resolver: &mut FnStringIdResolver<'a, 'b>,
args: impl Iterator<Item = &'b ast::FnArgument<&'a str>>,
-) -> Vec<ast::FnArgument<spirv::Word>> {
+) -> Result<Vec<ast::FnArgument<spirv::Word>>, 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<Vec<ast::Statement<ast::ParsedArgParams<'input>>>>,
) -> Result<Function<'input>, 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<TypedStatement>,
id_def: &mut MutableNumericIdResolver,
- mut f_args: ast::MethodDecl<'a, spirv::Word>,
-) -> Result<(ast::MethodDecl<'a, spirv::Word>, Vec<TypedStatement>), TranslateError> {
+ fn_decl: &mut SpirvMethodDecl,
+) -> Result<Vec<TypedStatement>, 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());
- let new_id = id_def.new_id(typ.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: 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,
- src2: new_id,
- },
- typ,
- ));
- p.name = new_id;
}
- None
+ None => return Err(TranslateError::Unreachable),
}
- ast::MethodDecl::Func(out_params, _, in_params) => {
- for p in in_params.iter_mut() {
- let typ = ast::Type::from(p.v_type.clone());
+ }
+ 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());
- 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(),
+ 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.clone(),
+ typ,
));
- 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!(),
+ arg.name = new_id;
}
+ None => {}
}
- };
+ }
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<Option<ast::VariableType>, 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<ExpandedStatement, TranslateError>;
}
+struct VisitArgumentDescriptor<'a, Ctor: FnOnce(spirv::Word) -> ExpandedStatement> {
+ desc: ArgumentDescriptor<spirv::Word>,
+ 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<spirv::Word>,
+ Option<&ast::Type>,
+ ) -> Result<spirv::Word, TranslateError>,
+ >(
+ self,
+ f: &mut F,
+ ) -> Result<ExpandedStatement, TranslateError> {
+ 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<TypedStatement>,
@@ -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<spirv::Word>, 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<spirv::Word, TranslateError> {
- 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::<ExpandedArgParams>::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::<ExpandedArgParams>::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::<ExpandedArgParams>::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::<ExpandedArgParams>::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::Word>,
+ spirv_input: &[ast::Variable<ast::Type, spirv::Word>],
+ spirv_output: &[ast::Variable<ast::Type, spirv::Word>],
) -> (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 {
- name,
- in_args,
- uses_shared_mem,
- } => ast::MethodDecl::Kernel {
+ ast::MethodDecl::Kernel { name, in_args } => 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<I, P: ast::ArgParams> { 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<From: ArgParamsEx<Id = spirv::Word>> ResolvedCall<From> { 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<From: ArgParamsEx<Id = spirv::Word>> ResolvedCall<From> { 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<ast::Variable<ast::VariableType, spirv::Word>>,
pub body: Option<Vec<ExpandedStatement>>,
import_as: Option<String>,
@@ -4574,9 +4798,7 @@ impl<T: ArgParamsEx> ast::Instruction<T> { // 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<T: ArgParamsEx> ast::Instruction<T> { 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<T: ArgParamsEx> ast::Instruction<T> { 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::<usize>(),
}
}
}
@@ -4911,6 +5202,8 @@ enum TypeKind { Array,
PointerScalar,
PointerVector,
+ PointerArray,
+ PointerPointer,
}
impl ast::Instruction<ExpandedArgParams> {
@@ -5002,6 +5295,12 @@ impl ast::Instruction<ExpandedArgParams> { 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::<f32>() as u8))
+ }
}
}
}
@@ -5139,28 +5438,6 @@ impl<'a> ast::Instruction<ast::ParsedArgParams<'a>> { }
}
-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::<usize>(),
- }
- }
-}
-
-impl ast::KernelArgumentType {
- fn width(&self) -> usize {
- match self {
- ast::KernelArgumentType::Normal(t) => t.width(),
- ast::KernelArgumentType::Shared => mem::size_of::<usize>(),
- }
- }
-}
-
impl From<ast::KernelArgumentType> for ast::Type {
fn from(this: ast::KernelArgumentType) -> Self {
match this {
@@ -5173,18 +5450,6 @@ impl From<ast::KernelArgumentType> 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<T: ArgParamsEx> ast::Arg1<T> {
fn cast<U: ArgParamsEx<Id = T::Id>>(self) -> ast::Arg1<U> {
ast::Arg1 { src: self.src }
@@ -5303,7 +5568,6 @@ impl<T: ArgParamsEx> ast::Arg2Ld<T> { self,
visitor: &mut V,
details: &ast::LdDetails,
- is_param: bool,
) -> Result<ast::Arg2Ld<U>, TranslateError> {
let dst = visitor.id_or_vector(
ArgumentDescriptor {
@@ -5313,24 +5577,22 @@ impl<T: ArgParamsEx> ast::Arg2Ld<T> { },
&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<T: ArgParamsEx> ast::Arg2St<T> { self,
visitor: &mut V,
details: &ast::StData,
- is_param: bool,
) -> Result<ast::Arg2St<U>, 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<ast::LdStateSpace>,
-) -> Result<Option<ConversionKind>, 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<ast::LdStateSpace>,
+) -> Result<Option<ConversionKind>, 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<spirv::Word>)) {
- 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<ast::Variable<ast::Type, spirv::Word>>,
+ output: Vec<ast::Variable<ast::Type, spirv::Word>>,
+ 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,
}
}
}
|