Applied-Software/ZLUDA
5
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-07-18 01:36:24 +03:00
Code Issues Packages Projects Releases Wiki Activity

Implement support for PTX call instruction

Browse Source
This commit is contained in:
Andrzej Janik
2020-09-06 21:38:01 +02:00
parent bbb3a6c5cb
commit 76afbeba63
20 changed files with 1691 additions and 887 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
doc/NOTES.md
View File

@ -75,3 +75,9 @@ CUDA <-> L0
* context ~= context (1.0+) * context ~= context (1.0+)
* graph ~= command list * graph ~= command list
* module ~= module * module ~= module
IGC
---
* IGC is extremely brittle and segfaults on fairly innocent code:
* OpBitcast of pointer to uint
* OpCopyMemory of alloca'd variable

44
ptx/src/ast.rs
View File

@ -51,23 +51,34 @@ pub struct Module<'a> {
pub functions: Vec<ParsedFunction<'a>>, pub functions: Vec<ParsedFunction<'a>>,
} }
pub enum FunctionHeader<'a, P: ArgParams> { pub enum MethodDecl<'a, P: ArgParams> {
Func(Vec<Argument<P>>, P::ID), Func(Vec<FnArgument<P>>, P::ID, Vec<FnArgument<P>>),
Kernel(&'a str), Kernel(&'a str, Vec<KernelArgument<P>>),
} }
pub struct Function<'a, P: ArgParams, S> { pub struct Function<'a, P: ArgParams, S> {
pub func_directive: FunctionHeader<'a, P>, pub func_directive: MethodDecl<'a, P>,
pub args: Vec<Argument<P>>,
pub body: Option<Vec<S>>, pub body: Option<Vec<S>>,
} }
pub type ParsedFunction<'a> = Function<'a, ParsedArgParams<'a>, Statement<ParsedArgParams<'a>>>; pub type ParsedFunction<'a> = Function<'a, ParsedArgParams<'a>, Statement<ParsedArgParams<'a>>>;
#[derive(Default)] pub struct FnArgument<P: ArgParams> {
pub struct Argument<P: ArgParams> { pub base: KernelArgument<P>,
pub state_space: FnArgStateSpace,
}
#[derive(Copy, Clone, PartialEq, Eq)]
pub enum FnArgStateSpace {
Reg,
Param,
}
#[derive(Default, Copy, Clone)]
pub struct KernelArgument<P: ArgParams> {
pub name: P::ID, pub name: P::ID,
pub a_type: ScalarType, pub a_type: ScalarType,
// TODO: turn length into part of type definition
pub length: u32, pub length: u32,
} }
@ -222,28 +233,26 @@ pub enum Instruction<P: ArgParams> {
Shl(ShlType, Arg3<P>), Shl(ShlType, Arg3<P>),
St(StData, Arg2St<P>), St(StData, Arg2St<P>),
Ret(RetData), Ret(RetData),
Call(CallData, ArgCall<P>), Call(CallInst<P>),
Abs(AbsDetails, Arg2<P>), Abs(AbsDetails, Arg2<P>),
} }
pub struct CallData {
pub uniform: bool,
}
pub struct AbsDetails { pub struct AbsDetails {
pub flush_to_zero: bool, pub flush_to_zero: bool,
pub typ: ScalarType, pub typ: ScalarType,
} }
pub struct ArgCall<P: ArgParams> { pub struct CallInst<P: ArgParams> {
pub uniform: bool,
pub ret_params: Vec<P::ID>, pub ret_params: Vec<P::ID>,
pub func: P::ID, pub func: P::ID,
pub param_list: Vec<P::ID>, pub param_list: Vec<P::CallOperand>,
} }
pub trait ArgParams { pub trait ArgParams {
type ID; type ID;
type Operand; type Operand;
type CallOperand;
type MovOperand; type MovOperand;
} }
@ -254,6 +263,7 @@ pub struct ParsedArgParams<'a> {
impl<'a> ArgParams for ParsedArgParams<'a> { impl<'a> ArgParams for ParsedArgParams<'a> {
type ID = &'a str; type ID = &'a str;
type Operand = Operand<&'a str>; type Operand = Operand<&'a str>;
type CallOperand = CallOperand<&'a str>;
type MovOperand = MovOperand<&'a str>; type MovOperand = MovOperand<&'a str>;
} }
@ -304,6 +314,12 @@ pub enum Operand<ID> {
Imm(i128), Imm(i128),
} }
#[derive(Copy, Clone)]
pub enum CallOperand<ID> {
Reg(ID),
Imm(i128),
}
pub enum MovOperand<ID> { pub enum MovOperand<ID> {
Op(Operand<ID>), Op(Operand<ID>),
Vec(String, String), Vec(String, String),

64
ptx/src/ptx.lalrpop
View File

@ -202,8 +202,7 @@ AddressSize = {
Function: ast::Function<'input, ast::ParsedArgParams<'input>, ast::Statement<ast::ParsedArgParams<'input>>> = { Function: ast::Function<'input, ast::ParsedArgParams<'input>, ast::Statement<ast::ParsedArgParams<'input>>> = {
LinkingDirective* LinkingDirective*
<func_directive:FunctionHeader> <func_directive:MethodDecl>
<args:Arguments>
<body:FunctionBody> => ast::Function{<>} <body:FunctionBody> => ast::Function{<>}
}; };
@ -213,24 +212,43 @@ LinkingDirective = {
".weak" ".weak"
}; };
FunctionHeader: ast::FunctionHeader<'input, ast::ParsedArgParams<'input>> = { MethodDecl: ast::MethodDecl<'input, ast::ParsedArgParams<'input>> = {
".entry" <name:ExtendedID> => ast::FunctionHeader::Kernel(name), ".entry" <name:ExtendedID> <params:KernelArguments> => ast::MethodDecl::Kernel(name, params),
".func" <args:Arguments?> <name:ExtendedID> => ast::FunctionHeader::Func(args.unwrap_or_else(|| Vec::new()), name) ".func" <ret_vals:FnArguments?> <name:ExtendedID> <params:FnArguments> => ast::MethodDecl::Func(ret_vals.unwrap_or_else(|| Vec::new()), name, params)
}; };
Arguments: Vec<ast::Argument<ast::ParsedArgParams<'input>>> = { KernelArguments: Vec<ast::KernelArgument<ast::ParsedArgParams<'input>>> = {
"(" <args:Comma<FunctionInput>> ")" => args "(" <args:Comma<KernelInput>> ")" => args
} };
FnArguments: Vec<ast::FnArgument<ast::ParsedArgParams<'input>>> = {
"(" <args:Comma<FnInput>> ")" => args
};
FnInput: ast::FnArgument<ast::ParsedArgParams<'input>> = {
".reg" <_type:ScalarType> <name:ExtendedID> => {
ast::FnArgument {
base: ast::KernelArgument {a_type: _type, name: name, length: 1 },
state_space: ast::FnArgStateSpace::Reg,
}
},
<p:KernelInput> => {
ast::FnArgument {
base: p,
state_space: ast::FnArgStateSpace::Param,
}
}
};
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parameter-state-space // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parameter-state-space
FunctionInput: ast::Argument<ast::ParsedArgParams<'input>> = { KernelInput: ast::KernelArgument<ast::ParsedArgParams<'input>> = {
".param" <_type:ScalarType> <name:ExtendedID> => { ".param" <_type:ScalarType> <name:ExtendedID> => {
ast::Argument {a_type: _type, name: name, length: 1 } ast::KernelArgument {a_type: _type, name: name, length: 1 }
}, },
".param" <a_type:ScalarType> <name:ExtendedID> "[" <length:Num> "]" => { ".param" <a_type:ScalarType> <name:ExtendedID> "[" <length:Num> "]" => {
let length = length.parse::<u32>(); let length = length.parse::<u32>();
let length = length.unwrap_with(errors); let length = length.unwrap_with(errors);
ast::Argument { a_type: a_type, name: name, length: length } ast::KernelArgument { a_type: a_type, name: name, length: length }
} }
}; };
@ -856,7 +874,10 @@ CvtaSize: ast::CvtaSize = {
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-call // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-call
InstCall: ast::Instruction<ast::ParsedArgParams<'input>> = { InstCall: ast::Instruction<ast::ParsedArgParams<'input>> = {
"call" <u:".uni"?> <a:ArgCall> => ast::Instruction::Call(ast::CallData { uniform: u.is_some() }, a) "call" <u:".uni"?> <args:ArgCall> => {
let (ret_params, func, param_list) = args;
ast::Instruction::Call(ast::CallInst { uniform: u.is_some(), ret_params, func, param_list })
}
}; };
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-abs // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-abs
@ -900,6 +921,15 @@ Operand: ast::Operand<&'input str> = {
} }
}; };
CallOperand: ast::CallOperand<&'input str> = {
<r:ExtendedID> => ast::CallOperand::Reg(r),
<o:Num> => {
let offset = o.parse::<i128>();
let offset = offset.unwrap_with(errors);
ast::CallOperand::Imm(offset)
}
};
MovOperand: ast::MovOperand<&'input str> = { MovOperand: ast::MovOperand<&'input str> = {
<o:Operand> => ast::MovOperand::Op(o), <o:Operand> => ast::MovOperand::Op(o),
<o:VectorOperand> => { <o:VectorOperand> => {
@ -938,10 +968,12 @@ Arg5: ast::Arg5<ast::ParsedArgParams<'input>> = {
<dst1:ExtendedID> <dst2:OptionalDst?> "," <src1:Operand> "," <src2:Operand> "," "!"? <src3:Operand> => ast::Arg5{<>} <dst1:ExtendedID> <dst2:OptionalDst?> "," <src1:Operand> "," <src2:Operand> "," "!"? <src3:Operand> => ast::Arg5{<>}
}; };
ArgCall: ast::ArgCall<ast::ParsedArgParams<'input>> = { ArgCall: (Vec<&'input str>, &'input str, Vec<ast::CallOperand<&'input str>>) = {
"(" <ret_params:Comma<ExtendedID>> ")" "," <func:ExtendedID> "," "(" <param_list:Comma<ExtendedID>> ")" => ast::ArgCall{<>}, "(" <ret_params:Comma<ExtendedID>> ")" "," <func:ExtendedID> "," "(" <param_list:Comma<CallOperand>> ")" => {
<func:ExtendedID> "," "(" <param_list:Comma<ExtendedID>> ")" => ast::ArgCall{ret_params: Vec::new(), func, param_list}, (ret_params, func, param_list)
<func:ExtendedID> => ast::ArgCall{ret_params: Vec::new(), func, param_list: Vec::new()}, },
<func:ExtendedID> "," "(" <param_list:Comma<CallOperand>> ")" => (Vec::new(), func, param_list),
<func:ExtendedID> => (Vec::new(), func, Vec::<ast::CallOperand<_>>::new()),
}; };
OptionalDst: &'input str = { OptionalDst: &'input str = {

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

@ -4,35 +4,43 @@
OpCapability Kernel OpCapability Kernel
OpCapability Int64 OpCapability Int64
OpCapability Int8 OpCapability Int8
%1 = OpExtInstImport "OpenCL.std" %25 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %5 "add" OpEntryPoint Kernel %1 "add"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%4 = 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
%ulong_1 = OpConstant %ulong 1 %ulong_1 = OpConstant %ulong 1
%5 = OpFunction %void None %4 %1 = OpFunction %void None %28
%6 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%7 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%21 = OpLabel %23 = OpLabel
%8 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%9 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%10 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%11 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
OpStore %8 %6 %6 = OpVariable %_ptr_Function_ulong Function
OpStore %9 %7 %7 = OpVariable %_ptr_Function_ulong Function
%12 = OpLoad %ulong %8 OpStore %2 %8
%19 = OpConvertUToPtr %_ptr_Generic_ulong %12 OpStore %3 %9
%13 = OpLoad %ulong %19 %11 = OpLoad %ulong %2
OpStore %10 %13 %10 = OpCopyObject %ulong %11
%14 = OpLoad %ulong %10 OpStore %4 %10
%15 = OpIAdd %ulong %14 %ulong_1 %13 = OpLoad %ulong %3
OpStore %11 %15 %12 = OpCopyObject %ulong %13
%16 = OpLoad %ulong %9 OpStore %5 %12
%17 = OpLoad %ulong %11 %15 = OpLoad %ulong %4
%20 = OpConvertUToPtr %_ptr_Generic_ulong %16 %21 = OpConvertUToPtr %_ptr_Generic_ulong %15
OpStore %20 %17 %14 = OpLoad %ulong %21
OpStore %6 %14
%17 = OpLoad %ulong %6
%16 = OpIAdd %ulong %17 %ulong_1
OpStore %7 %16
%18 = OpLoad %ulong %5
%19 = OpLoad %ulong %7
%22 = OpConvertUToPtr %_ptr_Generic_ulong %18
OpStore %22 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

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

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

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

@ -4,46 +4,53 @@
OpCapability Kernel OpCapability Kernel
OpCapability Int64 OpCapability Int64
OpCapability Int8 OpCapability Int8
%1 = OpExtInstImport "OpenCL.std" %31 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %5 "bra" OpEntryPoint Kernel %1 "bra"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%4 = 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_1 = OpConstant %ulong 1 %ulong_1 = OpConstant %ulong 1
%ulong_2 = OpConstant %ulong 2 %ulong_2 = OpConstant %ulong 2
%5 = OpFunction %void None %4 %1 = OpFunction %void None %34
%6 = OpFunctionParameter %ulong %11 = OpFunctionParameter %ulong
%7 = OpFunctionParameter %ulong %12 = OpFunctionParameter %ulong
%27 = OpLabel %29 = OpLabel
%11 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%12 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%13 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function
%14 = OpVariable %_ptr_Function_ulong Function %8 = OpVariable %_ptr_Function_ulong Function
OpStore %11 %6 %9 = OpVariable %_ptr_Function_ulong Function
OpStore %12 %7 %10 = OpVariable %_ptr_Function_ulong Function
%16 = OpLoad %ulong %11 OpStore %2 %11
%25 = OpConvertUToPtr %_ptr_Generic_ulong %16 OpStore %3 %12
%15 = OpLoad %ulong %25 %14 = OpLoad %ulong %2
OpStore %13 %15 %13 = OpCopyObject %ulong %14
OpBranch %8 OpStore %7 %13
%8 = OpLabel %16 = OpLoad %ulong %3
%18 = OpLoad %ulong %13 %15 = OpCopyObject %ulong %16
%17 = OpIAdd %ulong %18 %ulong_1 OpStore %8 %15
OpStore %14 %17 %18 = OpLoad %ulong %7
OpBranch %10 %27 = OpConvertUToPtr %_ptr_Generic_ulong %18
%30 = OpLabel %17 = OpLoad %ulong %27
%20 = OpLoad %ulong %13 OpStore %9 %17
%19 = OpIAdd %ulong %20 %ulong_2 OpBranch %4
OpStore %14 %19 %4 = OpLabel
OpBranch %10 %20 = OpLoad %ulong %9
%10 = OpLabel %19 = OpIAdd %ulong %20 %ulong_1
%21 = OpLoad %ulong %12 OpStore %10 %19
%22 = OpLoad %ulong %14 OpBranch %6
%26 = OpConvertUToPtr %_ptr_Generic_ulong %21 %37 = OpLabel
OpStore %26 %22 %22 = OpLoad %ulong %9
%21 = OpIAdd %ulong %22 %ulong_2
OpStore %10 %21
OpBranch %6
%6 = OpLabel
%23 = OpLoad %ulong %8
%24 = OpLoad %ulong %10
%28 = OpConvertUToPtr %_ptr_Generic_ulong %23
OpStore %28 %24
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

38
ptx/src/test/spirv_run/call.ptx Normal file
View File

@ -0,0 +1,38 @@
.version 6.5
.target sm_30
.address_size 64
.func (.param.u64 output) incr (.param.u64 input);
.visible .entry call(
.param .u64 input,
.param .u64 output
)
{
.reg .u64 in_addr;
.reg .u64 out_addr;
.reg .u64 temp;
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
ld.global.u64 temp, [in_addr];
.param.u64 incr_in;
.param.u64 incr_out;
st.param.b64 [incr_in], temp;
call (incr_out), incr, (incr_in);
ld.param.u64 temp, [incr_out];
st.global.u64 [out_addr], temp;
ret;
}
.func (.param .u64 output) incr(
.param .u64 input
)
{
.reg .u64 temp;
ld.param.u64 temp, [input];
add.u64 temp, temp, 1;
st.param.u64 [output], temp;
ret;
}

73
ptx/src/test/spirv_run/call.spvtxt Normal file
View File

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

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

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

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

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

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

@ -1,38 +1,41 @@
; SPIR-V
; Version: 1.5
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 20
; Schema: 0
OpCapability GenericPointer OpCapability GenericPointer
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64 OpCapability Int64
OpCapability Int8 OpCapability Int8
%1 = OpExtInstImport "OpenCL.std" %21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "ld_st" OpEntryPoint Kernel %1 "ld_st"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%5 = 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 %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%2 = OpFunction %void None %5 %1 = OpFunction %void None %24
%7 = OpFunctionParameter %ulong
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %19 = OpLabel
%10 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function
%11 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%12 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%13 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
OpStore %11 %8 %6 = OpVariable %_ptr_Function_ulong Function
OpStore %12 %9 OpStore %2 %7
%14 = OpLoad %ulong %11 OpStore %3 %8
%15 = OpConvertUToPtr %_ptr_Generic_ulong %14 %10 = OpLoad %ulong %2
%16 = OpLoad %ulong %15 %9 = OpCopyObject %ulong %10
OpStore %13 %16 OpStore %4 %9
%17 = OpLoad %ulong %12 %12 = OpLoad %ulong %3
%18 = OpLoad %ulong %13 %11 = OpCopyObject %ulong %12
%19 = OpConvertUToPtr %_ptr_Generic_ulong %17 OpStore %5 %11
OpStore %19 %18 %14 = OpLoad %ulong %4
%17 = OpConvertUToPtr %_ptr_Generic_ulong %14
%13 = OpLoad %ulong %17
OpStore %6 %13
%15 = OpLoad %ulong %5
%16 = OpLoad %ulong %6
%18 = OpConvertUToPtr %_ptr_Generic_ulong %15
OpStore %18 %16
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

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

@ -4,35 +4,43 @@
OpCapability Kernel OpCapability Kernel
OpCapability Int64 OpCapability Int64
OpCapability Int8 OpCapability Int8
%1 = OpExtInstImport "OpenCL.std" %22 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %5 "local_align" OpEntryPoint Kernel %1 "local_align"
OpDecorate %8 Alignment 8 OpDecorate %4 Alignment 8
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%4 = OpTypeFunction %void %ulong %ulong %25 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%uchar = OpTypeInt 8 0 %uchar = OpTypeInt 8 0
%_arr_uchar_8 = OpTypeArray %uchar %8 %_arr_uchar_8 = OpTypeArray %uchar %8
%_ptr_Function__arr_uchar_8 = OpTypePointer Function %_arr_uchar_8 %_ptr_Function__arr_uchar_8 = OpTypePointer Function %_arr_uchar_8
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%5 = OpFunction %void None %4 %1 = OpFunction %void None %25
%6 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%7 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%18 = OpLabel %20 = OpLabel
%8 = OpVariable %_ptr_Function__arr_uchar_8 Workgroup %2 = OpVariable %_ptr_Function_ulong Function
%9 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%10 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function__arr_uchar_8 Workgroup
%11 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
OpStore %9 %6 %6 = OpVariable %_ptr_Function_ulong Function
OpStore %10 %7 %7 = OpVariable %_ptr_Function_ulong Function
%13 = OpLoad %ulong %9 OpStore %2 %8
%16 = OpConvertUToPtr %_ptr_Generic_ulong %13 OpStore %3 %9
%12 = OpLoad %ulong %16 %11 = OpLoad %ulong %2
OpStore %11 %12 %10 = OpCopyObject %ulong %11
%14 = OpLoad %ulong %10 OpStore %5 %10
%15 = OpLoad %ulong %11 %13 = OpLoad %ulong %3
%17 = OpConvertUToPtr %_ptr_Generic_ulong %14 %12 = OpCopyObject %ulong %13
OpStore %17 %15 OpStore %6 %12
%15 = OpLoad %ulong %5
%18 = OpConvertUToPtr %_ptr_Generic_ulong %15
%14 = OpLoad %ulong %18
OpStore %7 %14
%16 = OpLoad %ulong %6
%17 = OpLoad %ulong %7
%19 = OpConvertUToPtr %_ptr_Generic_ulong %16
OpStore %19 %17
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

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

@ -14,7 +14,7 @@ use std::fmt;
use std::fmt::{Debug, Display, Formatter}; use std::fmt::{Debug, Display, Formatter};
use std::mem; use std::mem;
use std::slice; use std::slice;
use std::{collections::HashMap, ptr, str}; use std::{borrow::Cow, collections::HashMap, env, fs, path::PathBuf, ptr, str};
macro_rules! test_ptx { macro_rules! test_ptx {
($fn_name:ident, $input:expr, $output:expr) => { ($fn_name:ident, $input:expr, $output:expr) => {
@ -32,8 +32,9 @@ macro_rules! test_ptx {
#[test] #[test]
fn [<$fn_name _spvtxt>]() -> Result<(), Box<dyn std::error::Error>> { fn [<$fn_name _spvtxt>]() -> Result<(), Box<dyn std::error::Error>> {
let ptx_txt = include_str!(concat!(stringify!($fn_name), ".ptx")); let ptx_txt = include_str!(concat!(stringify!($fn_name), ".ptx"));
let spirv_file_name = concat!(stringify!($fn_name), ".spvtxt");
let spirv_txt = include_bytes!(concat!(stringify!($fn_name), ".spvtxt")); let spirv_txt = include_bytes!(concat!(stringify!($fn_name), ".spvtxt"));
test_spvtxt_assert(ptx_txt, spirv_txt) test_spvtxt_assert(ptx_txt, spirv_txt, spirv_file_name)
} }
} }
}; };
@ -140,6 +141,7 @@ fn run_spirv<T: From<u8> + ze::SafeRepr + Copy + Debug>(
fn test_spvtxt_assert<'a>( fn test_spvtxt_assert<'a>(
ptx_txt: &'a str, ptx_txt: &'a str,
spirv_txt: &'a [u8], spirv_txt: &'a [u8],
spirv_file_name: &'a str,
) -> Result<(), Box<dyn error::Error + 'a>> { ) -> Result<(), Box<dyn error::Error + 'a>> {
let mut errors = Vec::new(); let mut errors = Vec::new();
let ast = ptx::ModuleParser::new().parse(&mut errors, ptx_txt)?; let ast = ptx::ModuleParser::new().parse(&mut errors, ptx_txt)?;
@ -191,16 +193,27 @@ fn test_spvtxt_assert<'a>(
) )
}; };
unsafe { spirv_tools::spvContextDestroy(spv_context) }; unsafe { spirv_tools::spvContextDestroy(spv_context) };
if result == spv_result_t::SPV_SUCCESS { let spirv_text = if result == spv_result_t::SPV_SUCCESS {
let raw_text = unsafe { let raw_text = unsafe {
std::slice::from_raw_parts((*spv_text).str_ as *const u8, (*spv_text).length) std::slice::from_raw_parts((*spv_text).str_ as *const u8, (*spv_text).length)
}; };
let spv_from_ptx_text = unsafe { str::from_utf8_unchecked(raw_text) }; let spv_from_ptx_text = unsafe { str::from_utf8_unchecked(raw_text) };
// TODO: stop leaking kernel text // TODO: stop leaking kernel text
panic!(spv_from_ptx_text); Cow::Borrowed(spv_from_ptx_text)
} else { } else {
panic!(ptx_mod.disassemble()); Cow::Owned(ptx_mod.disassemble())
};
if let Ok(dump_path) = env::var("NOTCUDA_TEST_SPIRV_DUMP_DIR") {
let mut path = PathBuf::from(dump_path);
if let Ok(()) = fs::create_dir_all(&path) {
path.push(spirv_file_name);
#[allow(unused_must_use)]
{
fs::write(path, spirv_text.as_bytes());
}
}
} }
panic!(spirv_text);
} }
unsafe { spirv_tools::spvContextDestroy(spv_context) }; unsafe { spirv_tools::spvContextDestroy(spv_context) };
Ok(()) Ok(())

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

