Applied-Software/ZLUDA
5
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-07-21 19:26:22 +03:00
Code Issues Packages Projects Releases Wiki Activity

Fix issues with .param/.local and implement sin, cos, ex2, lg2

Browse Source
This commit is contained in:
Andrzej Janik
2020-11-05 00:27:46 +01:00
parent e5a53ed5d3
commit 96702d86c9
60 changed files with 2724 additions and 2173 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

63
ptx/src/ast.rs
View File

@ -113,6 +113,8 @@ sub_type! {
VariableRegType { VariableRegType {
Scalar(ScalarType), Scalar(ScalarType),
Vector(SizedScalarType, u8), 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 // Pointer variant is used when passing around SLM pointer between
// function calls for dynamic SLM // function calls for dynamic SLM
Pointer(SizedScalarType, PointerStateSpace) Pointer(SizedScalarType, PointerStateSpace)
@ -267,7 +269,6 @@ pub enum MethodDecl<'a, ID> {
Kernel { Kernel {
name: &'a str, name: &'a str,
in_args: Vec<KernelArgument<ID>>, in_args: Vec<KernelArgument<ID>>,
uses_shared_mem: bool,
}, },
} }
@ -293,16 +294,51 @@ pub enum KernelArgumentType {
Shared, Shared,
} }
impl From<FnArgumentType> for Type { impl FnArgumentType {
fn from(t: FnArgumentType) -> Self { pub fn to_type(&self, is_kernel: bool) -> Type {
match t { if is_kernel {
FnArgumentType::Reg(x) => x.into(), self.to_kernel_type()
FnArgumentType::Param(x) => x.into(), } 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 => { FnArgumentType::Shared => {
Type::Pointer(PointerType::Scalar(ScalarType::B8), LdStateSpace::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!( sub_enum!(
@ -323,11 +359,12 @@ pub enum Type {
Pointer(PointerType, LdStateSpace), Pointer(PointerType, LdStateSpace),
} }
sub_type! { #[derive(PartialEq, Eq, Clone)]
PointerType { pub enum PointerType {
Scalar(ScalarType), Scalar(ScalarType),
Vector(ScalarType, u8), Vector(ScalarType, u8),
} Array(ScalarType, VecU32),
Pointer(ScalarType, LdStateSpace),
} }
impl From<SizedScalarType> for PointerType { impl From<SizedScalarType> for PointerType {
@ -343,6 +380,8 @@ impl TryFrom<PointerType> for SizedScalarType {
match value { match value {
PointerType::Scalar(t) => Ok(t.try_into()?), PointerType::Scalar(t) => Ok(t.try_into()?),
PointerType::Vector(_, _) => Err(()), PointerType::Vector(_, _) => Err(()),
PointerType::Array(_, _) => Err(()),
PointerType::Pointer(_, _) => Err(()),
} }
} }
} }
@ -456,6 +495,7 @@ pub struct MultiVariable<ID> {
pub count: Option<u32>, pub count: Option<u32>,
} }
#[derive(Clone)]
pub struct Variable<T, ID> { pub struct Variable<T, ID> {
pub align: Option<u32>, pub align: Option<u32>,
pub v_type: T, pub v_type: T,
@ -543,6 +583,10 @@ pub enum Instruction<P: ArgParams> {
Sqrt(SqrtDetails, Arg2<P>), Sqrt(SqrtDetails, Arg2<P>),
Rsqrt(RsqrtDetails, Arg2<P>), Rsqrt(RsqrtDetails, Arg2<P>),
Neg(NegDetails, 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)] #[derive(Copy, Clone)]
@ -744,6 +788,7 @@ pub enum MemScope {
} }
#[derive(Copy, Clone, PartialEq, Eq, Debug)] #[derive(Copy, Clone, PartialEq, Eq, Debug)]
#[repr(u8)]
pub enum LdStateSpace { pub enum LdStateSpace {
Generic, Generic,
Const, Const,

42
ptx/src/ptx.lalrpop
View File

@ -144,12 +144,15 @@ match {
"barrier", "barrier",
"bra", "bra",
"call", "call",
"cos",
"cvt", "cvt",
"cvta", "cvta",
"debug", "debug",
"div", "div",
"ex2",
"fma", "fma",
"ld", "ld",
"lg2",
"mad", "mad",
"map_f64_to_f32", "map_f64_to_f32",
"max", "max",
@ -166,6 +169,7 @@ match {
"setp", "setp",
"shl", "shl",
"shr", "shr",
"sin",
r"sm_[0-9]+" => ShaderModel, r"sm_[0-9]+" => ShaderModel,
"sqrt", "sqrt",
"st", "st",
@ -187,12 +191,15 @@ ExtendedID : &'input str = {
"barrier", "barrier",
"bra", "bra",
"call", "call",
"cos",
"cvt", "cvt",
"cvta", "cvta",
"debug", "debug",
"div", "div",
"ex2",
"fma", "fma",
"ld", "ld",
"lg2",
"mad", "mad",
"map_f64_to_f32", "map_f64_to_f32",
"max", "max",
@ -209,6 +216,7 @@ ExtendedID : &'input str = {
"setp", "setp",
"shl", "shl",
"shr", "shr",
"sin",
ShaderModel, ShaderModel,
"sqrt", "sqrt",
"st", "st",
@ -346,7 +354,7 @@ LinkingDirectives: ast::LinkingDirective = {
MethodDecl: ast::MethodDecl<'input, &'input str> = { MethodDecl: ast::MethodDecl<'input, &'input str> = {
".entry" <name:ExtendedID> <in_args:KernelArguments> => ".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> => { ".func" <ret_vals:FnArguments?> <name:ExtendedID> <params:FnArguments> => {
ast::MethodDecl::Func(ret_vals.unwrap_or_else(|| Vec::new()), name, params) ast::MethodDecl::Func(ret_vals.unwrap_or_else(|| Vec::new()), name, params)
} }
@ -687,6 +695,10 @@ Instruction: ast::Instruction<ast::ParsedArgParams<'input>> = {
InstSqrt, InstSqrt,
InstRsqrt, InstRsqrt,
InstNeg, InstNeg,
InstSin,
InstCos,
InstLg2,
InstEx2,
}; };
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld // 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 = { NegTypeFtz: ast::ScalarType = {
".f16" => ast::ScalarType::F16, ".f16" => ast::ScalarType::F16,
".f16x2" => ast::ScalarType::F16x2, ".f16x2" => ast::ScalarType::F16x2,

43
ptx/src/test/spirv_run/add.spvtxt
View File

@ -2,21 +2,24 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%25 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%23 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "add" OpEntryPoint Kernel %1 "add"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%28 = OpTypeFunction %void %ulong %ulong %26 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_1 = OpConstant %ulong 1 %ulong_1 = OpConstant %ulong 1
%1 = OpFunction %void None %28 %1 = OpFunction %void None %26
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%23 = OpLabel %21 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -25,22 +28,20 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%15 = OpLoad %ulong %4 %19 = OpConvertUToPtr %_ptr_Generic_ulong %13
%21 = OpConvertUToPtr %_ptr_Generic_ulong %15 %12 = OpLoad %ulong %19
%14 = OpLoad %ulong %21 OpStore %6 %12
OpStore %6 %14 %15 = OpLoad %ulong %6
%17 = OpLoad %ulong %6 %14 = OpIAdd %ulong %15 %ulong_1
%16 = OpIAdd %ulong %17 %ulong_1 OpStore %7 %14
OpStore %7 %16 %16 = OpLoad %ulong %5
%18 = OpLoad %ulong %5 %17 = OpLoad %ulong %7
%19 = OpLoad %ulong %7 %20 = OpConvertUToPtr %_ptr_Generic_ulong %16
%22 = OpConvertUToPtr %_ptr_Generic_ulong %18 OpStore %20 %17
OpStore %22 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

124
ptx/src/test/spirv_run/and.spvtxt
View File

@ -1,66 +1,58 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 41 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %31 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "and"
OpCapability Float64 %void = OpTypeVoid
; OpCapability FunctionFloatControlINTEL %ulong = OpTypeInt 64 0
; OpExtension "SPV_INTEL_float_controls2" %34 = OpTypeFunction %void %ulong %ulong
%33 = OpExtInstImport "OpenCL.std" %_ptr_Function_ulong = OpTypePointer Function %ulong
OpMemoryModel Physical64 OpenCL %uint = OpTypeInt 32 0
OpEntryPoint Kernel %1 "and" %_ptr_Function_uint = OpTypePointer Function %uint
%34 = OpTypeVoid %_ptr_Generic_uint = OpTypePointer Generic %uint
%35 = OpTypeInt 64 0 %ulong_4 = OpConstant %ulong 4
%36 = OpTypeFunction %34 %35 %35 %1 = OpFunction %void None %34
%37 = OpTypePointer Function %35 %8 = OpFunctionParameter %ulong
%38 = OpTypeInt 32 0 %9 = OpFunctionParameter %ulong
%39 = OpTypePointer Function %38 %29 = OpLabel
%40 = OpTypePointer Generic %38 %2 = OpVariable %_ptr_Function_ulong Function
%23 = OpConstant %35 4 %3 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %34 None %36 %4 = OpVariable %_ptr_Function_ulong Function
%8 = OpFunctionParameter %35 %5 = OpVariable %_ptr_Function_ulong Function
%9 = OpFunctionParameter %35 %6 = OpVariable %_ptr_Function_uint Function
%31 = OpLabel %7 = OpVariable %_ptr_Function_uint Function
%2 = OpVariable %37 Function OpStore %2 %8
%3 = OpVariable %37 Function OpStore %3 %9
%4 = OpVariable %37 Function %10 = OpLoad %ulong %2
%5 = OpVariable %37 Function OpStore %4 %10
%6 = OpVariable %39 Function %11 = OpLoad %ulong %3
%7 = OpVariable %39 Function OpStore %5 %11
OpStore %2 %8 %13 = OpLoad %ulong %4
OpStore %3 %9 %23 = OpConvertUToPtr %_ptr_Generic_uint %13
%11 = OpLoad %35 %2 %12 = OpLoad %uint %23
%10 = OpCopyObject %35 %11 OpStore %6 %12
OpStore %4 %10 %15 = OpLoad %ulong %4
%13 = OpLoad %35 %3 %22 = OpIAdd %ulong %15 %ulong_4
%12 = OpCopyObject %35 %13 %24 = OpConvertUToPtr %_ptr_Generic_uint %22
OpStore %5 %12 %14 = OpLoad %uint %24
%15 = OpLoad %35 %4 OpStore %7 %14
%25 = OpConvertUToPtr %40 %15 %17 = OpLoad %uint %6
%14 = OpLoad %38 %25 %18 = OpLoad %uint %7
OpStore %6 %14 %26 = OpCopyObject %uint %17
%17 = OpLoad %35 %4 %27 = OpCopyObject %uint %18
%24 = OpIAdd %35 %17 %23 %25 = OpBitwiseAnd %uint %26 %27
%26 = OpConvertUToPtr %40 %24 %16 = OpCopyObject %uint %25
%16 = OpLoad %38 %26 OpStore %6 %16
OpStore %7 %16 %19 = OpLoad %ulong %5
%19 = OpLoad %38 %6 %20 = OpLoad %uint %6
%20 = OpLoad %38 %7 %28 = OpConvertUToPtr %_ptr_Generic_uint %19
%28 = OpCopyObject %38 %19 OpStore %28 %20
%29 = OpCopyObject %38 %20 OpReturn
%27 = OpBitwiseAnd %38 %28 %29 OpFunctionEnd
%18 = OpCopyObject %38 %27
OpStore %6 %18
%21 = OpLoad %35 %5
%22 = OpLoad %38 %6
%30 = OpConvertUToPtr %40 %21
OpStore %30 %22
OpReturn
OpFunctionEnd

160
ptx/src/test/spirv_run/atom_add.spvtxt
View File

@ -1,84 +1,76 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 55 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %38 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "atom_add" %4
OpCapability Float64 OpDecorate %4 Alignment 4
; OpCapability FunctionFloatControlINTEL %void = OpTypeVoid
; OpExtension "SPV_INTEL_float_controls2" %uint = OpTypeInt 32 0
%40 = OpExtInstImport "OpenCL.std" %uchar = OpTypeInt 8 0
OpMemoryModel Physical64 OpenCL %uint_1024 = OpConstant %uint 1024
OpEntryPoint Kernel %1 "atom_add" %4 %_arr_uchar_uint_1024 = OpTypeArray %uchar %uint_1024
OpDecorate %4 Alignment 4 %_ptr_Workgroup__arr_uchar_uint_1024 = OpTypePointer Workgroup %_arr_uchar_uint_1024
%41 = OpTypeVoid %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
%42 = OpTypeInt 32 0 %ulong = OpTypeInt 64 0
%43 = OpTypeInt 8 0 %46 = OpTypeFunction %void %ulong %ulong
%44 = OpConstant %42 1024 %_ptr_Function_ulong = OpTypePointer Function %ulong
%45 = OpTypeArray %43 %44 %_ptr_Function_uint = OpTypePointer Function %uint
%46 = OpTypePointer Workgroup %45 %_ptr_Generic_uint = OpTypePointer Generic %uint
%4 = OpVariable %46 Workgroup %ulong_4 = OpConstant %ulong 4
%47 = OpTypeInt 64 0 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
%48 = OpTypeFunction %41 %47 %47 %uint_1 = OpConstant %uint 1
%49 = OpTypePointer Function %47 %uint_0 = OpConstant %uint 0
%50 = OpTypePointer Function %42 %ulong_4_0 = OpConstant %ulong 4
%51 = OpTypePointer Generic %42 %1 = OpFunction %void None %46
%27 = OpConstant %47 4 %9 = OpFunctionParameter %ulong
%52 = OpTypePointer Workgroup %42 %10 = OpFunctionParameter %ulong
%53 = OpConstant %42 1 %36 = OpLabel
%54 = OpConstant %42 0 %2 = OpVariable %_ptr_Function_ulong Function
%29 = OpConstant %47 4 %3 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %41 None %48 %5 = OpVariable %_ptr_Function_ulong Function
%9 = OpFunctionParameter %47 %6 = OpVariable %_ptr_Function_ulong Function
%10 = OpFunctionParameter %47 %7 = OpVariable %_ptr_Function_uint Function
%38 = OpLabel %8 = OpVariable %_ptr_Function_uint Function
%2 = OpVariable %49 Function OpStore %2 %9
%3 = OpVariable %49 Function OpStore %3 %10
%5 = OpVariable %49 Function %11 = OpLoad %ulong %2
%6 = OpVariable %49 Function OpStore %5 %11
%7 = OpVariable %50 Function %12 = OpLoad %ulong %3
%8 = OpVariable %50 Function OpStore %6 %12
OpStore %2 %9 %14 = OpLoad %ulong %5
OpStore %3 %10 %29 = OpConvertUToPtr %_ptr_Generic_uint %14
%12 = OpLoad %47 %2 %13 = OpLoad %uint %29
%11 = OpCopyObject %47 %12 OpStore %7 %13
OpStore %5 %11 %16 = OpLoad %ulong %5
%14 = OpLoad %47 %3 %26 = OpIAdd %ulong %16 %ulong_4
%13 = OpCopyObject %47 %14 %30 = OpConvertUToPtr %_ptr_Generic_uint %26
OpStore %6 %13 %15 = OpLoad %uint %30
%16 = OpLoad %47 %5 OpStore %8 %15
%31 = OpConvertUToPtr %51 %16 %17 = OpLoad %uint %7
%15 = OpLoad %42 %31 %31 = OpBitcast %_ptr_Workgroup_uint %4
OpStore %7 %15 OpStore %31 %17
%18 = OpLoad %47 %5 %19 = OpLoad %uint %8
%28 = OpIAdd %47 %18 %27 %32 = OpBitcast %_ptr_Workgroup_uint %4
%32 = OpConvertUToPtr %51 %28 %18 = OpAtomicIAdd %uint %32 %uint_1 %uint_0 %19
%17 = OpLoad %42 %32 OpStore %7 %18
OpStore %8 %17 %33 = OpBitcast %_ptr_Workgroup_uint %4
%19 = OpLoad %42 %7 %20 = OpLoad %uint %33
%33 = OpBitcast %52 %4 OpStore %8 %20
OpStore %33 %19 %21 = OpLoad %ulong %6
%21 = OpLoad %42 %8 %22 = OpLoad %uint %7
%34 = OpBitcast %52 %4 %34 = OpConvertUToPtr %_ptr_Generic_uint %21
%20 = OpAtomicIAdd %42 %34 %53 %54 %21 OpStore %34 %22
OpStore %7 %20 %23 = OpLoad %ulong %6
%35 = OpBitcast %52 %4 %24 = OpLoad %uint %8
%22 = OpLoad %42 %35 %28 = OpIAdd %ulong %23 %ulong_4_0
OpStore %8 %22 %35 = OpConvertUToPtr %_ptr_Generic_uint %28
%23 = OpLoad %47 %6 OpStore %35 %24
%24 = OpLoad %42 %7 OpReturn
%36 = OpConvertUToPtr %51 %23 OpFunctionEnd
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

146
ptx/src/test/spirv_run/atom_cas.spvtxt
View File

@ -1,77 +1,69 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 51 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %39 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "atom_cas"
OpCapability Float64 %void = OpTypeVoid
; OpCapability FunctionFloatControlINTEL %ulong = OpTypeInt 64 0
; OpExtension "SPV_INTEL_float_controls2" %42 = OpTypeFunction %void %ulong %ulong
%41 = OpExtInstImport "OpenCL.std" %_ptr_Function_ulong = OpTypePointer Function %ulong
OpMemoryModel Physical64 OpenCL %uint = OpTypeInt 32 0
OpEntryPoint Kernel %1 "atom_cas" %_ptr_Function_uint = OpTypePointer Function %uint
%42 = OpTypeVoid %_ptr_Generic_uint = OpTypePointer Generic %uint
%43 = OpTypeInt 64 0 %ulong_4 = OpConstant %ulong 4
%44 = OpTypeFunction %42 %43 %43 %uint_100 = OpConstant %uint 100
%45 = OpTypePointer Function %43 %uint_1 = OpConstant %uint 1
%46 = OpTypeInt 32 0 %uint_0 = OpConstant %uint 0
%47 = OpTypePointer Function %46 %ulong_4_0 = OpConstant %ulong 4
%48 = OpTypePointer Generic %46 %ulong_4_1 = OpConstant %ulong 4
%25 = OpConstant %43 4 %1 = OpFunction %void None %42
%27 = OpConstant %46 100 %8 = OpFunctionParameter %ulong
%49 = OpConstant %46 1 %9 = OpFunctionParameter %ulong
%50 = OpConstant %46 0 %37 = OpLabel
%28 = OpConstant %43 4 %2 = OpVariable %_ptr_Function_ulong Function
%30 = OpConstant %43 4 %3 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %42 None %44 %4 = OpVariable %_ptr_Function_ulong Function
%8 = OpFunctionParameter %43 %5 = OpVariable %_ptr_Function_ulong Function
%9 = OpFunctionParameter %43 %6 = OpVariable %_ptr_Function_uint Function
%39 = OpLabel %7 = OpVariable %_ptr_Function_uint Function
%2 = OpVariable %45 Function OpStore %2 %8
%3 = OpVariable %45 Function OpStore %3 %9
%4 = OpVariable %45 Function %10 = OpLoad %ulong %2
%5 = OpVariable %45 Function OpStore %4 %10
%6 = OpVariable %47 Function %11 = OpLoad %ulong %3
%7 = OpVariable %47 Function OpStore %5 %11
OpStore %2 %8 %13 = OpLoad %ulong %4
OpStore %3 %9 %30 = OpConvertUToPtr %_ptr_Generic_uint %13
%11 = OpLoad %43 %2 %12 = OpLoad %uint %30
%10 = OpCopyObject %43 %11 OpStore %6 %12
OpStore %4 %10 %15 = OpLoad %ulong %4
%13 = OpLoad %43 %3 %16 = OpLoad %uint %6
%12 = OpCopyObject %43 %13 %24 = OpIAdd %ulong %15 %ulong_4
OpStore %5 %12 %32 = OpConvertUToPtr %_ptr_Generic_uint %24
%15 = OpLoad %43 %4 %33 = OpCopyObject %uint %16
%32 = OpConvertUToPtr %48 %15 %31 = OpAtomicCompareExchange %uint %32 %uint_1 %uint_0 %uint_0 %uint_100 %33
%14 = OpLoad %46 %32 %14 = OpCopyObject %uint %31
OpStore %6 %14 OpStore %6 %14
%17 = OpLoad %43 %4 %18 = OpLoad %ulong %4
%18 = OpLoad %46 %6 %27 = OpIAdd %ulong %18 %ulong_4_0
%26 = OpIAdd %43 %17 %25 %34 = OpConvertUToPtr %_ptr_Generic_uint %27
%34 = OpConvertUToPtr %48 %26 %17 = OpLoad %uint %34
%35 = OpCopyObject %46 %18 OpStore %7 %17
%33 = OpAtomicCompareExchange %46 %34 %49 %50 %50 %27 %35 %19 = OpLoad %ulong %5
%16 = OpCopyObject %46 %33 %20 = OpLoad %uint %6
OpStore %6 %16 %35 = OpConvertUToPtr %_ptr_Generic_uint %19
%20 = OpLoad %43 %4 OpStore %35 %20
%29 = OpIAdd %43 %20 %28 %21 = OpLoad %ulong %5
%36 = OpConvertUToPtr %48 %29 %22 = OpLoad %uint %7
%19 = OpLoad %46 %36 %29 = OpIAdd %ulong %21 %ulong_4_1
OpStore %7 %19 %36 = OpConvertUToPtr %_ptr_Generic_uint %29
%21 = OpLoad %43 %5 OpStore %36 %22
%22 = OpLoad %46 %6 OpReturn
%37 = OpConvertUToPtr %48 %21 OpFunctionEnd
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

46
ptx/src/test/spirv_run/b64tof64.spvtxt
View File

@ -2,23 +2,25 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64 OpCapability Float64
%26 = OpExtInstImport "OpenCL.std" %24 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "b64tof64" OpEntryPoint Kernel %1 "b64tof64"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%29 = OpTypeFunction %void %ulong %ulong %27 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%double = OpTypeFloat 64 %double = OpTypeFloat 64
%_ptr_Function_double = OpTypePointer Function %double %_ptr_Function_double = OpTypePointer Function %double
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%1 = OpFunction %void None %29 %1 = OpFunction %void None %27
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%24 = OpLabel %22 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_double Function %4 = OpVariable %_ptr_Function_double Function
@ -27,24 +29,22 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %18 = OpBitcast %_ptr_Function_double %2
%20 = OpBitcast %double %11 %10 = OpLoad %double %18
%10 = OpCopyObject %double %20
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %6 %11
OpStore %6 %12 %13 = OpLoad %double %4
%15 = OpLoad %double %4 %19 = OpBitcast %ulong %13
%21 = OpBitcast %ulong %15 %12 = OpCopyObject %ulong %19
%14 = OpCopyObject %ulong %21 OpStore %5 %12
OpStore %5 %14 %15 = OpLoad %ulong %5
%17 = OpLoad %ulong %5 %20 = OpConvertUToPtr %_ptr_Generic_ulong %15
%22 = OpConvertUToPtr %_ptr_Generic_ulong %17 %14 = OpLoad %ulong %20
%16 = OpLoad %ulong %22 OpStore %7 %14
OpStore %7 %16 %16 = OpLoad %ulong %6
%18 = OpLoad %ulong %6 %17 = OpLoad %ulong %7
%19 = OpLoad %ulong %7 %21 = OpConvertUToPtr %_ptr_Generic_ulong %16
%23 = OpConvertUToPtr %_ptr_Generic_ulong %18 OpStore %21 %17
OpStore %23 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

49
ptx/src/test/spirv_run/block.spvtxt
View File

@ -2,22 +2,25 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%29 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%27 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "block" OpEntryPoint Kernel %1 "block"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%32 = OpTypeFunction %void %ulong %ulong %30 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_1 = OpConstant %ulong 1 %ulong_1 = OpConstant %ulong 1
%ulong_1_0 = OpConstant %ulong 1 %ulong_1_0 = OpConstant %ulong 1
%1 = OpFunction %void None %32 %1 = OpFunction %void None %30
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong
%27 = OpLabel %25 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -27,25 +30,23 @@
%8 = OpVariable %_ptr_Function_ulong Function %8 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %9 OpStore %2 %9
OpStore %3 %10 OpStore %3 %10
%12 = OpLoad %ulong %2 %11 = OpLoad %ulong %2
%11 = OpCopyObject %ulong %12
OpStore %4 %11 OpStore %4 %11
%14 = OpLoad %ulong %3 %12 = OpLoad %ulong %3
%13 = OpCopyObject %ulong %14 OpStore %5 %12
OpStore %5 %13 %14 = OpLoad %ulong %4
%16 = OpLoad %ulong %4 %23 = OpConvertUToPtr %_ptr_Generic_ulong %14
%25 = OpConvertUToPtr %_ptr_Generic_ulong %16 %13 = OpLoad %ulong %23
%15 = OpLoad %ulong %25 OpStore %6 %13
OpStore %6 %15 %16 = OpLoad %ulong %6
%18 = OpLoad %ulong %6 %15 = OpIAdd %ulong %16 %ulong_1
%17 = OpIAdd %ulong %18 %ulong_1 OpStore %7 %15
OpStore %7 %17 %18 = OpLoad %ulong %8
%20 = OpLoad %ulong %8 %17 = OpIAdd %ulong %18 %ulong_1_0
%19 = OpIAdd %ulong %20 %ulong_1_0 OpStore %8 %17
OpStore %8 %19 %19 = OpLoad %ulong %5
%21 = OpLoad %ulong %5 %20 = OpLoad %ulong %7
%22 = OpLoad %ulong %7 %24 = OpConvertUToPtr %_ptr_Generic_ulong %19
%26 = OpConvertUToPtr %_ptr_Generic_ulong %21 OpStore %24 %20
OpStore %26 %22
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

49
ptx/src/test/spirv_run/bra.spvtxt
View File

@ -2,22 +2,25 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%31 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%29 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "bra" OpEntryPoint Kernel %1 "bra"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%34 = OpTypeFunction %void %ulong %ulong %32 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_1 = OpConstant %ulong 1 %ulong_1 = OpConstant %ulong 1
%ulong_2 = OpConstant %ulong 2 %ulong_2 = OpConstant %ulong 2
%1 = OpFunction %void None %34 %1 = OpFunction %void None %32
%11 = OpFunctionParameter %ulong %11 = OpFunctionParameter %ulong
%12 = OpFunctionParameter %ulong %12 = OpFunctionParameter %ulong
%29 = OpLabel %27 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
@ -26,31 +29,29 @@
%10 = OpVariable %_ptr_Function_ulong Function %10 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %11 OpStore %2 %11
OpStore %3 %12 OpStore %3 %12
%14 = OpLoad %ulong %2 %13 = OpLoad %ulong %2
%13 = OpCopyObject %ulong %14
OpStore %7 %13 OpStore %7 %13
%16 = OpLoad %ulong %3 %14 = OpLoad %ulong %3
%15 = OpCopyObject %ulong %16 OpStore %8 %14
OpStore %8 %15 %16 = OpLoad %ulong %7
%18 = OpLoad %ulong %7 %25 = OpConvertUToPtr %_ptr_Generic_ulong %16
%27 = OpConvertUToPtr %_ptr_Generic_ulong %18 %15 = OpLoad %ulong %25
%17 = OpLoad %ulong %27 OpStore %9 %15
OpStore %9 %17
OpBranch %4 OpBranch %4
%4 = OpLabel %4 = OpLabel
%18 = OpLoad %ulong %9
%17 = OpIAdd %ulong %18 %ulong_1
OpStore %10 %17
OpBranch %6
%35 = OpLabel
%20 = OpLoad %ulong %9 %20 = OpLoad %ulong %9
%19 = OpIAdd %ulong %20 %ulong_1 %19 = OpIAdd %ulong %20 %ulong_2
OpStore %10 %19 OpStore %10 %19
OpBranch %6 OpBranch %6
%37 = OpLabel
%22 = OpLoad %ulong %9
%21 = OpIAdd %ulong %22 %ulong_2
OpStore %10 %21
OpBranch %6
%6 = OpLabel %6 = OpLabel
%23 = OpLoad %ulong %8 %21 = OpLoad %ulong %8
%24 = OpLoad %ulong %10 %22 = OpLoad %ulong %10
%28 = OpConvertUToPtr %_ptr_Generic_ulong %23 %26 = OpConvertUToPtr %_ptr_Generic_ulong %21
OpStore %28 %24 OpStore %26 %22
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

86
ptx/src/test/spirv_run/call.spvtxt
View File

@ -2,22 +2,25 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%47 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%37 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %4 "call" OpEntryPoint Kernel %4 "call"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%50 = OpTypeFunction %void %ulong %ulong %40 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %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 %ulong_1 = OpConstant %ulong 1
%4 = OpFunction %void None %50 %4 = OpFunction %void None %40
%12 = OpFunctionParameter %ulong %12 = OpFunctionParameter %ulong
%13 = OpFunctionParameter %ulong %13 = OpFunctionParameter %ulong
%32 = OpLabel %26 = OpLabel
%5 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
%6 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
@ -27,49 +30,38 @@
%11 = OpVariable %_ptr_Function_ulong Function %11 = OpVariable %_ptr_Function_ulong Function
OpStore %5 %12 OpStore %5 %12
OpStore %6 %13 OpStore %6 %13
%15 = OpLoad %ulong %5 %14 = OpLoad %ulong %5
%14 = OpCopyObject %ulong %15
OpStore %7 %14 OpStore %7 %14
%17 = OpLoad %ulong %6 %15 = OpLoad %ulong %6
%16 = OpCopyObject %ulong %17 OpStore %8 %15
OpStore %8 %16 %17 = OpLoad %ulong %7
%19 = OpLoad %ulong %7 %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %17
%28 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %19 %16 = OpLoad %ulong %22
%18 = OpLoad %ulong %28 OpStore %9 %16
OpStore %9 %18 %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 %21 = OpLoad %ulong %9
%29 = OpCopyObject %ulong %21 %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %20
%30 = OpCopyObject %ulong %29 OpStore %25 %21
%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
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%1 = OpFunction %ulong None %53 %1 = OpFunction %void None %44
%36 = OpFunctionParameter %ulong %27 = OpFunctionParameter %_ptr_Function_ulong
%45 = OpLabel %28 = OpFunctionParameter %_ptr_Function_ulong
%34 = OpVariable %_ptr_Function_ulong Function %35 = OpLabel
%33 = OpVariable %_ptr_Function_ulong Function %29 = OpVariable %_ptr_Function_ulong Function
%35 = OpVariable %_ptr_Function_ulong Function %30 = OpLoad %ulong %28
OpStore %34 %36 OpStore %29 %30
%38 = OpLoad %ulong %34 %32 = OpLoad %ulong %29
%37 = OpCopyObject %ulong %38 %31 = OpIAdd %ulong %32 %ulong_1
OpStore %35 %37 OpStore %29 %31
%40 = OpLoad %ulong %35 %33 = OpLoad %ulong %29
%39 = OpIAdd %ulong %40 %ulong_1 OpStore %27 %33
OpStore %35 %39 OpReturn
%42 = OpLoad %ulong %35
%41 = OpCopyObject %ulong %42
OpStore %33 %41
%43 = OpLoad %ulong %33
OpReturnValue %43
OpFunctionEnd OpFunctionEnd

105
ptx/src/test/spirv_run/constant_f32.spvtxt
View File

@ -1,57 +1,48 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 32 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %22 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "constant_f32"
OpCapability Float64 %void = OpTypeVoid
; OpCapability FunctionFloatControlINTEL %ulong = OpTypeInt 64 0
; OpExtension "SPV_INTEL_float_controls2" %25 = OpTypeFunction %void %ulong %ulong
%24 = OpExtInstImport "OpenCL.std" %_ptr_Function_ulong = OpTypePointer Function %ulong
OpMemoryModel Physical64 OpenCL %float = OpTypeFloat 32
OpEntryPoint Kernel %1 "constant_f32" %_ptr_Function_float = OpTypePointer Function %float
; OpDecorate %1 FunctionDenormModeINTEL 32 Preserve %_ptr_Generic_float = OpTypePointer Generic %float
%25 = OpTypeVoid %float_0_5 = OpConstant %float 0.5
%26 = OpTypeInt 64 0 %1 = OpFunction %void None %25
%27 = OpTypeFunction %25 %26 %26 %7 = OpFunctionParameter %ulong
%28 = OpTypePointer Function %26 %8 = OpFunctionParameter %ulong
%29 = OpTypeFloat 32 %20 = OpLabel
%30 = OpTypePointer Function %29 %2 = OpVariable %_ptr_Function_ulong Function
%31 = OpTypePointer Generic %29 %3 = OpVariable %_ptr_Function_ulong Function
%19 = OpConstant %29 0.5 %4 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %25 None %27 %5 = OpVariable %_ptr_Function_ulong Function
%7 = OpFunctionParameter %26 %6 = OpVariable %_ptr_Function_float Function
%8 = OpFunctionParameter %26 OpStore %2 %7
%22 = OpLabel OpStore %3 %8
%2 = OpVariable %28 Function %9 = OpLoad %ulong %2
%3 = OpVariable %28 Function OpStore %4 %9
%4 = OpVariable %28 Function %10 = OpLoad %ulong %3
%5 = OpVariable %28 Function OpStore %5 %10
%6 = OpVariable %30 Function %12 = OpLoad %ulong %4
OpStore %2 %7 %18 = OpConvertUToPtr %_ptr_Generic_float %12
OpStore %3 %8 %11 = OpLoad %float %18
%10 = OpLoad %26 %2 OpStore %6 %11
%9 = OpCopyObject %26 %10 %14 = OpLoad %float %6
OpStore %4 %9 %13 = OpFMul %float %14 %float_0_5
%12 = OpLoad %26 %3 OpStore %6 %13
%11 = OpCopyObject %26 %12 %15 = OpLoad %ulong %5
OpStore %5 %11 %16 = OpLoad %float %6
%14 = OpLoad %26 %4 %19 = OpConvertUToPtr %_ptr_Generic_float %15
%20 = OpConvertUToPtr %31 %14 OpStore %19 %16
%13 = OpLoad %29 %20 OpReturn
OpStore %6 %13 OpFunctionEnd
%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

104
ptx/src/test/spirv_run/constant_negative.spvtxt
View File

@ -1,56 +1,48 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 32 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %22 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "constant_negative"
OpCapability Float64 %void = OpTypeVoid
; OpCapability FunctionFloatControlINTEL %ulong = OpTypeInt 64 0
; OpExtension "SPV_INTEL_float_controls2" %25 = OpTypeFunction %void %ulong %ulong
%24 = OpExtInstImport "OpenCL.std" %_ptr_Function_ulong = OpTypePointer Function %ulong
OpMemoryModel Physical64 OpenCL %uint = OpTypeInt 32 0
OpEntryPoint Kernel %1 "constant_negative" %_ptr_Function_uint = OpTypePointer Function %uint
%25 = OpTypeVoid %_ptr_Generic_uint = OpTypePointer Generic %uint
%26 = OpTypeInt 64 0 %uint_4294967295 = OpConstant %uint 4294967295
%27 = OpTypeFunction %25 %26 %26 %1 = OpFunction %void None %25
%28 = OpTypePointer Function %26 %7 = OpFunctionParameter %ulong
%29 = OpTypeInt 32 0 %8 = OpFunctionParameter %ulong
%30 = OpTypePointer Function %29 %20 = OpLabel
%31 = OpTypePointer Generic %29 %2 = OpVariable %_ptr_Function_ulong Function
%19 = OpConstant %29 4294967295 %3 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %25 None %27 %4 = OpVariable %_ptr_Function_ulong Function
%7 = OpFunctionParameter %26 %5 = OpVariable %_ptr_Function_ulong Function
%8 = OpFunctionParameter %26 %6 = OpVariable %_ptr_Function_uint Function
%22 = OpLabel OpStore %2 %7
%2 = OpVariable %28 Function OpStore %3 %8
%3 = OpVariable %28 Function %9 = OpLoad %ulong %2
%4 = OpVariable %28 Function OpStore %4 %9
%5 = OpVariable %28 Function %10 = OpLoad %ulong %3
%6 = OpVariable %30 Function OpStore %5 %10
OpStore %2 %7 %12 = OpLoad %ulong %4
OpStore %3 %8 %18 = OpConvertUToPtr %_ptr_Generic_uint %12
%10 = OpLoad %26 %2 %11 = OpLoad %uint %18
%9 = OpCopyObject %26 %10 OpStore %6 %11
OpStore %4 %9 %14 = OpLoad %uint %6
%12 = OpLoad %26 %3 %13 = OpIMul %uint %14 %uint_4294967295
%11 = OpCopyObject %26 %12 OpStore %6 %13
OpStore %5 %11 %15 = OpLoad %ulong %5
%14 = OpLoad %26 %4 %16 = OpLoad %uint %6
%20 = OpConvertUToPtr %31 %14 %19 = OpConvertUToPtr %_ptr_Generic_uint %15
%13 = OpLoad %29 %20 OpStore %19 %16
OpStore %6 %13 OpReturn
%16 = OpLoad %29 %6 OpFunctionEnd
%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

21
ptx/src/test/spirv_run/cos.ptx Normal file
View File

@ -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;
}

47
ptx/src/test/spirv_run/cos.spvtxt Normal file
View File

@ -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

49
ptx/src/test/spirv_run/cvt_sat_s_u.spvtxt
View File

@ -2,22 +2,25 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%27 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%25 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "cvt_sat_s_u" OpEntryPoint Kernel %1 "cvt_sat_s_u"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%30 = OpTypeFunction %void %ulong %ulong %28 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Generic_uint = OpTypePointer Generic %uint %_ptr_Generic_uint = OpTypePointer Generic %uint
%1 = OpFunction %void None %30 %1 = OpFunction %void None %28
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong
%25 = OpLabel %23 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -27,25 +30,23 @@
%8 = OpVariable %_ptr_Function_uint Function %8 = OpVariable %_ptr_Function_uint Function
OpStore %2 %9 OpStore %2 %9
OpStore %3 %10 OpStore %3 %10
%12 = OpLoad %ulong %2 %11 = OpLoad %ulong %2
%11 = OpCopyObject %ulong %12
OpStore %4 %11 OpStore %4 %11
%14 = OpLoad %ulong %3 %12 = OpLoad %ulong %3
%13 = OpCopyObject %ulong %14 OpStore %5 %12
OpStore %5 %13 %14 = OpLoad %ulong %4
%16 = OpLoad %ulong %4 %21 = OpConvertUToPtr %_ptr_Generic_uint %14
%23 = OpConvertUToPtr %_ptr_Generic_uint %16 %13 = OpLoad %uint %21
%15 = OpLoad %uint %23 OpStore %6 %13
OpStore %6 %15 %16 = OpLoad %uint %6
%18 = OpLoad %uint %6 %15 = OpSatConvertSToU %uint %16
%17 = OpSatConvertSToU %uint %18 OpStore %7 %15
OpStore %7 %17 %18 = OpLoad %uint %7
%20 = OpLoad %uint %7 %17 = OpBitcast %uint %18
%19 = OpBitcast %uint %20 OpStore %8 %17
OpStore %8 %19 %19 = OpLoad %ulong %5
%21 = OpLoad %ulong %5 %20 = OpLoad %uint %8
%22 = OpLoad %uint %8 %22 = OpConvertUToPtr %_ptr_Generic_uint %19
%24 = OpConvertUToPtr %_ptr_Generic_uint %21 OpStore %22 %20
OpStore %24 %22
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

51
ptx/src/test/spirv_run/cvta.spvtxt
View File

@ -2,22 +2,25 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%29 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%27 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "cvta" OpEntryPoint Kernel %1 "cvta"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%32 = OpTypeFunction %void %ulong %ulong %30 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%float = OpTypeFloat 32 %float = OpTypeFloat 32
%_ptr_Function_float = OpTypePointer Function %float %_ptr_Function_float = OpTypePointer Function %float
%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float %_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
%1 = OpFunction %void None %32 %1 = OpFunction %void None %30
%7 = OpFunctionParameter %ulong %7 = OpFunctionParameter %ulong
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%27 = OpLabel %25 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -25,29 +28,27 @@
%6 = OpVariable %_ptr_Function_float Function %6 = OpVariable %_ptr_Function_float Function
OpStore %2 %7 OpStore %2 %7
OpStore %3 %8 OpStore %3 %8
%10 = OpLoad %ulong %2 %9 = OpLoad %ulong %2
%9 = OpCopyObject %ulong %10
OpStore %4 %9 OpStore %4 %9
%12 = OpLoad %ulong %3 %10 = OpLoad %ulong %3
%11 = OpCopyObject %ulong %12 OpStore %5 %10
OpStore %5 %11 %12 = OpLoad %ulong %4
%14 = 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 %22 = OpCopyObject %ulong %14
%21 = OpCopyObject %ulong %22 %21 = OpCopyObject %ulong %22
%13 = OpCopyObject %ulong %21 %13 = OpCopyObject %ulong %21
OpStore %4 %13 OpStore %5 %13
%16 = OpLoad %ulong %5 %16 = OpLoad %ulong %4
%24 = OpCopyObject %ulong %16 %23 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %16
%23 = OpCopyObject %ulong %24 %15 = OpLoad %float %23
%15 = OpCopyObject %ulong %23 OpStore %6 %15
OpStore %5 %15 %17 = OpLoad %ulong %5
%18 = OpLoad %ulong %4 %18 = OpLoad %float %6
%25 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %18 %24 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %17
%17 = OpLoad %float %25 OpStore %24 %18
OpStore %6 %17
%19 = OpLoad %ulong %5
%20 = OpLoad %float %6
%26 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %19
OpStore %26 %20
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

121
ptx/src/test/spirv_run/div_approx.spvtxt
View File

@ -1,65 +1,56 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 38 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %28 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "div_approx"
OpCapability Float64 OpDecorate %16 FPFastMathMode AllowRecip
; OpCapability FunctionFloatControlINTEL %void = OpTypeVoid
; OpExtension "SPV_INTEL_float_controls2" %ulong = OpTypeInt 64 0
%30 = OpExtInstImport "OpenCL.std" %31 = OpTypeFunction %void %ulong %ulong
OpMemoryModel Physical64 OpenCL %_ptr_Function_ulong = OpTypePointer Function %ulong
OpEntryPoint Kernel %1 "div_approx" %float = OpTypeFloat 32
OpDecorate %1 FunctionDenormModeINTEL 32 Preserve %_ptr_Function_float = OpTypePointer Function %float
OpDecorate %18 FPFastMathMode AllowRecip %_ptr_Generic_float = OpTypePointer Generic %float
%31 = OpTypeVoid %ulong_4 = OpConstant %ulong 4
%32 = OpTypeInt 64 0 %1 = OpFunction %void None %31
%33 = OpTypeFunction %31 %32 %32 %8 = OpFunctionParameter %ulong
%34 = OpTypePointer Function %32 %9 = OpFunctionParameter %ulong
%35 = OpTypeFloat 32 %26 = OpLabel
%36 = OpTypePointer Function %35 %2 = OpVariable %_ptr_Function_ulong Function
%37 = OpTypePointer Generic %35 %3 = OpVariable %_ptr_Function_ulong Function
%23 = OpConstant %32 4 %4 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %31 None %33 %5 = OpVariable %_ptr_Function_ulong Function
%8 = OpFunctionParameter %32 %6 = OpVariable %_ptr_Function_float Function
%9 = OpFunctionParameter %32 %7 = OpVariable %_ptr_Function_float Function
%28 = OpLabel OpStore %2 %8
%2 = OpVariable %34 Function OpStore %3 %9
%3 = OpVariable %34 Function %10 = OpLoad %ulong %2
%4 = OpVariable %34 Function OpStore %4 %10
%5 = OpVariable %34 Function %11 = OpLoad %ulong %3
%6 = OpVariable %36 Function OpStore %5 %11
%7 = OpVariable %36 Function %13 = OpLoad %ulong %4
OpStore %2 %8 %23 = OpConvertUToPtr %_ptr_Generic_float %13
OpStore %3 %9 %12 = OpLoad %float %23
%11 = OpLoad %32 %2 OpStore %6 %12
%10 = OpCopyObject %32 %11 %15 = OpLoad %ulong %4
OpStore %4 %10 %22 = OpIAdd %ulong %15 %ulong_4
%13 = OpLoad %32 %3 %24 = OpConvertUToPtr %_ptr_Generic_float %22
%12 = OpCopyObject %32 %13 %14 = OpLoad %float %24
OpStore %5 %12 OpStore %7 %14
%15 = OpLoad %32 %4 %17 = OpLoad %float %6
%25 = OpConvertUToPtr %37 %15 %18 = OpLoad %float %7
%14 = OpLoad %35 %25 %16 = OpFDiv %float %17 %18
OpStore %6 %14 OpStore %6 %16
%17 = OpLoad %32 %4 %19 = OpLoad %ulong %5
%24 = OpIAdd %32 %17 %23 %20 = OpLoad %float %6
%26 = OpConvertUToPtr %37 %24 %25 = OpConvertUToPtr %_ptr_Generic_float %19
%16 = OpLoad %35 %26 OpStore %25 %20
OpStore %7 %16 OpReturn
%19 = OpLoad %35 %6 OpFunctionEnd
%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

21
ptx/src/test/spirv_run/ex2.ptx Normal file
View File

@ -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;
}

47
ptx/src/test/spirv_run/ex2.spvtxt Normal file
View File

@ -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

62
ptx/src/test/spirv_run/extern_shared.spvtxt
View File

@ -7,7 +7,7 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%32 = OpExtInstImport "OpenCL.std" %30 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "extern_shared" %1 OpEntryPoint Kernel %2 "extern_shared" %1
%void = OpTypeVoid %void = OpTypeVoid
@ -18,51 +18,49 @@
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%uchar = OpTypeInt 8 0 %uchar = OpTypeInt 8 0
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar %_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__ptr_Workgroup_uchar = OpTypePointer Function %_ptr_Workgroup_uchar
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%_ptr_Function__ptr_Workgroup_uint = OpTypePointer Function %_ptr_Workgroup_uint %_ptr_Function__ptr_Workgroup_uint = OpTypePointer Function %_ptr_Workgroup_uint
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%2 = OpFunction %void None %40 %2 = OpFunction %void None %38
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%28 = OpFunctionParameter %_ptr_Workgroup_uchar %26 = OpFunctionParameter %_ptr_Workgroup_uchar
%41 = OpLabel %39 = OpLabel
%29 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function %27 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%5 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
%6 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %29 %28 OpStore %27 %26
OpBranch %26 OpBranch %24
%26 = OpLabel %24 = OpLabel
OpStore %3 %8 OpStore %3 %8
OpStore %4 %9 OpStore %4 %9
%11 = OpLoad %ulong %3 %10 = OpLoad %ulong %3
%10 = OpCopyObject %ulong %11
OpStore %5 %10 OpStore %5 %10
%13 = OpLoad %ulong %4 %11 = OpLoad %ulong %4
%12 = OpCopyObject %ulong %13 OpStore %6 %11
OpStore %6 %12 %13 = OpLoad %ulong %5
%15 = OpLoad %ulong %5 %20 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %13
%22 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %15 %12 = OpLoad %ulong %20
%14 = OpLoad %ulong %22 OpStore %7 %12
OpStore %7 %14 %28 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %27
%30 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %29 %14 = OpLoad %_ptr_Workgroup_uint %28
%16 = OpLoad %_ptr_Workgroup_uint %30 %15 = OpLoad %ulong %7
%17 = OpLoad %ulong %7 %21 = OpBitcast %_ptr_Workgroup_ulong %14
%23 = OpBitcast %_ptr_Workgroup_ulong %16 OpStore %21 %15
OpStore %23 %17 %29 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %27
%31 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %29 %17 = OpLoad %_ptr_Workgroup_uint %29
%19 = OpLoad %_ptr_Workgroup_uint %31 %22 = OpBitcast %_ptr_Workgroup_ulong %17
%24 = OpBitcast %_ptr_Workgroup_ulong %19 %16 = OpLoad %ulong %22
%18 = OpLoad %ulong %24 OpStore %7 %16
OpStore %7 %18 %18 = OpLoad %ulong %6
%20 = OpLoad %ulong %6 %19 = OpLoad %ulong %7
%21 = OpLoad %ulong %7 %23 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %18
%25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %20 OpStore %23 %19
OpStore %25 %21
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

135
ptx/src/test/spirv_run/fma.spvtxt
View File

@ -1,72 +1,63 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 45 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %35 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "fma"
OpCapability Float64 %void = OpTypeVoid
; OpCapability FunctionFloatControlINTEL %ulong = OpTypeInt 64 0
; OpExtension "SPV_INTEL_float_controls2" %38 = OpTypeFunction %void %ulong %ulong
%37 = OpExtInstImport "OpenCL.std" %_ptr_Function_ulong = OpTypePointer Function %ulong
OpMemoryModel Physical64 OpenCL %float = OpTypeFloat 32
OpEntryPoint Kernel %1 "fma" %_ptr_Function_float = OpTypePointer Function %float
; OpDecorate %1 FunctionDenormModeINTEL 32 Preserve %_ptr_Generic_float = OpTypePointer Generic %float
%38 = OpTypeVoid %ulong_4 = OpConstant %ulong 4
%39 = OpTypeInt 64 0 %ulong_8 = OpConstant %ulong 8
%40 = OpTypeFunction %38 %39 %39 %1 = OpFunction %void None %38
%41 = OpTypePointer Function %39 %9 = OpFunctionParameter %ulong
%42 = OpTypeFloat 32 %10 = OpFunctionParameter %ulong
%43 = OpTypePointer Function %42 %33 = OpLabel
%44 = OpTypePointer Generic %42 %2 = OpVariable %_ptr_Function_ulong Function
%27 = OpConstant %39 4 %3 = OpVariable %_ptr_Function_ulong Function
%29 = OpConstant %39 8 %4 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %38 None %40 %5 = OpVariable %_ptr_Function_ulong Function
%9 = OpFunctionParameter %39 %6 = OpVariable %_ptr_Function_float Function
%10 = OpFunctionParameter %39 %7 = OpVariable %_ptr_Function_float Function
%35 = OpLabel %8 = OpVariable %_ptr_Function_float Function
%2 = OpVariable %41 Function OpStore %2 %9
%3 = OpVariable %41 Function OpStore %3 %10
%4 = OpVariable %41 Function %11 = OpLoad %ulong %2
%5 = OpVariable %41 Function OpStore %4 %11
%6 = OpVariable %43 Function %12 = OpLoad %ulong %3
%7 = OpVariable %43 Function OpStore %5 %12
%8 = OpVariable %43 Function %14 = OpLoad %ulong %4
OpStore %2 %9 %29 = OpConvertUToPtr %_ptr_Generic_float %14
OpStore %3 %10 %13 = OpLoad %float %29
%12 = OpLoad %39 %2 OpStore %6 %13
%11 = OpCopyObject %39 %12 %16 = OpLoad %ulong %4
OpStore %4 %11 %26 = OpIAdd %ulong %16 %ulong_4
%14 = OpLoad %39 %3 %30 = OpConvertUToPtr %_ptr_Generic_float %26
%13 = OpCopyObject %39 %14 %15 = OpLoad %float %30
OpStore %5 %13 OpStore %7 %15
%16 = OpLoad %39 %4 %18 = OpLoad %ulong %4
%31 = OpConvertUToPtr %44 %16 %28 = OpIAdd %ulong %18 %ulong_8
%15 = OpLoad %42 %31 %31 = OpConvertUToPtr %_ptr_Generic_float %28
OpStore %6 %15 %17 = OpLoad %float %31
%18 = OpLoad %39 %4 OpStore %8 %17
%28 = OpIAdd %39 %18 %27 %20 = OpLoad %float %6
%32 = OpConvertUToPtr %44 %28 %21 = OpLoad %float %7
%17 = OpLoad %42 %32 %22 = OpLoad %float %8
OpStore %7 %17 %19 = OpExtInst %float %35 mad %20 %21 %22
%20 = OpLoad %39 %4 OpStore %6 %19
%30 = OpIAdd %39 %20 %29 %23 = OpLoad %ulong %5
%33 = OpConvertUToPtr %44 %30 %24 = OpLoad %float %6
%19 = OpLoad %42 %33 %32 = OpConvertUToPtr %_ptr_Generic_float %23
OpStore %8 %19 OpStore %32 %24
%22 = OpLoad %42 %6 OpReturn
%23 = OpLoad %42 %7 OpFunctionEnd
%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

37
ptx/src/test/spirv_run/global_array.spvtxt
View File

@ -7,28 +7,28 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%22 = OpExtInstImport "OpenCL.std" %21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "global_array" %1 OpEntryPoint Kernel %2 "global_array" %1
%void = OpTypeVoid %void = OpTypeVoid
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4 %uint_4 = OpConstant %uint 4
%_arr_uint_uint_4 = OpTypeArray %uint %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_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%31 = OpConstantComposite %_arr_uint_uint_4 %uint_1 %uint_0 %uint_0 %uint_0 %28 = OpConstantComposite %_arr_uint_uint_4 %uint_1 %uint_0 %uint_0 %uint_0
%1 = OpVariable %_ptr_CrossWorkgroup__arr_uint_uint_4 CrossWorkgroup %31 %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 %ulong = OpTypeInt 64 0
%33 = OpTypeFunction %void %ulong %ulong %32 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint %_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%2 = OpFunction %void None %33 %2 = OpFunction %void None %32
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%20 = OpLabel %19 = OpLabel
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%5 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
@ -36,19 +36,18 @@
%7 = OpVariable %_ptr_Function_uint Function %7 = OpVariable %_ptr_Function_uint Function
OpStore %3 %8 OpStore %3 %8
OpStore %4 %9 OpStore %4 %9
%17 = OpConvertPtrToU %ulong %1 %16 = OpConvertPtrToU %ulong %1
%10 = OpCopyObject %ulong %17 %10 = OpCopyObject %ulong %16
OpStore %5 %10 OpStore %5 %10
%12 = OpLoad %ulong %4 %11 = OpLoad %ulong %4
%11 = OpCopyObject %ulong %12
OpStore %6 %11 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 %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %14
%13 = OpLoad %uint %18 OpStore %18 %15
OpStore %7 %13
%15 = OpLoad %ulong %6
%16 = OpLoad %uint %7
%19 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %15
OpStore %19 %16
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

48
ptx/src/test/spirv_run/implicit_param.spvtxt
View File

@ -2,25 +2,27 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64 OpCapability Float64
%28 = OpExtInstImport "OpenCL.std" %24 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "implicit_param" OpEntryPoint Kernel %1 "implicit_param"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%31 = OpTypeFunction %void %ulong %ulong %27 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%float = OpTypeFloat 32 %float = OpTypeFloat 32
%_ptr_Function_float = OpTypePointer Function %float %_ptr_Function_float = OpTypePointer Function %float
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float %_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
%1 = OpFunction %void None %31 %1 = OpFunction %void None %27
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%26 = OpLabel %22 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -29,27 +31,23 @@
%7 = OpVariable %_ptr_Function_uint Function %7 = OpVariable %_ptr_Function_uint Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%15 = OpLoad %ulong %4 %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %13
%22 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %15 %12 = OpLoad %float %18
%14 = OpLoad %float %22 OpStore %6 %12
OpStore %6 %14 %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 %17 = OpLoad %float %6
%23 = OpCopyObject %float %17 %21 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %16
%16 = OpBitcast %uint %23 OpStore %21 %17
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
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

37
ptx/src/test/spirv_run/ld_st.spvtxt
View File

@ -2,20 +2,23 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%21 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%19 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "ld_st" OpEntryPoint Kernel %1 "ld_st"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%24 = OpTypeFunction %void %ulong %ulong %22 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%1 = OpFunction %void None %24 %1 = OpFunction %void None %22
%7 = OpFunctionParameter %ulong %7 = OpFunctionParameter %ulong
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%19 = OpLabel %17 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -23,19 +26,17 @@
%6 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %7 OpStore %2 %7
OpStore %3 %8 OpStore %3 %8
%10 = OpLoad %ulong %2 %9 = OpLoad %ulong %2
%9 = OpCopyObject %ulong %10
OpStore %4 %9 OpStore %4 %9
%12 = OpLoad %ulong %3 %10 = OpLoad %ulong %3
%11 = OpCopyObject %ulong %12 OpStore %5 %10
OpStore %5 %11 %12 = OpLoad %ulong %4
%14 = OpLoad %ulong %4 %15 = OpConvertUToPtr %_ptr_Generic_ulong %12
%17 = OpConvertUToPtr %_ptr_Generic_ulong %14 %11 = OpLoad %ulong %15
%13 = OpLoad %ulong %17 OpStore %6 %11
OpStore %6 %13 %13 = OpLoad %ulong %5
%15 = OpLoad %ulong %5 %14 = OpLoad %ulong %6
%16 = OpLoad %ulong %6 %16 = OpConvertUToPtr %_ptr_Generic_ulong %13
%18 = OpConvertUToPtr %_ptr_Generic_ulong %15 OpStore %16 %14
OpStore %18 %16
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

42
ptx/src/test/spirv_run/ld_st_implicit.spvtxt
View File

@ -7,20 +7,20 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%23 = OpExtInstImport "OpenCL.std" %21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "ld_st_implicit" OpEntryPoint Kernel %1 "ld_st_implicit"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%26 = OpTypeFunction %void %ulong %ulong %24 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%float = OpTypeFloat 32 %float = OpTypeFloat 32
%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float %_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%1 = OpFunction %void None %26 %1 = OpFunction %void None %24
%7 = OpFunctionParameter %ulong %7 = OpFunctionParameter %ulong
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%21 = OpLabel %19 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -28,24 +28,22 @@
%6 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %7 OpStore %2 %7
OpStore %3 %8 OpStore %3 %8
%10 = OpLoad %ulong %2 %9 = OpLoad %ulong %2
%9 = OpCopyObject %ulong %10
OpStore %4 %9 OpStore %4 %9
%12 = OpLoad %ulong %3 %10 = OpLoad %ulong %3
%11 = OpCopyObject %ulong %12 OpStore %5 %10
OpStore %5 %11 %12 = OpLoad %ulong %4
%14 = OpLoad %ulong %4 %16 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %12
%18 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %14 %15 = OpLoad %float %16
%17 = OpLoad %float %18 %29 = OpBitcast %uint %15
%31 = OpBitcast %uint %17 %11 = OpUConvert %ulong %29
%13 = OpUConvert %ulong %31 OpStore %6 %11
OpStore %6 %13 %13 = OpLoad %ulong %5
%15 = OpLoad %ulong %5 %14 = OpLoad %ulong %6
%16 = OpLoad %ulong %6 %17 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %13
%19 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %15 %30 = OpBitcast %ulong %14
%32 = OpBitcast %ulong %16 %31 = OpUConvert %uint %30
%33 = OpUConvert %uint %32 %18 = OpBitcast %float %31
%20 = OpBitcast %float %33 OpStore %17 %18
OpStore %19 %20
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

53
ptx/src/test/spirv_run/ld_st_offset.spvtxt
View File

@ -2,24 +2,27 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%32 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%30 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "ld_st_offset" OpEntryPoint Kernel %1 "ld_st_offset"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%35 = OpTypeFunction %void %ulong %ulong %33 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Generic_uint = OpTypePointer Generic %uint %_ptr_Generic_uint = OpTypePointer Generic %uint
%ulong_4 = OpConstant %ulong 4 %ulong_4 = OpConstant %ulong 4
%ulong_4_0 = OpConstant %ulong 4 %ulong_4_0 = OpConstant %ulong 4
%1 = OpFunction %void None %35 %1 = OpFunction %void None %33
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%30 = OpLabel %28 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -28,29 +31,27 @@
%7 = OpVariable %_ptr_Function_uint Function %7 = OpVariable %_ptr_Function_uint Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%24 = OpConvertUToPtr %_ptr_Generic_uint %13
%12 = OpLoad %uint %24
OpStore %6 %12
%15 = OpLoad %ulong %4 %15 = OpLoad %ulong %4
%26 = OpConvertUToPtr %_ptr_Generic_uint %15 %21 = OpIAdd %ulong %15 %ulong_4
%14 = OpLoad %uint %26 %25 = OpConvertUToPtr %_ptr_Generic_uint %21
OpStore %6 %14 %14 = OpLoad %uint %25
%17 = OpLoad %ulong %4 OpStore %7 %14
%23 = OpIAdd %ulong %17 %ulong_4 %16 = OpLoad %ulong %5
%27 = OpConvertUToPtr %_ptr_Generic_uint %23 %17 = OpLoad %uint %7
%16 = OpLoad %uint %27 %26 = OpConvertUToPtr %_ptr_Generic_uint %16
OpStore %7 %16 OpStore %26 %17
%18 = OpLoad %ulong %5 %18 = OpLoad %ulong %5
%19 = OpLoad %uint %7 %19 = OpLoad %uint %6
%28 = OpConvertUToPtr %_ptr_Generic_uint %18 %23 = OpIAdd %ulong %18 %ulong_4_0
OpStore %28 %19 %27 = OpConvertUToPtr %_ptr_Generic_uint %23
%20 = OpLoad %ulong %5 OpStore %27 %19
%21 = OpLoad %uint %6
%25 = OpIAdd %ulong %20 %ulong_4_0
%29 = OpConvertUToPtr %_ptr_Generic_uint %25
OpStore %29 %21
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

21
ptx/src/test/spirv_run/lg2.ptx Normal file
View File

@ -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;
}

47
ptx/src/test/spirv_run/lg2.spvtxt Normal file
View File

@ -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

39
ptx/src/test/spirv_run/local_align.spvtxt
View File

@ -2,26 +2,29 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%22 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%20 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "local_align" OpEntryPoint Kernel %1 "local_align"
OpDecorate %4 Alignment 8 OpDecorate %4 Alignment 8
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%25 = OpTypeFunction %void %ulong %ulong %23 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uchar = OpTypeInt 8 0
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%uint_8 = OpConstant %uint 8 %uint_8 = OpConstant %uint 8
%_arr_uchar_uint_8 = OpTypeArray %uchar %uint_8 %_arr_uchar_uint_8 = OpTypeArray %uchar %uint_8
%_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8 %_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%1 = OpFunction %void None %25 %1 = OpFunction %void None %23
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%20 = OpLabel %18 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function__arr_uchar_uint_8 Function %4 = OpVariable %_ptr_Function__arr_uchar_uint_8 Function
@ -30,19 +33,17 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %5 %10 OpStore %5 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %6 %11
OpStore %6 %12 %13 = OpLoad %ulong %5
%15 = OpLoad %ulong %5 %16 = OpConvertUToPtr %_ptr_Generic_ulong %13
%18 = OpConvertUToPtr %_ptr_Generic_ulong %15 %12 = OpLoad %ulong %16
%14 = OpLoad %ulong %18 OpStore %7 %12
OpStore %7 %14 %14 = OpLoad %ulong %6
%16 = OpLoad %ulong %6 %15 = OpLoad %ulong %7
%17 = OpLoad %ulong %7 %17 = OpConvertUToPtr %_ptr_Generic_ulong %14
%19 = OpConvertUToPtr %_ptr_Generic_ulong %16 OpStore %17 %15
OpStore %19 %17
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

76
ptx/src/test/spirv_run/mad_s32.spvtxt
View File

@ -2,15 +2,17 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64 OpCapability Float64
%48 = OpExtInstImport "OpenCL.std" %46 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "mad_s32" OpEntryPoint Kernel %1 "mad_s32"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%51 = OpTypeFunction %void %ulong %ulong %49 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
@ -19,10 +21,10 @@
%ulong_8 = OpConstant %ulong 8 %ulong_8 = OpConstant %ulong 8
%ulong_4_0 = OpConstant %ulong 4 %ulong_4_0 = OpConstant %ulong 4
%ulong_8_0 = OpConstant %ulong 8 %ulong_8_0 = OpConstant %ulong 8
%1 = OpFunction %void None %51 %1 = OpFunction %void None %49
%10 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong
%11 = OpFunctionParameter %ulong %11 = OpFunctionParameter %ulong
%46 = OpLabel %44 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -33,45 +35,43 @@
%9 = OpVariable %_ptr_Function_uint Function %9 = OpVariable %_ptr_Function_uint Function
OpStore %2 %10 OpStore %2 %10
OpStore %3 %11 OpStore %3 %11
%13 = OpLoad %ulong %2 %12 = OpLoad %ulong %2
%12 = OpCopyObject %ulong %13
OpStore %4 %12 OpStore %4 %12
%15 = OpLoad %ulong %3 %13 = OpLoad %ulong %3
%14 = OpCopyObject %ulong %15 OpStore %5 %13
OpStore %5 %14 %15 = OpLoad %ulong %4
%38 = OpConvertUToPtr %_ptr_Generic_uint %15
%14 = OpLoad %uint %38
OpStore %7 %14
%17 = OpLoad %ulong %4 %17 = OpLoad %ulong %4
%40 = OpConvertUToPtr %_ptr_Generic_uint %17 %31 = OpIAdd %ulong %17 %ulong_4
%16 = OpLoad %uint %40 %39 = OpConvertUToPtr %_ptr_Generic_uint %31
OpStore %7 %16 %16 = OpLoad %uint %39
OpStore %8 %16
%19 = OpLoad %ulong %4 %19 = OpLoad %ulong %4
%33 = OpIAdd %ulong %19 %ulong_4 %33 = OpIAdd %ulong %19 %ulong_8
%41 = OpConvertUToPtr %_ptr_Generic_uint %33 %40 = OpConvertUToPtr %_ptr_Generic_uint %33
%18 = OpLoad %uint %41 %18 = OpLoad %uint %40
OpStore %8 %18 OpStore %9 %18
%21 = OpLoad %ulong %4 %21 = OpLoad %uint %7
%35 = OpIAdd %ulong %21 %ulong_8 %22 = OpLoad %uint %8
%42 = OpConvertUToPtr %_ptr_Generic_uint %35 %23 = OpLoad %uint %9
%20 = OpLoad %uint %42 %54 = OpIMul %uint %21 %22
OpStore %9 %20 %20 = OpIAdd %uint %23 %54
%23 = OpLoad %uint %7 OpStore %6 %20
%24 = OpLoad %uint %8 %24 = OpLoad %ulong %5
%25 = OpLoad %uint %9 %25 = OpLoad %uint %6
%56 = OpIMul %uint %23 %24 %41 = OpConvertUToPtr %_ptr_Generic_uint %24
%22 = OpIAdd %uint %25 %56 OpStore %41 %25
OpStore %6 %22
%26 = OpLoad %ulong %5 %26 = OpLoad %ulong %5
%27 = OpLoad %uint %6 %27 = OpLoad %uint %6
%43 = OpConvertUToPtr %_ptr_Generic_uint %26 %35 = OpIAdd %ulong %26 %ulong_4_0
OpStore %43 %27 %42 = OpConvertUToPtr %_ptr_Generic_uint %35
OpStore %42 %27
%28 = OpLoad %ulong %5 %28 = OpLoad %ulong %5
%29 = OpLoad %uint %6 %29 = OpLoad %uint %6
%37 = OpIAdd %ulong %28 %ulong_4_0 %37 = OpIAdd %ulong %28 %ulong_8_0
%44 = OpConvertUToPtr %_ptr_Generic_uint %37 %43 = OpConvertUToPtr %_ptr_Generic_uint %37
OpStore %44 %29 OpStore %43 %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
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

48
ptx/src/test/spirv_run/max.spvtxt
View File

@ -7,21 +7,21 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%30 = OpExtInstImport "OpenCL.std" %28 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "max" OpEntryPoint Kernel %1 "max"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%33 = OpTypeFunction %void %ulong %ulong %31 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Generic_uint = OpTypePointer Generic %uint %_ptr_Generic_uint = OpTypePointer Generic %uint
%ulong_4 = OpConstant %ulong 4 %ulong_4 = OpConstant %ulong 4
%1 = OpFunction %void None %33 %1 = OpFunction %void None %31
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%28 = OpLabel %26 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -30,28 +30,26 @@
%7 = OpVariable %_ptr_Function_uint Function %7 = OpVariable %_ptr_Function_uint Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%23 = OpConvertUToPtr %_ptr_Generic_uint %13
%12 = OpLoad %uint %23
OpStore %6 %12
%15 = OpLoad %ulong %4 %15 = OpLoad %ulong %4
%25 = OpConvertUToPtr %_ptr_Generic_uint %15 %22 = OpIAdd %ulong %15 %ulong_4
%14 = OpLoad %uint %25 %24 = OpConvertUToPtr %_ptr_Generic_uint %22
OpStore %6 %14 %14 = OpLoad %uint %24
%17 = OpLoad %ulong %4 OpStore %7 %14
%24 = OpIAdd %ulong %17 %ulong_4 %17 = OpLoad %uint %6
%26 = OpConvertUToPtr %_ptr_Generic_uint %24 %18 = OpLoad %uint %7
%16 = OpLoad %uint %26 %16 = OpExtInst %uint %28 s_max %17 %18
OpStore %7 %16 OpStore %6 %16
%19 = OpLoad %uint %6 %19 = OpLoad %ulong %5
%20 = OpLoad %uint %7 %20 = OpLoad %uint %6
%18 = OpExtInst %uint %30 s_max %19 %20 %25 = OpConvertUToPtr %_ptr_Generic_uint %19
OpStore %6 %18 OpStore %25 %20
%21 = OpLoad %ulong %5
%22 = OpLoad %uint %6
%27 = OpConvertUToPtr %_ptr_Generic_uint %21
OpStore %27 %22
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

48
ptx/src/test/spirv_run/min.spvtxt
View File

@ -7,21 +7,21 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%30 = OpExtInstImport "OpenCL.std" %28 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "min" OpEntryPoint Kernel %1 "min"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%33 = OpTypeFunction %void %ulong %ulong %31 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Generic_uint = OpTypePointer Generic %uint %_ptr_Generic_uint = OpTypePointer Generic %uint
%ulong_4 = OpConstant %ulong 4 %ulong_4 = OpConstant %ulong 4
%1 = OpFunction %void None %33 %1 = OpFunction %void None %31
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%28 = OpLabel %26 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -30,28 +30,26 @@
%7 = OpVariable %_ptr_Function_uint Function %7 = OpVariable %_ptr_Function_uint Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%23 = OpConvertUToPtr %_ptr_Generic_uint %13
%12 = OpLoad %uint %23
OpStore %6 %12
%15 = OpLoad %ulong %4 %15 = OpLoad %ulong %4
%25 = OpConvertUToPtr %_ptr_Generic_uint %15 %22 = OpIAdd %ulong %15 %ulong_4
%14 = OpLoad %uint %25 %24 = OpConvertUToPtr %_ptr_Generic_uint %22
OpStore %6 %14 %14 = OpLoad %uint %24
%17 = OpLoad %ulong %4 OpStore %7 %14
%24 = OpIAdd %ulong %17 %ulong_4 %17 = OpLoad %uint %6
%26 = OpConvertUToPtr %_ptr_Generic_uint %24 %18 = OpLoad %uint %7
%16 = OpLoad %uint %26 %16 = OpExtInst %uint %28 s_min %17 %18
OpStore %7 %16 OpStore %6 %16
%19 = OpLoad %uint %6 %19 = OpLoad %ulong %5
%20 = OpLoad %uint %7 %20 = OpLoad %uint %6
%18 = OpExtInst %uint %30 s_min %19 %20 %25 = OpConvertUToPtr %_ptr_Generic_uint %19
OpStore %6 %18 OpStore %25 %20
%21 = OpLoad %ulong %5
%22 = OpLoad %uint %6
%27 = OpConvertUToPtr %_ptr_Generic_uint %21
OpStore %27 %22
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

7
ptx/src/test/spirv_run/mod.rs
View File

@ -60,8 +60,7 @@ test_ptx!(call, [1u64], [2u64]);
test_ptx!(vector, [1u32, 2u32], [3u32, 3u32]); test_ptx!(vector, [1u32, 2u32], [3u32, 3u32]);
test_ptx!(ld_st_offset, [1u32, 2u32], [2u32, 1u32]); test_ptx!(ld_st_offset, [1u32, 2u32], [2u32, 1u32]);
test_ptx!(ntid, [3u32], [4u32]); 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!(mov_address, [0xDEADu64], [0u64]);
test_ptx!(b64tof64, [111u64], [111u64]); test_ptx!(b64tof64, [111u64], [111u64]);
test_ptx!(implicit_param, [34u32], [34u32]); 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!(sqrt, [0.25f32], [0.5f32]);
test_ptx!(rsqrt, [0.25f64], [2f64]); test_ptx!(rsqrt, [0.25f64], [2f64]);
test_ptx!(neg, [181i32], [-181i32]); 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> { struct DisplayError<T: Debug> {
err: T, err: T,

43
ptx/src/test/spirv_run/mov.spvtxt
View File

@ -2,20 +2,23 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%24 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%22 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "mov" OpEntryPoint Kernel %1 "mov"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%27 = OpTypeFunction %void %ulong %ulong %25 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%1 = OpFunction %void None %27 %1 = OpFunction %void None %25
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%22 = OpLabel %20 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -24,22 +27,20 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%15 = OpLoad %ulong %4 %18 = OpConvertUToPtr %_ptr_Generic_ulong %13
%20 = OpConvertUToPtr %_ptr_Generic_ulong %15 %12 = OpLoad %ulong %18
%14 = OpLoad %ulong %20 OpStore %6 %12
OpStore %6 %14 %15 = OpLoad %ulong %6
%17 = OpLoad %ulong %6 %14 = OpCopyObject %ulong %15
%16 = OpCopyObject %ulong %17 OpStore %7 %14
OpStore %7 %16 %16 = OpLoad %ulong %5
%18 = OpLoad %ulong %5 %17 = OpLoad %ulong %7
%19 = OpLoad %ulong %7 %19 = OpConvertUToPtr %_ptr_Generic_ulong %16
%21 = OpConvertUToPtr %_ptr_Generic_ulong %18 OpStore %19 %17
OpStore %21 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

43
ptx/src/test/spirv_run/mul_hi.spvtxt
View File

@ -2,21 +2,24 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%25 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%23 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "mul_hi" OpEntryPoint Kernel %1 "mul_hi"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%28 = OpTypeFunction %void %ulong %ulong %26 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_2 = OpConstant %ulong 2 %ulong_2 = OpConstant %ulong 2
%1 = OpFunction %void None %28 %1 = OpFunction %void None %26
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%23 = OpLabel %21 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -25,22 +28,20 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%15 = OpLoad %ulong %4 %19 = OpConvertUToPtr %_ptr_Generic_ulong %13
%21 = OpConvertUToPtr %_ptr_Generic_ulong %15 %12 = OpLoad %ulong %19
%14 = OpLoad %ulong %21 OpStore %6 %12
OpStore %6 %14 %15 = OpLoad %ulong %6
%17 = OpLoad %ulong %6 %14 = OpExtInst %ulong %23 u_mul_hi %15 %ulong_2
%16 = OpExtInst %ulong %25 u_mul_hi %17 %ulong_2 OpStore %7 %14
OpStore %7 %16 %16 = OpLoad %ulong %5
%18 = OpLoad %ulong %5 %17 = OpLoad %ulong %7
%19 = OpLoad %ulong %7 %20 = OpConvertUToPtr %_ptr_Generic_ulong %16
%22 = OpConvertUToPtr %_ptr_Generic_ulong %18 OpStore %20 %17
OpStore %22 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

43
ptx/src/test/spirv_run/mul_lo.spvtxt
View File

@ -2,21 +2,24 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%25 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%23 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "mul_lo" OpEntryPoint Kernel %1 "mul_lo"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%28 = OpTypeFunction %void %ulong %ulong %26 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_2 = OpConstant %ulong 2 %ulong_2 = OpConstant %ulong 2
%1 = OpFunction %void None %28 %1 = OpFunction %void None %26
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%23 = OpLabel %21 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -25,22 +28,20 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%15 = OpLoad %ulong %4 %19 = OpConvertUToPtr %_ptr_Generic_ulong %13
%21 = OpConvertUToPtr %_ptr_Generic_ulong %15 %12 = OpLoad %ulong %19
%14 = OpLoad %ulong %21 OpStore %6 %12
OpStore %6 %14 %15 = OpLoad %ulong %6
%17 = OpLoad %ulong %6 %14 = OpIMul %ulong %15 %ulong_2
%16 = OpIMul %ulong %17 %ulong_2 OpStore %7 %14
OpStore %7 %16 %16 = OpLoad %ulong %5
%18 = OpLoad %ulong %5 %17 = OpLoad %ulong %7
%19 = OpLoad %ulong %7 %20 = OpConvertUToPtr %_ptr_Generic_ulong %16
%22 = OpConvertUToPtr %_ptr_Generic_ulong %18 OpStore %20 %17
OpStore %22 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

52
ptx/src/test/spirv_run/mul_non_ftz.spvtxt
View File

@ -7,25 +7,21 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
OpCapability DenormFlushToZero %28 = OpExtInstImport "OpenCL.std"
OpCapability DenormPreserve
OpExtension "SPV_KHR_float_controls"
%30 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "mul_non_ftz" OpEntryPoint Kernel %1 "mul_non_ftz"
OpExecutionMode %1 DenormPreserve 32
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%33 = OpTypeFunction %void %ulong %ulong %31 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%float = OpTypeFloat 32 %float = OpTypeFloat 32
%_ptr_Function_float = OpTypePointer Function %float %_ptr_Function_float = OpTypePointer Function %float
%_ptr_Generic_float = OpTypePointer Generic %float %_ptr_Generic_float = OpTypePointer Generic %float
%ulong_4 = OpConstant %ulong 4 %ulong_4 = OpConstant %ulong 4
%1 = OpFunction %void None %33 %1 = OpFunction %void None %31
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%28 = OpLabel %26 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -34,28 +30,26 @@
%7 = OpVariable %_ptr_Function_float Function %7 = OpVariable %_ptr_Function_float Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%23 = OpConvertUToPtr %_ptr_Generic_float %13
%12 = OpLoad %float %23
OpStore %6 %12
%15 = OpLoad %ulong %4 %15 = OpLoad %ulong %4
%25 = OpConvertUToPtr %_ptr_Generic_float %15 %22 = OpIAdd %ulong %15 %ulong_4
%14 = OpLoad %float %25 %24 = OpConvertUToPtr %_ptr_Generic_float %22
OpStore %6 %14 %14 = OpLoad %float %24
%17 = OpLoad %ulong %4 OpStore %7 %14
%24 = OpIAdd %ulong %17 %ulong_4 %17 = OpLoad %float %6
%26 = OpConvertUToPtr %_ptr_Generic_float %24 %18 = OpLoad %float %7
%16 = OpLoad %float %26 %16 = OpFMul %float %17 %18
OpStore %7 %16 OpStore %6 %16
%19 = OpLoad %float %6 %19 = OpLoad %ulong %5
%20 = OpLoad %float %7 %20 = OpLoad %float %6
%18 = OpFMul %float %19 %20 %25 = OpConvertUToPtr %_ptr_Generic_float %19
OpStore %6 %18 OpStore %25 %20
%21 = OpLoad %ulong %5
%22 = OpLoad %float %6
%27 = OpConvertUToPtr %_ptr_Generic_float %21
OpStore %27 %22
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

60
ptx/src/test/spirv_run/mul_wide.spvtxt
View File

@ -7,24 +7,24 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%32 = OpExtInstImport "OpenCL.std" %30 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "mul_wide" OpEntryPoint Kernel %1 "mul_wide"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%35 = OpTypeFunction %void %ulong %ulong %33 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint %_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%ulong_4 = OpConstant %ulong 4 %ulong_4 = OpConstant %ulong 4
%_struct_40 = OpTypeStruct %uint %uint %_struct_38 = OpTypeStruct %uint %uint
%v2uint = OpTypeVector %uint 2 %v2uint = OpTypeVector %uint 2
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%1 = OpFunction %void None %35 %1 = OpFunction %void None %33
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong
%30 = OpLabel %28 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -34,33 +34,31 @@
%8 = OpVariable %_ptr_Function_ulong Function %8 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %9 OpStore %2 %9
OpStore %3 %10 OpStore %3 %10
%12 = OpLoad %ulong %2 %11 = OpLoad %ulong %2
%11 = OpCopyObject %ulong %12
OpStore %4 %11 OpStore %4 %11
%14 = OpLoad %ulong %3 %12 = OpLoad %ulong %3
%13 = OpCopyObject %ulong %14 OpStore %5 %12
OpStore %5 %13 %14 = OpLoad %ulong %4
%24 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %14
%13 = OpLoad %uint %24
OpStore %6 %13
%16 = OpLoad %ulong %4 %16 = OpLoad %ulong %4
%26 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %16 %23 = OpIAdd %ulong %16 %ulong_4
%15 = OpLoad %uint %26 %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %23
OpStore %6 %15 %15 = OpLoad %uint %25
%18 = OpLoad %ulong %4 OpStore %7 %15
%25 = OpIAdd %ulong %18 %ulong_4 %18 = OpLoad %uint %6
%27 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %25 %19 = OpLoad %uint %7
%17 = OpLoad %uint %27 %39 = OpSMulExtended %_struct_38 %18 %19
OpStore %7 %17 %40 = OpCompositeExtract %uint %39 0
%20 = OpLoad %uint %6 %41 = OpCompositeExtract %uint %39 1
%21 = OpLoad %uint %7 %43 = OpCompositeConstruct %v2uint %40 %41
%41 = OpSMulExtended %_struct_40 %20 %21 %17 = OpBitcast %ulong %43
%42 = OpCompositeExtract %uint %41 0 OpStore %8 %17
%43 = OpCompositeExtract %uint %41 1 %20 = OpLoad %ulong %5
%45 = OpCompositeConstruct %v2uint %42 %43 %21 = OpLoad %ulong %8
%19 = OpBitcast %ulong %45 %26 = OpConvertUToPtr %_ptr_Generic_ulong %20
OpStore %8 %19 %27 = OpCopyObject %ulong %21
%22 = OpLoad %ulong %5 OpStore %26 %27
%23 = OpLoad %ulong %8
%28 = OpConvertUToPtr %_ptr_Generic_ulong %22
%29 = OpCopyObject %ulong %23
OpStore %28 %29
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

62
ptx/src/test/spirv_run/neg.spvtxt
View File

@ -2,46 +2,46 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%26 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "not" OpEntryPoint Kernel %1 "neg"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%29 = OpTypeFunction %void %ulong %ulong %24 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %uint = OpTypeInt 32 0
%1 = OpFunction %void None %29 %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Generic_uint = OpTypePointer Generic %uint
%1 = OpFunction %void None %24
%7 = OpFunctionParameter %ulong
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %19 = OpLabel
%24 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%5 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
%6 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_uint Function
%7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %7
OpStore %2 %8 OpStore %3 %8
OpStore %3 %9 %9 = OpLoad %ulong %2
%11 = OpLoad %ulong %2 OpStore %4 %9
%10 = OpCopyObject %ulong %11 %10 = OpLoad %ulong %3
OpStore %4 %10 OpStore %5 %10
%13 = OpLoad %ulong %3 %12 = OpLoad %ulong %4
%12 = OpCopyObject %ulong %13 %17 = OpConvertUToPtr %_ptr_Generic_uint %12
OpStore %5 %12 %11 = OpLoad %uint %17
%15 = OpLoad %ulong %4 OpStore %6 %11
%20 = OpConvertUToPtr %_ptr_Generic_ulong %15 %14 = OpLoad %uint %6
%14 = OpLoad %ulong %20 %13 = OpSNegate %uint %14
OpStore %6 %14 OpStore %6 %13
%17 = OpLoad %ulong %6 %15 = OpLoad %ulong %5
%22 = OpCopyObject %ulong %17 %16 = OpLoad %uint %6
%21 = OpNot %ulong %22 %18 = OpConvertUToPtr %_ptr_Generic_uint %15
%16 = OpCopyObject %ulong %21 OpStore %18 %16
OpStore %7 %16
%18 = OpLoad %ulong %5
%19 = OpLoad %ulong %7
%23 = OpConvertUToPtr %_ptr_Generic_ulong %18
OpStore %23 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

47
ptx/src/test/spirv_run/not.spvtxt
View File

@ -2,20 +2,23 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%26 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%24 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "not" OpEntryPoint Kernel %1 "not"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%29 = OpTypeFunction %void %ulong %ulong %27 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%1 = OpFunction %void None %29 %1 = OpFunction %void None %27
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%24 = OpLabel %22 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -24,24 +27,22 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%15 = OpLoad %ulong %4 %18 = OpConvertUToPtr %_ptr_Generic_ulong %13
%20 = OpConvertUToPtr %_ptr_Generic_ulong %15 %12 = OpLoad %ulong %18
%14 = OpLoad %ulong %20 OpStore %6 %12
OpStore %6 %14 %15 = OpLoad %ulong %6
%17 = OpLoad %ulong %6 %20 = OpCopyObject %ulong %15
%22 = OpCopyObject %ulong %17 %19 = OpNot %ulong %20
%21 = OpNot %ulong %22 %14 = OpCopyObject %ulong %19
%16 = OpCopyObject %ulong %21 OpStore %7 %14
OpStore %7 %16 %16 = OpLoad %ulong %5
%18 = OpLoad %ulong %5 %17 = OpLoad %ulong %7
%19 = OpLoad %ulong %7 %21 = OpConvertUToPtr %_ptr_Generic_ulong %16
%23 = OpConvertUToPtr %_ptr_Generic_ulong %18 OpStore %21 %17
OpStore %23 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

52
ptx/src/test/spirv_run/ntid.spvtxt
View File

@ -2,10 +2,12 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64 OpCapability Float64
%29 = OpExtInstImport "OpenCL.std" %27 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "ntid" %gl_WorkGroupSize OpEntryPoint Kernel %1 "ntid" %gl_WorkGroupSize
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
@ -15,14 +17,14 @@
%_ptr_UniformConstant_v4uint = OpTypePointer UniformConstant %v4uint %_ptr_UniformConstant_v4uint = OpTypePointer UniformConstant %v4uint
%gl_WorkGroupSize = OpVariable %_ptr_UniformConstant_v4uint UniformConstant %gl_WorkGroupSize = OpVariable %_ptr_UniformConstant_v4uint UniformConstant
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%35 = OpTypeFunction %void %ulong %ulong %33 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Generic_uint = OpTypePointer Generic %uint %_ptr_Generic_uint = OpTypePointer Generic %uint
%1 = OpFunction %void None %35 %1 = OpFunction %void None %33
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong
%27 = OpLabel %25 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -31,27 +33,25 @@
%7 = OpVariable %_ptr_Function_uint Function %7 = OpVariable %_ptr_Function_uint Function
OpStore %2 %9 OpStore %2 %9
OpStore %3 %10 OpStore %3 %10
%12 = OpLoad %ulong %2 %11 = OpLoad %ulong %2
%11 = OpCopyObject %ulong %12
OpStore %4 %11 OpStore %4 %11
%14 = OpLoad %ulong %3 %12 = OpLoad %ulong %3
%13 = OpCopyObject %ulong %14 OpStore %5 %12
OpStore %5 %13 %14 = OpLoad %ulong %4
%16 = OpLoad %ulong %4 %23 = OpConvertUToPtr %_ptr_Generic_uint %14
%25 = OpConvertUToPtr %_ptr_Generic_uint %16 %13 = OpLoad %uint %23
%15 = OpLoad %uint %25 OpStore %6 %13
OpStore %6 %15 %16 = OpLoad %v4uint %gl_WorkGroupSize
%18 = OpLoad %v4uint %gl_WorkGroupSize %22 = OpCompositeExtract %uint %16 0
%24 = OpCompositeExtract %uint %18 0 %15 = OpCopyObject %uint %22
%17 = OpCopyObject %uint %24 OpStore %7 %15
OpStore %7 %17 %18 = OpLoad %uint %6
%20 = OpLoad %uint %6 %19 = OpLoad %uint %7
%21 = OpLoad %uint %7 %17 = OpIAdd %uint %18 %19
%19 = OpIAdd %uint %20 %21 OpStore %6 %17
OpStore %6 %19 %20 = OpLoad %ulong %5
%22 = OpLoad %ulong %5 %21 = OpLoad %uint %6
%23 = OpLoad %uint %6 %24 = OpConvertUToPtr %_ptr_Generic_uint %20
%26 = OpConvertUToPtr %_ptr_Generic_uint %22 OpStore %24 %21
OpStore %26 %23
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

54
ptx/src/test/spirv_run/or.spvtxt
View File

@ -7,19 +7,19 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%33 = OpExtInstImport "OpenCL.std" %31 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "or" OpEntryPoint Kernel %1 "or"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%36 = OpTypeFunction %void %ulong %ulong %34 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_8 = OpConstant %ulong 8 %ulong_8 = OpConstant %ulong 8
%1 = OpFunction %void None %36 %1 = OpFunction %void None %34
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%31 = OpLabel %29 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -28,31 +28,29 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%23 = OpConvertUToPtr %_ptr_Generic_ulong %13
%12 = OpLoad %ulong %23
OpStore %6 %12
%15 = OpLoad %ulong %4 %15 = OpLoad %ulong %4
%25 = OpConvertUToPtr %_ptr_Generic_ulong %15 %22 = OpIAdd %ulong %15 %ulong_8
%14 = OpLoad %ulong %25 %24 = OpConvertUToPtr %_ptr_Generic_ulong %22
OpStore %6 %14 %14 = OpLoad %ulong %24
%17 = OpLoad %ulong %4 OpStore %7 %14
%24 = OpIAdd %ulong %17 %ulong_8 %17 = OpLoad %ulong %6
%26 = OpConvertUToPtr %_ptr_Generic_ulong %24 %18 = OpLoad %ulong %7
%16 = OpLoad %ulong %26 %26 = OpCopyObject %ulong %17
OpStore %7 %16 %27 = OpCopyObject %ulong %18
%19 = OpLoad %ulong %6 %25 = OpBitwiseOr %ulong %26 %27
%20 = OpLoad %ulong %7 %16 = OpCopyObject %ulong %25
%28 = OpCopyObject %ulong %19 OpStore %6 %16
%29 = OpCopyObject %ulong %20 %19 = OpLoad %ulong %5
%27 = OpBitwiseOr %ulong %28 %29 %20 = OpLoad %ulong %6
%18 = OpCopyObject %ulong %27 %28 = OpConvertUToPtr %_ptr_Generic_ulong %19
OpStore %6 %18 OpStore %28 %20
%21 = OpLoad %ulong %5
%22 = OpLoad %ulong %6
%30 = OpConvertUToPtr %_ptr_Generic_ulong %21
OpStore %30 %22
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

74
ptx/src/test/spirv_run/pred_not.spvtxt
View File

@ -2,15 +2,17 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64 OpCapability Float64
%44 = OpExtInstImport "OpenCL.std" %42 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "pred_not" OpEntryPoint Kernel %1 "pred_not"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%47 = OpTypeFunction %void %ulong %ulong %45 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%bool = OpTypeBool %bool = OpTypeBool
%_ptr_Function_bool = OpTypePointer Function %bool %_ptr_Function_bool = OpTypePointer Function %bool
@ -20,10 +22,10 @@
%false = OpConstantFalse %bool %false = OpConstantFalse %bool
%ulong_1 = OpConstant %ulong 1 %ulong_1 = OpConstant %ulong 1
%ulong_2 = OpConstant %ulong 2 %ulong_2 = OpConstant %ulong 2
%1 = OpFunction %void None %47 %1 = OpFunction %void None %45
%14 = OpFunctionParameter %ulong %14 = OpFunctionParameter %ulong
%15 = OpFunctionParameter %ulong %15 = OpFunctionParameter %ulong
%42 = OpLabel %40 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -34,45 +36,43 @@
%9 = OpVariable %_ptr_Function_bool Function %9 = OpVariable %_ptr_Function_bool Function
OpStore %2 %14 OpStore %2 %14
OpStore %3 %15 OpStore %3 %15
%17 = OpLoad %ulong %2 %16 = OpLoad %ulong %2
%16 = OpCopyObject %ulong %17
OpStore %4 %16 OpStore %4 %16
%19 = OpLoad %ulong %3 %17 = OpLoad %ulong %3
%18 = OpCopyObject %ulong %19 OpStore %5 %17
OpStore %5 %18 %19 = OpLoad %ulong %4
%37 = OpConvertUToPtr %_ptr_Generic_ulong %19
%18 = OpLoad %ulong %37
OpStore %6 %18
%21 = OpLoad %ulong %4 %21 = OpLoad %ulong %4
%39 = OpConvertUToPtr %_ptr_Generic_ulong %21 %34 = OpIAdd %ulong %21 %ulong_8
%20 = OpLoad %ulong %39 %38 = OpConvertUToPtr %_ptr_Generic_ulong %34
OpStore %6 %20 %20 = OpLoad %ulong %38
%23 = OpLoad %ulong %4 OpStore %7 %20
%36 = OpIAdd %ulong %23 %ulong_8 %23 = OpLoad %ulong %6
%40 = OpConvertUToPtr %_ptr_Generic_ulong %36 %24 = OpLoad %ulong %7
%22 = OpLoad %ulong %40 %22 = OpULessThan %bool %23 %24
OpStore %7 %22 OpStore %9 %22
%25 = OpLoad %ulong %6 %26 = OpLoad %bool %9
%26 = OpLoad %ulong %7 %25 = OpSelect %bool %26 %false %true
%24 = OpULessThan %bool %25 %26 OpStore %9 %25
OpStore %9 %24 %27 = OpLoad %bool %9
%28 = OpLoad %bool %9 OpBranchConditional %27 %10 %11
%27 = OpSelect %bool %28 %false %true
OpStore %9 %27
%29 = OpLoad %bool %9
OpBranchConditional %29 %10 %11
%10 = OpLabel %10 = OpLabel
%30 = OpCopyObject %ulong %ulong_1 %28 = OpCopyObject %ulong %ulong_1
OpStore %8 %30 OpStore %8 %28
OpBranch %11 OpBranch %11
%11 = OpLabel %11 = OpLabel
%31 = OpLoad %bool %9 %29 = OpLoad %bool %9
OpBranchConditional %31 %13 %12 OpBranchConditional %29 %13 %12
%12 = OpLabel %12 = OpLabel
%32 = OpCopyObject %ulong %ulong_2 %30 = OpCopyObject %ulong %ulong_2
OpStore %8 %32 OpStore %8 %30
OpBranch %13 OpBranch %13
%13 = OpLabel %13 = OpLabel
%33 = OpLoad %ulong %5 %31 = OpLoad %ulong %5
%34 = OpLoad %ulong %8 %32 = OpLoad %ulong %8
%41 = OpConvertUToPtr %_ptr_Generic_ulong %33 %39 = OpConvertUToPtr %_ptr_Generic_ulong %31
OpStore %41 %34 OpStore %39 %32
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

38
ptx/src/test/spirv_run/rcp.spvtxt
View File

@ -7,24 +7,22 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
OpExtension "SPV_KHR_float_controls" %21 = OpExtInstImport "OpenCL.std"
%23 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "rcp" OpEntryPoint Kernel %1 "rcp"
OpExecutionMode %1 DenormPreserve 32 OpDecorate %13 FPFastMathMode AllowRecip
OpDecorate %15 FPFastMathMode AllowRecip
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%26 = OpTypeFunction %void %ulong %ulong %24 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%float = OpTypeFloat 32 %float = OpTypeFloat 32
%_ptr_Function_float = OpTypePointer Function %float %_ptr_Function_float = OpTypePointer Function %float
%_ptr_Generic_float = OpTypePointer Generic %float %_ptr_Generic_float = OpTypePointer Generic %float
%float_1 = OpConstant %float 1 %float_1 = OpConstant %float 1
%1 = OpFunction %void None %26 %1 = OpFunction %void None %24
%7 = OpFunctionParameter %ulong %7 = OpFunctionParameter %ulong
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%21 = OpLabel %19 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -32,22 +30,20 @@
%6 = OpVariable %_ptr_Function_float Function %6 = OpVariable %_ptr_Function_float Function
OpStore %2 %7 OpStore %2 %7
OpStore %3 %8 OpStore %3 %8
%10 = OpLoad %ulong %2 %9 = OpLoad %ulong %2
%9 = OpCopyObject %ulong %10
OpStore %4 %9 OpStore %4 %9
%12 = OpLoad %ulong %3 %10 = OpLoad %ulong %3
%11 = OpCopyObject %ulong %12 OpStore %5 %10
OpStore %5 %11 %12 = OpLoad %ulong %4
%14 = OpLoad %ulong %4 %17 = OpConvertUToPtr %_ptr_Generic_float %12
%19 = OpConvertUToPtr %_ptr_Generic_float %14 %11 = OpLoad %float %17
%13 = OpLoad %float %19 OpStore %6 %11
%14 = OpLoad %float %6
%13 = OpFDiv %float %float_1 %14
OpStore %6 %13 OpStore %6 %13
%15 = OpLoad %ulong %5
%16 = OpLoad %float %6 %16 = OpLoad %float %6
%15 = OpFDiv %float %float_1 %16 %18 = OpConvertUToPtr %_ptr_Generic_float %15
OpStore %6 %15 OpStore %18 %16
%17 = OpLoad %ulong %5
%18 = OpLoad %float %6
%20 = OpConvertUToPtr %_ptr_Generic_float %17
OpStore %20 %18
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

5
ptx/src/test/spirv_run/reg_local.ptx
View File

@ -11,14 +11,13 @@
.reg .u64 in_addr; .reg .u64 in_addr;
.reg .u64 out_addr; .reg .u64 out_addr;
.reg .b64 temp; .reg .b64 temp;
.reg .s64 unused;
ld.param.u64 in_addr, [input]; ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output]; ld.param.u64 out_addr, [output];
ld.global.u64 temp, [in_addr]; ld.global.u64 temp, [in_addr];
st.u64 [local_x], temp + 1; st.u64 [local_x], temp + 1;
ld.u64 temp, [local_x]; ld.u64 temp, [local_x+0];
st.global.u64 [out_addr], temp; st.global.u64 [out_addr+0], temp;
ret; ret;
} }

76
ptx/src/test/spirv_run/reg_local.spvtxt
View File

@ -2,62 +2,66 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%35 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%34 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "reg_local" OpEntryPoint Kernel %1 "reg_local"
OpDecorate %4 Alignment 8 OpDecorate %4 Alignment 8
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%38 = OpTypeFunction %void %ulong %ulong %37 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uchar = OpTypeInt 8 0
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%uint_8 = OpConstant %uint 8 %uint_8 = OpConstant %uint 8
%_arr_uchar_uint_8 = OpTypeArray %uchar %uint_8 %_arr_uchar_uint_8 = OpTypeArray %uchar %uint_8
%_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8 %_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%ulong_1 = OpConstant %ulong 1 %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 %9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong %32 = OpLabel
%33 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function__arr_uchar_uint_8 Function %4 = OpVariable %_ptr_Function__arr_uchar_uint_8 Function
%5 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
%6 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
%8 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8
OpStore %2 %9 OpStore %3 %9
OpStore %3 %10 %10 = OpLoad %ulong %2
%12 = OpLoad %ulong %2 OpStore %5 %10
%11 = OpCopyObject %ulong %12 %11 = OpLoad %ulong %3
OpStore %5 %11 OpStore %6 %11
%14 = OpLoad %ulong %3 %13 = OpLoad %ulong %5
%13 = OpCopyObject %ulong %14 %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %13
OpStore %6 %13 %24 = OpLoad %ulong %25
%16 = OpLoad %ulong %5 %12 = OpCopyObject %ulong %24
%25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %16 OpStore %7 %12
%26 = OpLoad %ulong %25 %14 = OpLoad %ulong %7
%15 = OpCopyObject %ulong %26 %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 OpStore %7 %15
%18 = OpLoad %ulong %7 %16 = OpLoad %ulong %6
%27 = OpCopyObject %ulong %18 %17 = OpLoad %ulong %7
%24 = OpIAdd %ulong %27 %ulong_1 %23 = OpIAdd %ulong %16 %ulong_0_0
%28 = OpCopyObject %ulong %24 %30 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %23
%17 = OpBitcast %ulong %28 %31 = OpCopyObject %ulong %17
OpStore %4 %17 OpStore %30 %31
%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
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

103
ptx/src/test/spirv_run/rsqrt.spvtxt
View File

@ -1,56 +1,47 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 31 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %21 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "rsqrt"
OpCapability Float64 %void = OpTypeVoid
; OpCapability FunctionFloatControlINTEL %ulong = OpTypeInt 64 0
; OpExtension "SPV_INTEL_float_controls2" %24 = OpTypeFunction %void %ulong %ulong
%23 = OpExtInstImport "OpenCL.std" %_ptr_Function_ulong = OpTypePointer Function %ulong
OpMemoryModel Physical64 OpenCL %double = OpTypeFloat 64
OpEntryPoint Kernel %1 "rsqrt" %_ptr_Function_double = OpTypePointer Function %double
OpDecorate %1 FunctionDenormModeINTEL 64 Preserve %_ptr_Generic_double = OpTypePointer Generic %double
%24 = OpTypeVoid %1 = OpFunction %void None %24
%25 = OpTypeInt 64 0 %7 = OpFunctionParameter %ulong
%26 = OpTypeFunction %24 %25 %25 %8 = OpFunctionParameter %ulong
%27 = OpTypePointer Function %25 %19 = OpLabel
%28 = OpTypeFloat 64 %2 = OpVariable %_ptr_Function_ulong Function
%29 = OpTypePointer Function %28 %3 = OpVariable %_ptr_Function_ulong Function
%30 = OpTypePointer Generic %28 %4 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %24 None %26 %5 = OpVariable %_ptr_Function_ulong Function
%7 = OpFunctionParameter %25 %6 = OpVariable %_ptr_Function_double Function
%8 = OpFunctionParameter %25 OpStore %2 %7
%21 = OpLabel OpStore %3 %8
%2 = OpVariable %27 Function %9 = OpLoad %ulong %2
%3 = OpVariable %27 Function OpStore %4 %9
%4 = OpVariable %27 Function %10 = OpLoad %ulong %3
%5 = OpVariable %27 Function OpStore %5 %10
%6 = OpVariable %29 Function %12 = OpLoad %ulong %4
OpStore %2 %7 %17 = OpConvertUToPtr %_ptr_Generic_double %12
OpStore %3 %8 %11 = OpLoad %double %17
%10 = OpLoad %25 %2 OpStore %6 %11
%9 = OpCopyObject %25 %10 %14 = OpLoad %double %6
OpStore %4 %9 %13 = OpExtInst %double %21 native_rsqrt %14
%12 = OpLoad %25 %3 OpStore %6 %13
%11 = OpCopyObject %25 %12 %15 = OpLoad %ulong %5
OpStore %5 %11 %16 = OpLoad %double %6
%14 = OpLoad %25 %4 %18 = OpConvertUToPtr %_ptr_Generic_double %15
%19 = OpConvertUToPtr %30 %14 OpStore %18 %16
%13 = OpLoad %28 %19 OpReturn
OpStore %6 %13 OpFunctionEnd
%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

122
ptx/src/test/spirv_run/selp.spvtxt
View File

@ -1,65 +1,57 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 40 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %29 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "selp"
OpCapability Float64 %void = OpTypeVoid
; OpCapability FunctionFloatControlINTEL %ulong = OpTypeInt 64 0
; OpExtension "SPV_INTEL_float_controls2" %32 = OpTypeFunction %void %ulong %ulong
%31 = OpExtInstImport "OpenCL.std" %_ptr_Function_ulong = OpTypePointer Function %ulong
OpMemoryModel Physical64 OpenCL %ushort = OpTypeInt 16 0
OpEntryPoint Kernel %1 "selp" %_ptr_Function_ushort = OpTypePointer Function %ushort
%32 = OpTypeVoid %_ptr_Generic_ushort = OpTypePointer Generic %ushort
%33 = OpTypeInt 64 0 %ulong_2 = OpConstant %ulong 2
%34 = OpTypeFunction %32 %33 %33 %bool = OpTypeBool
%35 = OpTypePointer Function %33 %false = OpConstantFalse %bool
%36 = OpTypeInt 16 0 %1 = OpFunction %void None %32
%37 = OpTypePointer Function %36 %8 = OpFunctionParameter %ulong
%38 = OpTypePointer Generic %36 %9 = OpFunctionParameter %ulong
%23 = OpConstant %33 2 %27 = OpLabel
%39 = OpTypeBool %2 = OpVariable %_ptr_Function_ulong Function
%25 = OpConstantFalse %39 %3 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %32 None %34 %4 = OpVariable %_ptr_Function_ulong Function
%8 = OpFunctionParameter %33 %5 = OpVariable %_ptr_Function_ulong Function
%9 = OpFunctionParameter %33 %6 = OpVariable %_ptr_Function_ushort Function
%29 = OpLabel %7 = OpVariable %_ptr_Function_ushort Function
%2 = OpVariable %35 Function OpStore %2 %8
%3 = OpVariable %35 Function OpStore %3 %9
%4 = OpVariable %35 Function %10 = OpLoad %ulong %2
%5 = OpVariable %35 Function OpStore %4 %10
%6 = OpVariable %37 Function %11 = OpLoad %ulong %3
%7 = OpVariable %37 Function OpStore %5 %11
OpStore %2 %8 %13 = OpLoad %ulong %4
OpStore %3 %9 %24 = OpConvertUToPtr %_ptr_Generic_ushort %13
%11 = OpLoad %33 %2 %12 = OpLoad %ushort %24
%10 = OpCopyObject %33 %11 OpStore %6 %12
OpStore %4 %10 %15 = OpLoad %ulong %4
%13 = OpLoad %33 %3 %22 = OpIAdd %ulong %15 %ulong_2
%12 = OpCopyObject %33 %13 %25 = OpConvertUToPtr %_ptr_Generic_ushort %22
OpStore %5 %12 %14 = OpLoad %ushort %25
%15 = OpLoad %33 %4 OpStore %7 %14
%26 = OpConvertUToPtr %38 %15 %17 = OpLoad %ushort %6
%14 = OpLoad %36 %26 %18 = OpLoad %ushort %7
OpStore %6 %14 %16 = OpSelect %ushort %false %18 %18
%17 = OpLoad %33 %4 OpStore %6 %16
%24 = OpIAdd %33 %17 %23 %19 = OpLoad %ulong %5
%27 = OpConvertUToPtr %38 %24 %20 = OpLoad %ushort %6
%16 = OpLoad %36 %27 %26 = OpConvertUToPtr %_ptr_Generic_ushort %19
OpStore %7 %16 OpStore %26 %20
%19 = OpLoad %36 %6 OpReturn
%20 = OpLoad %36 %7 OpFunctionEnd
%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

69
ptx/src/test/spirv_run/setp.spvtxt
View File

@ -2,14 +2,17 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%42 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%40 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "setp" OpEntryPoint Kernel %1 "setp"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%45 = OpTypeFunction %void %ulong %ulong %43 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%bool = OpTypeBool %bool = OpTypeBool
%_ptr_Function_bool = OpTypePointer Function %bool %_ptr_Function_bool = OpTypePointer Function %bool
@ -17,10 +20,10 @@
%ulong_8 = OpConstant %ulong 8 %ulong_8 = OpConstant %ulong 8
%ulong_1 = OpConstant %ulong 1 %ulong_1 = OpConstant %ulong 1
%ulong_2 = OpConstant %ulong 2 %ulong_2 = OpConstant %ulong 2
%1 = OpFunction %void None %45 %1 = OpFunction %void None %43
%14 = OpFunctionParameter %ulong %14 = OpFunctionParameter %ulong
%15 = OpFunctionParameter %ulong %15 = OpFunctionParameter %ulong
%40 = OpLabel %38 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -31,42 +34,40 @@
%9 = OpVariable %_ptr_Function_bool Function %9 = OpVariable %_ptr_Function_bool Function
OpStore %2 %14 OpStore %2 %14
OpStore %3 %15 OpStore %3 %15
%17 = OpLoad %ulong %2 %16 = OpLoad %ulong %2
%16 = OpCopyObject %ulong %17
OpStore %4 %16 OpStore %4 %16
%19 = OpLoad %ulong %3 %17 = OpLoad %ulong %3
%18 = OpCopyObject %ulong %19 OpStore %5 %17
OpStore %5 %18 %19 = OpLoad %ulong %4
%35 = OpConvertUToPtr %_ptr_Generic_ulong %19
%18 = OpLoad %ulong %35
OpStore %6 %18
%21 = OpLoad %ulong %4 %21 = OpLoad %ulong %4
%37 = OpConvertUToPtr %_ptr_Generic_ulong %21 %32 = OpIAdd %ulong %21 %ulong_8
%20 = OpLoad %ulong %37 %36 = OpConvertUToPtr %_ptr_Generic_ulong %32
OpStore %6 %20 %20 = OpLoad %ulong %36
%23 = OpLoad %ulong %4 OpStore %7 %20
%34 = OpIAdd %ulong %23 %ulong_8 %23 = OpLoad %ulong %6
%38 = OpConvertUToPtr %_ptr_Generic_ulong %34 %24 = OpLoad %ulong %7
%22 = OpLoad %ulong %38 %22 = OpULessThan %bool %23 %24
OpStore %7 %22 OpStore %9 %22
%25 = OpLoad %ulong %6 %25 = OpLoad %bool %9
%26 = OpLoad %ulong %7 OpBranchConditional %25 %10 %11
%24 = OpULessThan %bool %25 %26
OpStore %9 %24
%27 = OpLoad %bool %9
OpBranchConditional %27 %10 %11
%10 = OpLabel %10 = OpLabel
%28 = OpCopyObject %ulong %ulong_1 %26 = OpCopyObject %ulong %ulong_1
OpStore %8 %28 OpStore %8 %26
OpBranch %11 OpBranch %11
%11 = OpLabel %11 = OpLabel
%29 = OpLoad %bool %9 %27 = OpLoad %bool %9
OpBranchConditional %29 %13 %12 OpBranchConditional %27 %13 %12
%12 = OpLabel %12 = OpLabel
%30 = OpCopyObject %ulong %ulong_2 %28 = OpCopyObject %ulong %ulong_2
OpStore %8 %30 OpStore %8 %28
OpBranch %13 OpBranch %13
%13 = OpLabel %13 = OpLabel
%31 = OpLoad %ulong %5 %29 = OpLoad %ulong %5
%32 = OpLoad %ulong %8 %30 = OpLoad %ulong %8
%39 = OpConvertUToPtr %_ptr_Generic_ulong %31 %37 = OpConvertUToPtr %_ptr_Generic_ulong %29
OpStore %39 %32 OpStore %37 %30
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

140
ptx/src/test/spirv_run/shared_ptr_32.spvtxt
View File

@ -1,74 +1,66 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 47 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %32 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "shared_ptr_32" %4
OpCapability Float64 OpDecorate %4 Alignment 4
; OpCapability FunctionFloatControlINTEL %void = OpTypeVoid
; OpExtension "SPV_INTEL_float_controls2" %uint = OpTypeInt 32 0
%34 = OpExtInstImport "OpenCL.std" %uchar = OpTypeInt 8 0
OpMemoryModel Physical64 OpenCL %uint_128 = OpConstant %uint 128
OpEntryPoint Kernel %1 "shared_ptr_32" %4 %_arr_uchar_uint_128 = OpTypeArray %uchar %uint_128
OpDecorate %4 Alignment 4 %_ptr_Workgroup__arr_uchar_uint_128 = OpTypePointer Workgroup %_arr_uchar_uint_128
%35 = OpTypeVoid %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
%36 = OpTypeInt 32 0 %ulong = OpTypeInt 64 0
%37 = OpTypeInt 8 0 %40 = OpTypeFunction %void %ulong %ulong
%38 = OpConstant %36 128 %_ptr_Function_ulong = OpTypePointer Function %ulong
%39 = OpTypeArray %37 %38 %_ptr_Function_uint = OpTypePointer Function %uint
%40 = OpTypePointer Workgroup %39 %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%4 = OpVariable %40 Workgroup %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%41 = OpTypeInt 64 0 %uint_0 = OpConstant %uint 0
%42 = OpTypeFunction %35 %41 %41 %1 = OpFunction %void None %40
%43 = OpTypePointer Function %41 %10 = OpFunctionParameter %ulong
%44 = OpTypePointer Function %36 %11 = OpFunctionParameter %ulong
%45 = OpTypePointer CrossWorkgroup %41 %30 = OpLabel
%46 = OpTypePointer Workgroup %41 %2 = OpVariable %_ptr_Function_ulong Function
%25 = OpConstant %36 0 %3 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %35 None %42 %5 = OpVariable %_ptr_Function_ulong Function
%10 = OpFunctionParameter %41 %6 = OpVariable %_ptr_Function_ulong Function
%11 = OpFunctionParameter %41 %7 = OpVariable %_ptr_Function_uint Function
%32 = OpLabel %8 = OpVariable %_ptr_Function_ulong Function
%2 = OpVariable %43 Function %9 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %43 Function OpStore %2 %10
%5 = OpVariable %43 Function OpStore %3 %11
%6 = OpVariable %43 Function %12 = OpLoad %ulong %2
%7 = OpVariable %44 Function OpStore %5 %12
%8 = OpVariable %43 Function %13 = OpLoad %ulong %3
%9 = OpVariable %43 Function OpStore %6 %13
OpStore %2 %10 %25 = OpConvertPtrToU %uint %4
OpStore %3 %11 %14 = OpCopyObject %uint %25
%13 = OpLoad %41 %2 OpStore %7 %14
%12 = OpCopyObject %41 %13 %16 = OpLoad %ulong %5
OpStore %5 %12 %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %16
%15 = OpLoad %41 %3 %15 = OpLoad %ulong %26
%14 = OpCopyObject %41 %15 OpStore %8 %15
OpStore %6 %14 %17 = OpLoad %uint %7
%27 = OpConvertPtrToU %36 %4 %18 = OpLoad %ulong %8
%16 = OpCopyObject %36 %27 %27 = OpConvertUToPtr %_ptr_Workgroup_ulong %17
OpStore %7 %16 OpStore %27 %18
%18 = OpLoad %41 %5 %20 = OpLoad %uint %7
%28 = OpConvertUToPtr %45 %18 %24 = OpIAdd %uint %20 %uint_0
%17 = OpLoad %41 %28 %28 = OpConvertUToPtr %_ptr_Workgroup_ulong %24
OpStore %8 %17 %19 = OpLoad %ulong %28
%19 = OpLoad %36 %7 OpStore %9 %19
%20 = OpLoad %41 %8 %21 = OpLoad %ulong %6
%29 = OpConvertUToPtr %46 %19 %22 = OpLoad %ulong %9
OpStore %29 %20 %29 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %21
%22 = OpLoad %36 %7 OpStore %29 %22
%26 = OpIAdd %36 %22 %25 OpReturn
%30 = OpConvertUToPtr %46 %26 OpFunctionEnd
%21 = OpLoad %41 %30
OpStore %9 %21
%23 = OpLoad %41 %6
%24 = OpLoad %41 %9
%31 = OpConvertUToPtr %45 %23
OpStore %31 %24
OpReturn
OpFunctionEnd

122
ptx/src/test/spirv_run/shared_variable.spvtxt
View File

@ -1,65 +1,57 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 39 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %25 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "shared_variable" %4
OpCapability Float64 OpDecorate %4 Alignment 4
; OpCapability FunctionFloatControlINTEL %void = OpTypeVoid
; OpExtension "SPV_INTEL_float_controls2" %uint = OpTypeInt 32 0
%27 = OpExtInstImport "OpenCL.std" %uchar = OpTypeInt 8 0
OpMemoryModel Physical64 OpenCL %uint_128 = OpConstant %uint 128
OpEntryPoint Kernel %1 "shared_variable" %4 %_arr_uchar_uint_128 = OpTypeArray %uchar %uint_128
OpDecorate %4 Alignment 4 %_ptr_Workgroup__arr_uchar_uint_128 = OpTypePointer Workgroup %_arr_uchar_uint_128
%28 = OpTypeVoid %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
%29 = OpTypeInt 32 0 %ulong = OpTypeInt 64 0
%30 = OpTypeInt 8 0 %33 = OpTypeFunction %void %ulong %ulong
%31 = OpConstant %29 128 %_ptr_Function_ulong = OpTypePointer Function %ulong
%32 = OpTypeArray %30 %31 %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%33 = OpTypePointer Workgroup %32 %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%4 = OpVariable %33 Workgroup %1 = OpFunction %void None %33
%34 = OpTypeInt 64 0 %9 = OpFunctionParameter %ulong
%35 = OpTypeFunction %28 %34 %34 %10 = OpFunctionParameter %ulong
%36 = OpTypePointer Function %34 %23 = OpLabel
%37 = OpTypePointer CrossWorkgroup %34 %2 = OpVariable %_ptr_Function_ulong Function
%38 = OpTypePointer Workgroup %34 %3 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %28 None %35 %5 = OpVariable %_ptr_Function_ulong Function
%9 = OpFunctionParameter %34 %6 = OpVariable %_ptr_Function_ulong Function
%10 = OpFunctionParameter %34 %7 = OpVariable %_ptr_Function_ulong Function
%25 = OpLabel %8 = OpVariable %_ptr_Function_ulong Function
%2 = OpVariable %36 Function OpStore %2 %9
%3 = OpVariable %36 Function OpStore %3 %10
%5 = OpVariable %36 Function %11 = OpLoad %ulong %2
%6 = OpVariable %36 Function OpStore %5 %11
%7 = OpVariable %36 Function %12 = OpLoad %ulong %3
%8 = OpVariable %36 Function OpStore %6 %12
OpStore %2 %9 %14 = OpLoad %ulong %5
OpStore %3 %10 %19 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %14
%12 = OpLoad %34 %2 %13 = OpLoad %ulong %19
%11 = OpCopyObject %34 %12 OpStore %7 %13
OpStore %5 %11 %15 = OpLoad %ulong %7
%14 = OpLoad %34 %3 %20 = OpBitcast %_ptr_Workgroup_ulong %4
%13 = OpCopyObject %34 %14 OpStore %20 %15
OpStore %6 %13 %21 = OpBitcast %_ptr_Workgroup_ulong %4
%16 = OpLoad %34 %5 %16 = OpLoad %ulong %21
%21 = OpConvertUToPtr %37 %16 OpStore %8 %16
%15 = OpLoad %34 %21 %17 = OpLoad %ulong %6
OpStore %7 %15 %18 = OpLoad %ulong %8
%17 = OpLoad %34 %7 %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %17
%22 = OpBitcast %38 %4 OpStore %22 %18
OpStore %22 %17 OpReturn
%23 = OpBitcast %38 %4 OpFunctionEnd
%18 = OpLoad %34 %23
OpStore %8 %18
%19 = OpLoad %34 %6
%20 = OpLoad %34 %8
%24 = OpConvertUToPtr %37 %19
OpStore %24 %20
OpReturn
OpFunctionEnd

47
ptx/src/test/spirv_run/shl.spvtxt
View File

@ -2,22 +2,25 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%27 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%25 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "shl" OpEntryPoint Kernel %1 "shl"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%30 = OpTypeFunction %void %ulong %ulong %28 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%1 = OpFunction %void None %30 %1 = OpFunction %void None %28
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%25 = OpLabel %23 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -26,24 +29,22 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%15 = OpLoad %ulong %4 %19 = OpConvertUToPtr %_ptr_Generic_ulong %13
%21 = OpConvertUToPtr %_ptr_Generic_ulong %15 %12 = OpLoad %ulong %19
%14 = OpLoad %ulong %21 OpStore %6 %12
OpStore %6 %14 %15 = OpLoad %ulong %6
%17 = OpLoad %ulong %6 %21 = OpCopyObject %ulong %15
%23 = OpCopyObject %ulong %17 %20 = OpShiftLeftLogical %ulong %21 %uint_2
%22 = OpShiftLeftLogical %ulong %23 %uint_2 %14 = OpCopyObject %ulong %20
%16 = OpCopyObject %ulong %22 OpStore %7 %14
OpStore %7 %16 %16 = OpLoad %ulong %5
%18 = OpLoad %ulong %5 %17 = OpLoad %ulong %7
%19 = OpLoad %ulong %7 %22 = OpConvertUToPtr %_ptr_Generic_ulong %16
%24 = OpConvertUToPtr %_ptr_Generic_ulong %18 OpStore %22 %17
OpStore %24 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

34
ptx/src/test/spirv_run/shr.spvtxt
View File

@ -7,21 +7,21 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%24 = OpExtInstImport "OpenCL.std" %22 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "shr" OpEntryPoint Kernel %1 "shr"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%27 = OpTypeFunction %void %ulong %ulong %25 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Generic_uint = OpTypePointer Generic %uint %_ptr_Generic_uint = OpTypePointer Generic %uint
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%1 = OpFunction %void None %27 %1 = OpFunction %void None %25
%7 = OpFunctionParameter %ulong %7 = OpFunctionParameter %ulong
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%22 = OpLabel %20 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -29,22 +29,20 @@
%6 = OpVariable %_ptr_Function_uint Function %6 = OpVariable %_ptr_Function_uint Function
OpStore %2 %7 OpStore %2 %7
OpStore %3 %8 OpStore %3 %8
%10 = OpLoad %ulong %2 %9 = OpLoad %ulong %2
%9 = OpCopyObject %ulong %10
OpStore %4 %9 OpStore %4 %9
%12 = OpLoad %ulong %3 %10 = OpLoad %ulong %3
%11 = OpCopyObject %ulong %12 OpStore %5 %10
OpStore %5 %11 %12 = OpLoad %ulong %4
%14 = OpLoad %ulong %4 %18 = OpConvertUToPtr %_ptr_Generic_uint %12
%20 = OpConvertUToPtr %_ptr_Generic_uint %14 %11 = OpLoad %uint %18
%13 = OpLoad %uint %20 OpStore %6 %11
%14 = OpLoad %uint %6
%13 = OpShiftRightArithmetic %uint %14 %uint_1
OpStore %6 %13 OpStore %6 %13
%15 = OpLoad %ulong %5
%16 = OpLoad %uint %6 %16 = OpLoad %uint %6
%15 = OpShiftRightArithmetic %uint %16 %uint_1 %19 = OpConvertUToPtr %_ptr_Generic_uint %15
OpStore %6 %15 OpStore %19 %16
%17 = OpLoad %ulong %5
%18 = OpLoad %uint %6
%21 = OpConvertUToPtr %_ptr_Generic_uint %17
OpStore %21 %18
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

21
ptx/src/test/spirv_run/sin.ptx Normal file
View File

@ -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;
}

47
ptx/src/test/spirv_run/sin.spvtxt Normal file
View File

@ -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

103
ptx/src/test/spirv_run/sqrt.spvtxt
View File

@ -1,56 +1,47 @@
; SPIR-V OpCapability GenericPointer
; Version: 1.3 OpCapability Linkage
; Generator: rspirv OpCapability Addresses
; Bound: 31 OpCapability Kernel
OpCapability GenericPointer OpCapability Int8
OpCapability Linkage OpCapability Int16
OpCapability Addresses OpCapability Int64
OpCapability Kernel OpCapability Float16
OpCapability Int8 OpCapability Float64
OpCapability Int16 %21 = OpExtInstImport "OpenCL.std"
OpCapability Int64 OpMemoryModel Physical64 OpenCL
OpCapability Float16 OpEntryPoint Kernel %1 "sqrt"
OpCapability Float64 %void = OpTypeVoid
; OpCapability FunctionFloatControlINTEL %ulong = OpTypeInt 64 0
; OpExtension "SPV_INTEL_float_controls2" %24 = OpTypeFunction %void %ulong %ulong
%23 = OpExtInstImport "OpenCL.std" %_ptr_Function_ulong = OpTypePointer Function %ulong
OpMemoryModel Physical64 OpenCL %float = OpTypeFloat 32
OpEntryPoint Kernel %1 "sqrt" %_ptr_Function_float = OpTypePointer Function %float
OpDecorate %1 FunctionDenormModeINTEL 32 Preserve %_ptr_Generic_float = OpTypePointer Generic %float
%24 = OpTypeVoid %1 = OpFunction %void None %24
%25 = OpTypeInt 64 0 %7 = OpFunctionParameter %ulong
%26 = OpTypeFunction %24 %25 %25 %8 = OpFunctionParameter %ulong
%27 = OpTypePointer Function %25 %19 = OpLabel
%28 = OpTypeFloat 32 %2 = OpVariable %_ptr_Function_ulong Function
%29 = OpTypePointer Function %28 %3 = OpVariable %_ptr_Function_ulong Function
%30 = OpTypePointer Generic %28 %4 = OpVariable %_ptr_Function_ulong Function
%1 = OpFunction %24 None %26 %5 = OpVariable %_ptr_Function_ulong Function
%7 = OpFunctionParameter %25 %6 = OpVariable %_ptr_Function_float Function
%8 = OpFunctionParameter %25 OpStore %2 %7
%21 = OpLabel OpStore %3 %8
%2 = OpVariable %27 Function %9 = OpLoad %ulong %2
%3 = OpVariable %27 Function OpStore %4 %9
%4 = OpVariable %27 Function %10 = OpLoad %ulong %3
%5 = OpVariable %27 Function OpStore %5 %10
%6 = OpVariable %29 Function %12 = OpLoad %ulong %4
OpStore %2 %7 %17 = OpConvertUToPtr %_ptr_Generic_float %12
OpStore %3 %8 %11 = OpLoad %float %17
%10 = OpLoad %25 %2 OpStore %6 %11
%9 = OpCopyObject %25 %10 %14 = OpLoad %float %6
OpStore %4 %9 %13 = OpExtInst %float %21 native_sqrt %14
%12 = OpLoad %25 %3 OpStore %6 %13
%11 = OpCopyObject %25 %12 %15 = OpLoad %ulong %5
OpStore %5 %11 %16 = OpLoad %float %6
%14 = OpLoad %25 %4 %18 = OpConvertUToPtr %_ptr_Generic_float %15
%19 = OpConvertUToPtr %30 %14 OpStore %18 %16
%13 = OpLoad %28 %19 OpReturn
OpStore %6 %13 OpFunctionEnd
%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

38
ptx/src/test/spirv_run/sub.spvtxt
View File

@ -7,19 +7,19 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%25 = OpExtInstImport "OpenCL.std" %23 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "sub" OpEntryPoint Kernel %1 "sub"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%28 = OpTypeFunction %void %ulong %ulong %26 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_1 = OpConstant %ulong 1 %ulong_1 = OpConstant %ulong 1
%1 = OpFunction %void None %28 %1 = OpFunction %void None %26
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%23 = OpLabel %21 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -28,22 +28,20 @@
%7 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %8 OpStore %2 %8
OpStore %3 %9 OpStore %3 %9
%11 = OpLoad %ulong %2 %10 = OpLoad %ulong %2
%10 = OpCopyObject %ulong %11
OpStore %4 %10 OpStore %4 %10
%13 = OpLoad %ulong %3 %11 = OpLoad %ulong %3
%12 = OpCopyObject %ulong %13 OpStore %5 %11
OpStore %5 %12 %13 = OpLoad %ulong %4
%15 = OpLoad %ulong %4 %19 = OpConvertUToPtr %_ptr_Generic_ulong %13
%21 = OpConvertUToPtr %_ptr_Generic_ulong %15 %12 = OpLoad %ulong %19
%14 = OpLoad %ulong %21 OpStore %6 %12
OpStore %6 %14 %15 = OpLoad %ulong %6
%17 = OpLoad %ulong %6 %14 = OpISub %ulong %15 %ulong_1
%16 = OpISub %ulong %17 %ulong_1 OpStore %7 %14
OpStore %7 %16 %16 = OpLoad %ulong %5
%18 = OpLoad %ulong %5 %17 = OpLoad %ulong %7
%19 = OpLoad %ulong %7 %20 = OpConvertUToPtr %_ptr_Generic_ulong %16
%22 = OpConvertUToPtr %_ptr_Generic_ulong %18 OpStore %20 %17
OpStore %22 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

52
ptx/src/test/spirv_run/vector.spvtxt
View File

@ -2,26 +2,29 @@
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64
OpCapability Int8 OpCapability Int8
%60 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%57 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %31 "vector" OpEntryPoint Kernel %31 "vector"
%void = OpTypeVoid %void = OpTypeVoid
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%v2uint = OpTypeVector %uint 2 %v2uint = OpTypeVector %uint 2
%64 = OpTypeFunction %v2uint %v2uint %61 = OpTypeFunction %v2uint %v2uint
%_ptr_Function_v2uint = OpTypePointer Function %v2uint %_ptr_Function_v2uint = OpTypePointer Function %v2uint
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%68 = OpTypeFunction %void %ulong %ulong %65 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_v2uint = OpTypePointer Generic %v2uint %_ptr_Generic_v2uint = OpTypePointer Generic %v2uint
%1 = OpFunction %v2uint None %64 %1 = OpFunction %v2uint None %61
%7 = OpFunctionParameter %v2uint %7 = OpFunctionParameter %v2uint
%30 = OpLabel %30 = OpLabel
%3 = OpVariable %_ptr_Function_v2uint Function
%2 = OpVariable %_ptr_Function_v2uint Function %2 = OpVariable %_ptr_Function_v2uint Function
%3 = OpVariable %_ptr_Function_v2uint Function
%4 = OpVariable %_ptr_Function_v2uint Function %4 = OpVariable %_ptr_Function_v2uint Function
%5 = OpVariable %_ptr_Function_uint Function %5 = OpVariable %_ptr_Function_uint Function
%6 = OpVariable %_ptr_Function_uint Function %6 = OpVariable %_ptr_Function_uint Function
@ -57,10 +60,10 @@
%26 = OpLoad %v2uint %2 %26 = OpLoad %v2uint %2
OpReturnValue %26 OpReturnValue %26
OpFunctionEnd OpFunctionEnd
%31 = OpFunction %void None %68 %31 = OpFunction %void None %65
%40 = OpFunctionParameter %ulong %40 = OpFunctionParameter %ulong
%41 = OpFunctionParameter %ulong %41 = OpFunctionParameter %ulong
%58 = OpLabel %55 = OpLabel
%32 = OpVariable %_ptr_Function_ulong Function %32 = OpVariable %_ptr_Function_ulong Function
%33 = OpVariable %_ptr_Function_ulong Function %33 = OpVariable %_ptr_Function_ulong Function
%34 = OpVariable %_ptr_Function_ulong Function %34 = OpVariable %_ptr_Function_ulong Function
@ -71,27 +74,24 @@
%39 = OpVariable %_ptr_Function_ulong Function %39 = OpVariable %_ptr_Function_ulong Function
OpStore %32 %40 OpStore %32 %40
OpStore %33 %41 OpStore %33 %41
%43 = OpLoad %ulong %32 %42 = OpLoad %ulong %32
%42 = OpCopyObject %ulong %43
OpStore %34 %42 OpStore %34 %42
%45 = OpLoad %ulong %33 %43 = OpLoad %ulong %33
%44 = OpCopyObject %ulong %45 OpStore %35 %43
OpStore %35 %44 %45 = OpLoad %ulong %34
%47 = OpLoad %ulong %34 %52 = OpConvertUToPtr %_ptr_Generic_v2uint %45
%54 = OpConvertUToPtr %_ptr_Generic_v2uint %47 %44 = OpLoad %v2uint %52
%46 = OpLoad %v2uint %54 OpStore %36 %44
%47 = OpLoad %v2uint %36
%46 = OpFunctionCall %v2uint %1 %47
OpStore %36 %46 OpStore %36 %46
%49 = OpLoad %v2uint %36 %49 = OpLoad %v2uint %36
%48 = OpFunctionCall %v2uint %1 %49 %53 = OpBitcast %ulong %49
OpStore %36 %48 %48 = OpCopyObject %ulong %53
OpStore %39 %48
%50 = OpLoad %ulong %35
%51 = OpLoad %v2uint %36 %51 = OpLoad %v2uint %36
%55 = OpBitcast %ulong %51 %54 = OpConvertUToPtr %_ptr_Generic_v2uint %50
%56 = OpCopyObject %ulong %55 OpStore %54 %51
%50 = OpCopyObject %ulong %56
OpStore %39 %50
%52 = OpLoad %ulong %35
%53 = OpLoad %v2uint %36
%57 = OpConvertUToPtr %_ptr_Generic_v2uint %52
OpStore %57 %53
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

156
ptx/src/test/spirv_run/vector_extract.spvtxt
View File

@ -7,12 +7,12 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%75 = OpExtInstImport "OpenCL.std" %73 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "vector_extract" OpEntryPoint Kernel %1 "vector_extract"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%78 = OpTypeFunction %void %ulong %ulong %76 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%ushort = OpTypeInt 16 0 %ushort = OpTypeInt 16 0
%_ptr_Function_ushort = OpTypePointer Function %ushort %_ptr_Function_ushort = OpTypePointer Function %ushort
@ -21,10 +21,10 @@
%uchar = OpTypeInt 8 0 %uchar = OpTypeInt 8 0
%v4uchar = OpTypeVector %uchar 4 %v4uchar = OpTypeVector %uchar 4
%_ptr_CrossWorkgroup_v4uchar = OpTypePointer CrossWorkgroup %v4uchar %_ptr_CrossWorkgroup_v4uchar = OpTypePointer CrossWorkgroup %v4uchar
%1 = OpFunction %void None %78 %1 = OpFunction %void None %76
%11 = OpFunctionParameter %ulong %11 = OpFunctionParameter %ulong
%12 = OpFunctionParameter %ulong %12 = OpFunctionParameter %ulong
%73 = OpLabel %71 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -36,89 +36,87 @@
%10 = OpVariable %_ptr_Function_v4ushort Function %10 = OpVariable %_ptr_Function_v4ushort Function
OpStore %2 %11 OpStore %2 %11
OpStore %3 %12 OpStore %3 %12
%14 = OpLoad %ulong %2 %13 = OpLoad %ulong %2
%13 = OpCopyObject %ulong %14
OpStore %4 %13 OpStore %4 %13
%16 = OpLoad %ulong %3 %14 = OpLoad %ulong %3
%15 = OpCopyObject %ulong %16 OpStore %5 %14
OpStore %5 %15 %19 = OpLoad %ulong %4
%21 = OpLoad %ulong %4 %61 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %19
%63 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %21 %43 = OpLoad %v4uchar %61
%45 = OpLoad %v4uchar %63 %62 = OpCompositeExtract %uchar %43 0
%64 = OpCompositeExtract %uchar %45 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 %87 = OpBitcast %uchar %64
%17 = OpUConvert %ushort %87 %17 = OpUConvert %ushort %87
%65 = OpCompositeExtract %uchar %45 1 %65 = OpCompositeExtract %uchar %43 3
%88 = OpBitcast %uchar %65 %88 = OpBitcast %uchar %65
%18 = OpUConvert %ushort %88 %18 = OpUConvert %ushort %88
%66 = OpCompositeExtract %uchar %45 2 OpStore %6 %15
%89 = OpBitcast %uchar %66 OpStore %7 %16
%19 = OpUConvert %ushort %89 OpStore %8 %17
%67 = OpCompositeExtract %uchar %45 3 OpStore %9 %18
%90 = OpBitcast %uchar %67 %21 = OpLoad %ushort %7
%20 = OpUConvert %ushort %90 %22 = OpLoad %ushort %8
OpStore %6 %17 %23 = OpLoad %ushort %9
OpStore %7 %18 %24 = OpLoad %ushort %6
OpStore %8 %19 %44 = OpUndef %v4ushort
OpStore %9 %20 %45 = OpCompositeInsert %v4ushort %21 %44 0
%23 = OpLoad %ushort %7 %46 = OpCompositeInsert %v4ushort %22 %45 1
%24 = OpLoad %ushort %8 %47 = OpCompositeInsert %v4ushort %23 %46 2
%25 = OpLoad %ushort %9 %48 = OpCompositeInsert %v4ushort %24 %47 3
%26 = OpLoad %ushort %6 %20 = OpCopyObject %v4ushort %48
%46 = OpUndef %v4ushort OpStore %10 %20
%47 = OpCompositeInsert %v4ushort %23 %46 0 %29 = OpLoad %v4ushort %10
%48 = OpCompositeInsert %v4ushort %24 %47 1 %49 = OpCopyObject %v4ushort %29
%49 = OpCompositeInsert %v4ushort %25 %48 2 %25 = OpCompositeExtract %ushort %49 0
%50 = OpCompositeInsert %v4ushort %26 %49 3 %26 = OpCompositeExtract %ushort %49 1
%22 = OpCopyObject %v4ushort %50 %27 = OpCompositeExtract %ushort %49 2
OpStore %10 %22 %28 = OpCompositeExtract %ushort %49 3
%31 = OpLoad %v4ushort %10 OpStore %8 %25
%51 = OpCopyObject %v4ushort %31 OpStore %9 %26
%27 = OpCompositeExtract %ushort %51 0 OpStore %6 %27
%28 = OpCompositeExtract %ushort %51 1 OpStore %7 %28
%29 = OpCompositeExtract %ushort %51 2 %34 = OpLoad %ushort %8
%30 = OpCompositeExtract %ushort %51 3 %35 = OpLoad %ushort %9
OpStore %8 %27 %36 = OpLoad %ushort %6
OpStore %9 %28 %37 = OpLoad %ushort %7
OpStore %6 %29 %51 = OpUndef %v4ushort
OpStore %7 %30 %52 = OpCompositeInsert %v4ushort %34 %51 0
%36 = OpLoad %ushort %8 %53 = OpCompositeInsert %v4ushort %35 %52 1
%37 = OpLoad %ushort %9 %54 = OpCompositeInsert %v4ushort %36 %53 2
%38 = OpLoad %ushort %6 %55 = OpCompositeInsert %v4ushort %37 %54 3
%39 = OpLoad %ushort %7 %50 = OpCopyObject %v4ushort %55
%53 = OpUndef %v4ushort %30 = OpCompositeExtract %ushort %50 0
%54 = OpCompositeInsert %v4ushort %36 %53 0 %31 = OpCompositeExtract %ushort %50 1
%55 = OpCompositeInsert %v4ushort %37 %54 1 %32 = OpCompositeExtract %ushort %50 2
%56 = OpCompositeInsert %v4ushort %38 %55 2 %33 = OpCompositeExtract %ushort %50 3
%57 = OpCompositeInsert %v4ushort %39 %56 3 OpStore %9 %30
%52 = OpCopyObject %v4ushort %57 OpStore %6 %31
%32 = OpCompositeExtract %ushort %52 0 OpStore %7 %32
%33 = OpCompositeExtract %ushort %52 1 OpStore %8 %33
%34 = OpCompositeExtract %ushort %52 2 %38 = OpLoad %ulong %5
%35 = OpCompositeExtract %ushort %52 3 %39 = OpLoad %ushort %6
OpStore %9 %32 %40 = OpLoad %ushort %7
OpStore %6 %33 %41 = OpLoad %ushort %8
OpStore %7 %34 %42 = OpLoad %ushort %9
OpStore %8 %35 %56 = OpUndef %v4uchar
%40 = OpLoad %ulong %5 %89 = OpBitcast %ushort %39
%41 = OpLoad %ushort %6 %66 = OpUConvert %uchar %89
%42 = OpLoad %ushort %7 %57 = OpCompositeInsert %v4uchar %66 %56 0
%43 = OpLoad %ushort %8 %90 = OpBitcast %ushort %40
%44 = OpLoad %ushort %9 %67 = OpUConvert %uchar %90
%58 = OpUndef %v4uchar %58 = OpCompositeInsert %v4uchar %67 %57 1
%91 = OpBitcast %ushort %41 %91 = OpBitcast %ushort %41
%68 = OpUConvert %uchar %91 %68 = OpUConvert %uchar %91
%59 = OpCompositeInsert %v4uchar %68 %58 0 %59 = OpCompositeInsert %v4uchar %68 %58 2
%92 = OpBitcast %ushort %42 %92 = OpBitcast %ushort %42
%69 = OpUConvert %uchar %92 %69 = OpUConvert %uchar %92
%60 = OpCompositeInsert %v4uchar %69 %59 1 %60 = OpCompositeInsert %v4uchar %69 %59 3
%93 = OpBitcast %ushort %43 %70 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %38
%70 = OpUConvert %uchar %93 OpStore %70 %60
%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
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

970
ptx/src/translate.rs
View File

File diff suppressed because it is too large Load Diff