@ -1,43 +1,45 @@
; SPIR-V
; Version: 1.5
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 23
; Schema: 0
OpCapability GenericPointer OpCapability GenericPointer
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64 OpCapability Int64
OpCapability Int8 OpCapability Int8
%1 = OpExtInstImport "OpenCL.std" %24 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "mov" OpEntryPoint Kernel %1 "mov"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%5 = 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
%2 = OpFunction %void None %5 %1 = OpFunction %void None %27
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%10 = OpLabel %22 = OpLabel
%11 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%12 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%13 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%14 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
OpStore %11 %8 %6 = OpVariable %_ptr_Function_ulong Function
OpStore %12 %9 %7 = OpVariable %_ptr_Function_ulong Function
%15 = OpLoad %ulong %11 OpStore %2 %8
%16 = OpConvertUToPtr %_ptr_Generic_ulong %15 OpStore %3 %9
%17 = OpLoad %ulong %16 %11 = OpLoad %ulong %2
OpStore %13 %17 %10 = OpCopyObject %ulong %11
%18 = OpLoad %ulong %13 OpStore %4 %10
%19 = OpCopyObject %ulong %18 %13 = OpLoad %ulong %3
OpStore %14 %19 %12 = OpCopyObject %ulong %13
%20 = OpLoad %ulong %12 OpStore %5 %12
%21 = OpLoad %ulong %14 %15 = OpLoad %ulong %4
%22 = OpConvertUToPtr %_ptr_Generic_ulong %20 %20 = OpConvertUToPtr %_ptr_Generic_ulong %15
OpStore %22 %21 %14 = OpLoad %ulong %20
OpStore %6 %14
%17 = OpLoad %ulong %6
%16 = OpCopyObject %ulong %17
OpStore %7 %16
%18 = OpLoad %ulong %5
%19 = OpLoad %ulong %7
%21 = OpConvertUToPtr %_ptr_Generic_ulong %18
OpStore %21 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

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

@ -1,43 +1,46 @@
; SPIR-V
; Version: 1.5
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 24
; Schema: 0
OpCapability GenericPointer OpCapability GenericPointer
OpCapability Linkage OpCapability Linkage
OpCapability Addresses OpCapability Addresses
OpCapability Kernel OpCapability Kernel
OpCapability Int64 OpCapability Int64
OpCapability Int8 OpCapability Int8
%1 = OpExtInstImport "OpenCL.std" %25 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "mul_hi" OpEntryPoint Kernel %1 "mul_hi"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%5 = 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
%ulong_2 = OpConstant %ulong 2 %ulong_2 = OpConstant %ulong 2
%2 = OpFunction %void None %5 %1 = OpFunction %void None %28
%8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong %23 = OpLabel
%11 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function
%12 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%13 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%14 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
%15 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function
OpStore %12 %9 %7 = OpVariable %_ptr_Function_ulong Function
OpStore %13 %10 OpStore %2 %8
%16 = OpLoad %ulong %12 OpStore %3 %9
%17 = OpConvertUToPtr %_ptr_Generic_ulong %16 %11 = OpLoad %ulong %2
%18 = OpLoad %ulong %17 %10 = OpCopyObject %ulong %11
OpStore %14 %18 OpStore %4 %10
%19 = OpLoad %ulong %14 %13 = OpLoad %ulong %3
%20 = OpExtInst %ulong %1 u_mul_hi %19 %ulong_2 %12 = OpCopyObject %ulong %13
OpStore %15 %20 OpStore %5 %12
%21 = OpLoad %ulong %13 %15 = OpLoad %ulong %4
%22 = OpLoad %ulong %15 %21 = OpConvertUToPtr %_ptr_Generic_ulong %15
%23 = OpConvertUToPtr %_ptr_Generic_ulong %21 %14 = OpLoad %ulong %21
OpStore %23 %22 OpStore %6 %14
%17 = OpLoad %ulong %6
%16 = OpExtInst %ulong %25 u_mul_hi %17 %ulong_2
OpStore %7 %16
%18 = OpLoad %ulong %5
%19 = OpLoad %ulong %7
%22 = OpConvertUToPtr %_ptr_Generic_ulong %18
OpStore %22 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

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

@ -4,35 +4,43 @@
OpCapability Kernel OpCapability Kernel
OpCapability Int64 OpCapability Int64
OpCapability Int8 OpCapability Int8
%1 = OpExtInstImport "OpenCL.std" %25 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %5 "mul_lo" OpEntryPoint Kernel %1 "mul_lo"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%4 = 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
%ulong_2 = OpConstant %ulong 2 %ulong_2 = OpConstant %ulong 2
%5 = OpFunction %void None %4 %1 = OpFunction %void None %28
%6 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%7 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%21 = OpLabel %23 = OpLabel
%8 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%9 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%10 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%11 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
OpStore %8 %6 %6 = OpVariable %_ptr_Function_ulong Function
OpStore %9 %7 %7 = OpVariable %_ptr_Function_ulong Function
%12 = OpLoad %ulong %8 OpStore %2 %8
%19 = OpConvertUToPtr %_ptr_Generic_ulong %12 OpStore %3 %9
%13 = OpLoad %ulong %19 %11 = OpLoad %ulong %2
OpStore %10 %13 %10 = OpCopyObject %ulong %11
%14 = OpLoad %ulong %10 OpStore %4 %10
%15 = OpIMul %ulong %14 %ulong_2 %13 = OpLoad %ulong %3
OpStore %11 %15 %12 = OpCopyObject %ulong %13
%16 = OpLoad %ulong %9 OpStore %5 %12
%17 = OpLoad %ulong %11 %15 = OpLoad %ulong %4
%20 = OpConvertUToPtr %_ptr_Generic_ulong %16 %21 = OpConvertUToPtr %_ptr_Generic_ulong %15
OpStore %20 %17 %14 = OpLoad %ulong %21
OpStore %6 %14
%17 = OpLoad %ulong %6
%16 = OpIMul %ulong %17 %ulong_2
OpStore %7 %16
%18 = OpLoad %ulong %5
%19 = OpLoad %ulong %7
%22 = OpConvertUToPtr %_ptr_Generic_ulong %18
OpStore %22 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

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

@ -4,34 +4,42 @@
OpCapability Kernel OpCapability Kernel
OpCapability Int64 OpCapability Int64
OpCapability Int8 OpCapability Int8
%1 = OpExtInstImport "OpenCL.std" %24 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %5 "not" OpEntryPoint Kernel %1 "not"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%4 = 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
%5 = OpFunction %void None %4 %1 = OpFunction %void None %27
%6 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%7 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%20 = OpLabel %22 = OpLabel
%8 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%9 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%10 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%11 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
OpStore %8 %6 %6 = OpVariable %_ptr_Function_ulong Function
OpStore %9 %7 %7 = OpVariable %_ptr_Function_ulong Function
%13 = OpLoad %ulong %8 OpStore %2 %8
%18 = OpConvertUToPtr %_ptr_Generic_ulong %13 OpStore %3 %9
%12 = OpLoad %ulong %18 %11 = OpLoad %ulong %2
OpStore %10 %12 %10 = OpCopyObject %ulong %11
%15 = OpLoad %ulong %10 OpStore %4 %10
%14 = OpNot %ulong %15 %13 = OpLoad %ulong %3
OpStore %11 %14 %12 = OpCopyObject %ulong %13
%16 = OpLoad %ulong %9 OpStore %5 %12
%17 = OpLoad %ulong %11 %15 = OpLoad %ulong %4
%19 = OpConvertUToPtr %_ptr_Generic_ulong %16 %20 = OpConvertUToPtr %_ptr_Generic_ulong %15
OpStore %19 %17 %14 = OpLoad %ulong %20
OpStore %6 %14
%17 = OpLoad %ulong %6
%16 = OpNot %ulong %17
OpStore %7 %16
%18 = OpLoad %ulong %5
%19 = OpLoad %ulong %7
%21 = OpConvertUToPtr %_ptr_Generic_ulong %18
OpStore %21 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

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

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

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

@ -4,36 +4,44 @@
OpCapability Kernel OpCapability Kernel
OpCapability Int64 OpCapability Int64
OpCapability Int8 OpCapability Int8
%1 = OpExtInstImport "OpenCL.std" %25 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %5 "shl" OpEntryPoint Kernel %1 "shl"
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%4 = 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
%5 = OpFunction %void None %4 %1 = OpFunction %void None %28
%6 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%7 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%21 = OpLabel %23 = OpLabel
%8 = OpVariable %_ptr_Function_ulong Function %2 = OpVariable %_ptr_Function_ulong Function
%9 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%10 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
%11 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function
OpStore %8 %6 %6 = OpVariable %_ptr_Function_ulong Function
OpStore %9 %7 %7 = OpVariable %_ptr_Function_ulong Function
%13 = OpLoad %ulong %8 OpStore %2 %8
%19 = OpConvertUToPtr %_ptr_Generic_ulong %13 OpStore %3 %9
%12 = OpLoad %ulong %19 %11 = OpLoad %ulong %2
OpStore %10 %12 %10 = OpCopyObject %ulong %11
%15 = OpLoad %ulong %10 OpStore %4 %10
%14 = OpShiftLeftLogical %ulong %15 %uint_2 %13 = OpLoad %ulong %3
OpStore %11 %14 %12 = OpCopyObject %ulong %13
%16 = OpLoad %ulong %9 OpStore %5 %12
%17 = OpLoad %ulong %11 %15 = OpLoad %ulong %4
%20 = OpConvertUToPtr %_ptr_Generic_ulong %16 %21 = OpConvertUToPtr %_ptr_Generic_ulong %15
OpStore %20 %17 %14 = OpLoad %ulong %21
OpStore %6 %14
%17 = OpLoad %ulong %6
%16 = OpShiftLeftLogical %ulong %17 %uint_2
OpStore %7 %16
%18 = OpLoad %ulong %5
%19 = OpLoad %ulong %7
%22 = OpConvertUToPtr %_ptr_Generic_ulong %18
OpStore %22 %19
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

1131
ptx/src/translate.rs
View File

File diff suppressed because it is too large Load Diff