mirrors/ZLUDA
1
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-04-20 08:24:44 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Implement stateless-to-stateful optimization

Browse source
This commit is contained in:
Andrzej Janik 2020-11-19 22:12:12 +01:00
parent eac5fbd806
commit f77b653d36
16 changed files with 1496 additions and 488 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

6
ptx/src/ast.rs
View file

@ -766,6 +766,8 @@ sub_type! {
LdStType {
Scalar(LdStScalarType),
Vector(LdStScalarType, u8),
// Used in generated code
Pointer(PointerType, LdStateSpace),
}
}
@ -774,6 +776,10 @@ impl From<LdStType> for PointerType {
match t {
LdStType::Scalar(t) => PointerType::Scalar(t.into()),
LdStType::Vector(t, len) => PointerType::Vector(t.into(), len),
LdStType::Pointer(PointerType::Scalar(scalar_type), space) => {
PointerType::Pointer(scalar_type, space)
}
LdStType::Pointer(..) => unreachable!(),
}
}
}

12
ptx/src/ptx.lalrpop
View file

@ -1237,18 +1237,18 @@ InstRet: ast::Instruction<ast::ParsedArgParams<'input>> = {
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvta
InstCvta: ast::Instruction<ast::ParsedArgParams<'input>> = {
"cvta" <to:CvtaStateSpace> <s:CvtaSize> <a:Arg2> => {
"cvta" <from:CvtaStateSpace> <s:CvtaSize> <a:Arg2> => {
ast::Instruction::Cvta(ast::CvtaDetails {
to: to,
from: ast::CvtaStateSpace::Generic,
to: ast::CvtaStateSpace::Generic,
from,
size: s
},
a)
},
"cvta" ".to" <from:CvtaStateSpace> <s:CvtaSize> <a:Arg2> => {
"cvta" ".to" <to:CvtaStateSpace> <s:CvtaSize> <a:Arg2> => {
ast::Instruction::Cvta(ast::CvtaDetails {
to: ast::CvtaStateSpace::Generic,
from: from,
to,
from: ast::CvtaStateSpace::Generic,
size: s
},
a)

170
ptx/src/test/spirv_run/atom_inc.spvtxt
View file

@ -1,89 +1,81 @@
; SPIR-V
; Version: 1.3
; Generator: rspirv
; Bound: 60
OpCapability GenericPointer
OpCapability Linkage
OpCapability Addresses
OpCapability Kernel
OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
; OpCapability FunctionFloatControlINTEL
; OpExtension "SPV_INTEL_float_controls2"
%49 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "atom_inc"
OpDecorate %40 LinkageAttributes "__notcuda_ptx_impl__atom_relaxed_gpu_generic_inc" Import
OpDecorate %44 LinkageAttributes "__notcuda_ptx_impl__atom_relaxed_gpu_global_inc" Import
%50 = OpTypeVoid
%51 = OpTypeInt 32 0
%52 = OpTypePointer Generic %51
%53 = OpTypeFunction %51 %52 %51
%54 = OpTypePointer CrossWorkgroup %51
%55 = OpTypeFunction %51 %54 %51
%56 = OpTypeInt 64 0
%57 = OpTypeFunction %50 %56 %56
%58 = OpTypePointer Function %56
%59 = OpTypePointer Function %51
%27 = OpConstant %51 101
%28 = OpConstant %51 101
%29 = OpConstant %56 4
%31 = OpConstant %56 8
%40 = OpFunction %51 None %53
%42 = OpFunctionParameter %52
%43 = OpFunctionParameter %51
OpFunctionEnd
%44 = OpFunction %51 None %55
%46 = OpFunctionParameter %54
%47 = OpFunctionParameter %51
OpFunctionEnd
%1 = OpFunction %50 None %57
%9 = OpFunctionParameter %56
%10 = OpFunctionParameter %56
%39 = OpLabel
%2 = OpVariable %58 Function
%3 = OpVariable %58 Function
%4 = OpVariable %58 Function
%5 = OpVariable %58 Function
%6 = OpVariable %59 Function
%7 = OpVariable %59 Function
%8 = OpVariable %59 Function
OpStore %2 %9
OpStore %3 %10
%12 = OpLoad %56 %2
%11 = OpCopyObject %56 %12
OpStore %4 %11
%14 = OpLoad %56 %3
%13 = OpCopyObject %56 %14
OpStore %5 %13
%16 = OpLoad %56 %4
%33 = OpConvertUToPtr %52 %16
%15 = OpFunctionCall %51 %40 %33 %27
OpStore %6 %15
%18 = OpLoad %56 %4
%34 = OpConvertUToPtr %54 %18
%17 = OpFunctionCall %51 %44 %34 %28
OpStore %7 %17
%20 = OpLoad %56 %4
%35 = OpConvertUToPtr %52 %20
%19 = OpLoad %51 %35
OpStore %8 %19
%21 = OpLoad %56 %5
%22 = OpLoad %51 %6
%36 = OpConvertUToPtr %52 %21
OpStore %36 %22
%23 = OpLoad %56 %5
%24 = OpLoad %51 %7
%30 = OpIAdd %56 %23 %29
%37 = OpConvertUToPtr %52 %30
OpStore %37 %24
%25 = OpLoad %56 %5
%26 = OpLoad %51 %8
%32 = OpIAdd %56 %25 %31
%38 = OpConvertUToPtr %52 %32
OpStore %38 %26
OpReturn
OpFunctionEnd
OpCapability GenericPointer
OpCapability Linkage
OpCapability Addresses
OpCapability Kernel
OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%47 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "atom_inc"
OpDecorate %38 LinkageAttributes "__notcuda_ptx_impl__atom_relaxed_gpu_generic_inc" Import
OpDecorate %42 LinkageAttributes "__notcuda_ptx_impl__atom_relaxed_gpu_global_inc" Import
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%_ptr_Generic_uint = OpTypePointer Generic %uint
%51 = OpTypeFunction %uint %_ptr_Generic_uint %uint
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%53 = OpTypeFunction %uint %_ptr_CrossWorkgroup_uint %uint
%ulong = OpTypeInt 64 0
%55 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Function_uint = OpTypePointer Function %uint
%uint_101 = OpConstant %uint 101
%uint_101_0 = OpConstant %uint 101
%ulong_4 = OpConstant %ulong 4
%ulong_8 = OpConstant %ulong 8
%38 = OpFunction %uint None %51
%40 = OpFunctionParameter %_ptr_Generic_uint
%41 = OpFunctionParameter %uint
OpFunctionEnd
%42 = OpFunction %uint None %53
%44 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%45 = OpFunctionParameter %uint
OpFunctionEnd
%1 = OpFunction %void None %55
%9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong
%37 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function
%5 = OpVariable %_ptr_Function_ulong Function
%6 = OpVariable %_ptr_Function_uint Function
%7 = OpVariable %_ptr_Function_uint Function
%8 = OpVariable %_ptr_Function_uint Function
OpStore %2 %9
OpStore %3 %10
%11 = OpLoad %ulong %2
OpStore %4 %11
%12 = OpLoad %ulong %3
OpStore %5 %12
%14 = OpLoad %ulong %4
%31 = OpConvertUToPtr %_ptr_Generic_uint %14
%13 = OpFunctionCall %uint %38 %31 %uint_101
OpStore %6 %13
%16 = OpLoad %ulong %4
%32 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %16
%15 = OpFunctionCall %uint %42 %32 %uint_101_0
OpStore %7 %15
%18 = OpLoad %ulong %4
%33 = OpConvertUToPtr %_ptr_Generic_uint %18
%17 = OpLoad %uint %33
OpStore %8 %17
%19 = OpLoad %ulong %5
%20 = OpLoad %uint %6
%34 = OpConvertUToPtr %_ptr_Generic_uint %19
OpStore %34 %20
%21 = OpLoad %ulong %5
%22 = OpLoad %uint %7
%28 = OpIAdd %ulong %21 %ulong_4
%35 = OpConvertUToPtr %_ptr_Generic_uint %28
OpStore %35 %22
%23 = OpLoad %ulong %5
%24 = OpLoad %uint %8
%30 = OpIAdd %ulong %23 %ulong_8
%36 = OpConvertUToPtr %_ptr_Generic_uint %30
OpStore %36 %24
OpReturn
OpFunctionEnd

83
ptx/src/test/spirv_run/cvta.spvtxt
View file

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

84
ptx/src/test/spirv_run/extern_shared_call.spvtxt
View file

@ -7,7 +7,7 @@
OpCapability Int64
OpCapability Float16
OpCapability Float64
%48 = OpExtInstImport "OpenCL.std"
%46 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %14 "extern_shared_call" %1
OpDecorate %1 Alignment 4
@ -18,78 +18,76 @@
%1 = OpVariable %_ptr_Workgroup__ptr_Workgroup_uint Workgroup
%uchar = OpTypeInt 8 0
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
%55 = OpTypeFunction %void %_ptr_Workgroup_uchar
%53 = OpTypeFunction %void %_ptr_Workgroup_uchar
%_ptr_Function__ptr_Workgroup_uchar = OpTypePointer Function %_ptr_Workgroup_uchar
%ulong = OpTypeInt 64 0
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Function__ptr_Workgroup_uint = OpTypePointer Function %_ptr_Workgroup_uint
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%ulong_2 = OpConstant %ulong 2
%62 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar
%60 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%2 = OpFunction %void None %55
%40 = OpFunctionParameter %_ptr_Workgroup_uchar
%56 = OpLabel
%41 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function
%2 = OpFunction %void None %53
%38 = OpFunctionParameter %_ptr_Workgroup_uchar
%54 = OpLabel
%39 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function
%3 = OpVariable %_ptr_Function_ulong Function
OpStore %41 %40
OpStore %39 %38
OpBranch %13
%13 = OpLabel
%42 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %41
%5 = OpLoad %_ptr_Workgroup_uint %42
%40 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %39
%5 = OpLoad %_ptr_Workgroup_uint %40
%11 = OpBitcast %_ptr_Workgroup_ulong %5
%4 = OpLoad %ulong %11
OpStore %3 %4
%7 = OpLoad %ulong %3
%6 = OpIAdd %ulong %7 %ulong_2
OpStore %3 %6
%43 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %41
%8 = OpLoad %_ptr_Workgroup_uint %43
%41 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %39
%8 = OpLoad %_ptr_Workgroup_uint %41
%9 = OpLoad %ulong %3
%12 = OpBitcast %_ptr_Workgroup_ulong %8
OpStore %12 %9
OpReturn
OpFunctionEnd
%14 = OpFunction %void None %62
%14 = OpFunction %void None %60
%20 = OpFunctionParameter %ulong
%21 = OpFunctionParameter %ulong
%44 = OpFunctionParameter %_ptr_Workgroup_uchar
%63 = OpLabel
%45 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function
%42 = OpFunctionParameter %_ptr_Workgroup_uchar
%61 = OpLabel
%43 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function
%15 = OpVariable %_ptr_Function_ulong Function
%16 = OpVariable %_ptr_Function_ulong Function
%17 = OpVariable %_ptr_Function_ulong Function
%18 = OpVariable %_ptr_Function_ulong Function
%19 = OpVariable %_ptr_Function_ulong Function
OpStore %45 %44
OpBranch %38
%38 = OpLabel
OpStore %43 %42
OpBranch %36
%36 = OpLabel
OpStore %15 %20
OpStore %16 %21
%23 = OpLoad %ulong %15
%22 = OpCopyObject %ulong %23
%22 = OpLoad %ulong %15
OpStore %17 %22
%25 = OpLoad %ulong %16
%24 = OpCopyObject %ulong %25
OpStore %18 %24
%27 = OpLoad %ulong %17
%34 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %27
%26 = OpLoad %ulong %34
OpStore %19 %26
%46 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %45
%28 = OpLoad %_ptr_Workgroup_uint %46
%29 = OpLoad %ulong %19
%35 = OpBitcast %_ptr_Workgroup_ulong %28
OpStore %35 %29
%65 = OpFunctionCall %void %2 %44
%47 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %45
%31 = OpLoad %_ptr_Workgroup_uint %47
%36 = OpBitcast %_ptr_Workgroup_ulong %31
%30 = OpLoad %ulong %36
OpStore %19 %30
%32 = OpLoad %ulong %18
%33 = OpLoad %ulong %19
%37 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %32
OpStore %37 %33
%23 = OpLoad %ulong %16
OpStore %18 %23
%25 = OpLoad %ulong %17
%32 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %25
%24 = OpLoad %ulong %32
OpStore %19 %24
%44 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %43
%26 = OpLoad %_ptr_Workgroup_uint %44
%27 = OpLoad %ulong %19
%33 = OpBitcast %_ptr_Workgroup_ulong %26
OpStore %33 %27
%63 = OpFunctionCall %void %2 %42
%45 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %43
%29 = OpLoad %_ptr_Workgroup_uint %45
%34 = OpBitcast %_ptr_Workgroup_ulong %29
%28 = OpLoad %ulong %34
OpStore %19 %28
%30 = OpLoad %ulong %18
%31 = OpLoad %ulong %19
%35 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %30
OpStore %35 %31
OpReturn
OpFunctionEnd

18
ptx/src/test/spirv_run/mod.rs
View file

@ -133,6 +133,10 @@ test_ptx!(
[0b11111000_11000001_00100010_10100000u32, 16u32, 8u32],
[0b11000001u32]
);
test_ptx!(stateful_ld_st_simple, [121u64], [121u64]);
test_ptx!(stateful_ld_st_ntid, [123u64], [123u64]);
test_ptx!(stateful_ld_st_ntid_chain, [12651u64], [12651u64]);
test_ptx!(stateful_ld_st_ntid_sub, [96311u64], [96311u64]);
struct DisplayError<T: Debug> {
err: T,
@ -292,7 +296,7 @@ fn test_spvtxt_assert<'a>(
rspirv::binary::parse_words(&parsed_spirv, &mut loader)?;
let spvtxt_mod = loader.module();
unsafe { spirv_tools::spvBinaryDestroy(spv_binary) };
if !is_spirv_fn_equal(&spirv_module.spirv.functions[0], &spvtxt_mod.functions[0]) {
if !is_spirv_fns_equal(&spirv_module.spirv.functions, &spvtxt_mod.functions) {
// We could simply use ptx_mod.disassemble, but SPIRV-Tools text formattinmg is so much nicer
let spv_from_ptx_binary = spirv_module.spirv.assemble();
let mut spv_text: spirv_tools::spv_text = ptr::null_mut();
@ -364,6 +368,18 @@ impl<T: Copy + Eq + Hash> EqMap<T> {
}
}
fn is_spirv_fns_equal(fns1: &[Function], fns2: &[Function]) -> bool {
if fns1.len() != fns2.len() {
return false;
}
for (fn1, fn2) in fns1.iter().zip(fns2.iter()) {
if !is_spirv_fn_equal(fn1, fn2) {
return false;
}
}
true
}
fn is_spirv_fn_equal(fn1: &Function, fn2: &Function) -> bool {
let mut map = EqMap::new();
if !is_option_equal(&fn1.def, &fn2.def, &mut map, is_instr_equal) {

12
ptx/src/test/spirv_run/reg_local.spvtxt
View file

@ -22,7 +22,9 @@
%_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%ulong_1 = OpConstant %ulong 1
%_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_0 = OpConstant %ulong 0
%_ptr_Generic_uchar = OpTypePointer Generic %uchar
%ulong_0_0 = OpConstant %ulong 0
%1 = OpFunction %void None %37
%8 = OpFunctionParameter %ulong
@ -48,12 +50,12 @@
%14 = OpLoad %ulong %7
%26 = OpCopyObject %ulong %14
%19 = OpIAdd %ulong %26 %ulong_1
%27 = OpBitcast %_ptr_Function_ulong %4
%27 = OpBitcast %_ptr_Generic_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
%28 = OpBitcast %_ptr_Generic_ulong %4
%47 = OpBitcast %_ptr_Generic_uchar %28
%48 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %47 %ulong_0
%21 = OpBitcast %_ptr_Generic_ulong %48
%29 = OpLoad %ulong %21
%15 = OpCopyObject %ulong %29
OpStore %7 %15

31
ptx/src/test/spirv_run/stateful_ld_st_ntid.ptx Normal file
View file

@ -0,0 +1,31 @@
.version 6.5
.target sm_30
.address_size 64
.visible .entry stateful_ld_st_ntid(
.param .u64 input,
.param .u64 output
)
{
.reg .b64 in_addr;
.reg .b64 out_addr;
.reg .u32 tid_32;
.reg .u64 tid_64;
.reg .u64 temp;
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
cvta.to.global.u64 in_addr, in_addr;
cvta.to.global.u64 out_addr, out_addr;
mov.u32 tid_32, %tid.x;
cvt.u64.u32 tid_64, tid_32;
add.u64 in_addr, in_addr, tid_64;
add.u64 out_addr, out_addr, tid_64;
ld.global.u64 temp, [in_addr];
st.global.u64 [out_addr], temp;
ret;
}

89
ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt Normal file
View file

@ -0,0 +1,89 @@
OpCapability GenericPointer
OpCapability Linkage
OpCapability Addresses
OpCapability Kernel
OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%49 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "stateful_ld_st_ntid" %gl_LocalInvocationID
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%v4uint = OpTypeVector %uint 4
%_ptr_Input_v4uint = OpTypePointer Input %v4uint
%gl_LocalInvocationID = OpVariable %_ptr_Input_v4uint Input
%uchar = OpTypeInt 8 0
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%56 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
%_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
%_ptr_Function_uint = OpTypePointer Function %uint
%ulong = OpTypeInt 64 0
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%1 = OpFunction %void None %56
%20 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%21 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%47 = OpLabel
%2 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%3 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%10 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%11 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%6 = OpVariable %_ptr_Function_uint Function
%7 = OpVariable %_ptr_Function_ulong Function
%8 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %20
OpStore %3 %21
%13 = OpBitcast %_ptr_Function_ulong %2
%43 = OpLoad %ulong %13
%12 = OpCopyObject %ulong %43
%22 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %12
OpStore %10 %22
%15 = OpBitcast %_ptr_Function_ulong %3
%44 = OpLoad %ulong %15
%14 = OpCopyObject %ulong %44
%23 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %14
OpStore %11 %23
%24 = OpLoad %_ptr_CrossWorkgroup_uchar %10
%17 = OpConvertPtrToU %ulong %24
%16 = OpCopyObject %ulong %17
%25 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %16
OpStore %10 %25
%26 = OpLoad %_ptr_CrossWorkgroup_uchar %11
%19 = OpConvertPtrToU %ulong %26
%18 = OpCopyObject %ulong %19
%27 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %18
OpStore %11 %27
%29 = OpLoad %v4uint %gl_LocalInvocationID
%42 = OpCompositeExtract %uint %29 0
%28 = OpCopyObject %uint %42
OpStore %6 %28
%31 = OpLoad %uint %6
%61 = OpBitcast %uint %31
%30 = OpUConvert %ulong %61
OpStore %7 %30
%33 = OpLoad %_ptr_CrossWorkgroup_uchar %10
%34 = OpLoad %ulong %7
%62 = OpBitcast %_ptr_CrossWorkgroup_uchar %33
%63 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %62 %34
%32 = OpBitcast %_ptr_CrossWorkgroup_uchar %63
OpStore %10 %32
%36 = OpLoad %_ptr_CrossWorkgroup_uchar %11
%37 = OpLoad %ulong %7
%64 = OpBitcast %_ptr_CrossWorkgroup_uchar %36
%65 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %64 %37
%35 = OpBitcast %_ptr_CrossWorkgroup_uchar %65
OpStore %11 %35
%39 = OpLoad %_ptr_CrossWorkgroup_uchar %10
%45 = OpBitcast %_ptr_CrossWorkgroup_ulong %39
%38 = OpLoad %ulong %45
OpStore %8 %38
%40 = OpLoad %_ptr_CrossWorkgroup_uchar %11
%41 = OpLoad %ulong %8
%46 = OpBitcast %_ptr_CrossWorkgroup_ulong %40
OpStore %46 %41
OpReturn
OpFunctionEnd

35
ptx/src/test/spirv_run/stateful_ld_st_ntid_chain.ptx Normal file
View file

@ -0,0 +1,35 @@
.version 6.5
.target sm_30
.address_size 64
.visible .entry stateful_ld_st_ntid_chain(
.param .u64 input,
.param .u64 output
)
{
.reg .b64 in_addr1;
.reg .b64 in_addr2;
.reg .b64 in_addr3;
.reg .b64 out_addr1;
.reg .b64 out_addr2;
.reg .b64 out_addr3;
.reg .u32 tid_32;
.reg .u64 tid_64;
.reg .u64 temp;
ld.param.u64 in_addr1, [input];
ld.param.u64 out_addr1, [output];
cvta.to.global.u64 in_addr2, in_addr1;
cvta.to.global.u64 out_addr2, out_addr1;
mov.u32 tid_32, %tid.x;
cvt.u64.u32 tid_64, tid_32;
add.u64 in_addr3, in_addr2, tid_64;
add.u64 out_addr3, out_addr2, tid_64;
ld.global.u64 temp, [in_addr3];
st.global.u64 [out_addr3], temp;
ret;
}

93
ptx/src/test/spirv_run/stateful_ld_st_ntid_chain.spvtxt Normal file
View file

@ -0,0 +1,93 @@
OpCapability GenericPointer
OpCapability Linkage
OpCapability Addresses
OpCapability Kernel
OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%57 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "stateful_ld_st_ntid_chain" %gl_LocalInvocationID
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%v4uint = OpTypeVector %uint 4
%_ptr_Input_v4uint = OpTypePointer Input %v4uint
%gl_LocalInvocationID = OpVariable %_ptr_Input_v4uint Input
%uchar = OpTypeInt 8 0
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%64 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
%_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
%_ptr_Function_uint = OpTypePointer Function %uint
%ulong = OpTypeInt 64 0
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%1 = OpFunction %void None %64
%28 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%29 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%55 = OpLabel
%2 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%3 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%14 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%15 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%16 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%17 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%18 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%19 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%10 = OpVariable %_ptr_Function_uint Function
%11 = OpVariable %_ptr_Function_ulong Function
%12 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %28
OpStore %3 %29
%21 = OpBitcast %_ptr_Function_ulong %2
%51 = OpLoad %ulong %21
%20 = OpCopyObject %ulong %51
%30 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %20
OpStore %14 %30
%23 = OpBitcast %_ptr_Function_ulong %3
%52 = OpLoad %ulong %23
%22 = OpCopyObject %ulong %52
%31 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %22
OpStore %17 %31
%32 = OpLoad %_ptr_CrossWorkgroup_uchar %14
%25 = OpConvertPtrToU %ulong %32
%24 = OpCopyObject %ulong %25
%33 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %24
OpStore %15 %33
%34 = OpLoad %_ptr_CrossWorkgroup_uchar %17
%27 = OpConvertPtrToU %ulong %34
%26 = OpCopyObject %ulong %27
%35 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %26
OpStore %18 %35
%37 = OpLoad %v4uint %gl_LocalInvocationID
%50 = OpCompositeExtract %uint %37 0
%36 = OpCopyObject %uint %50
OpStore %10 %36
%39 = OpLoad %uint %10
%69 = OpBitcast %uint %39
%38 = OpUConvert %ulong %69
OpStore %11 %38
%41 = OpLoad %_ptr_CrossWorkgroup_uchar %15
%42 = OpLoad %ulong %11
%70 = OpBitcast %_ptr_CrossWorkgroup_uchar %41
%71 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %70 %42
%40 = OpBitcast %_ptr_CrossWorkgroup_uchar %71
OpStore %16 %40
%44 = OpLoad %_ptr_CrossWorkgroup_uchar %18
%45 = OpLoad %ulong %11
%72 = OpBitcast %_ptr_CrossWorkgroup_uchar %44
%73 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %72 %45
%43 = OpBitcast %_ptr_CrossWorkgroup_uchar %73
OpStore %19 %43
%47 = OpLoad %_ptr_CrossWorkgroup_uchar %16
%53 = OpBitcast %_ptr_CrossWorkgroup_ulong %47
%46 = OpLoad %ulong %53
OpStore %12 %46
%48 = OpLoad %_ptr_CrossWorkgroup_uchar %19
%49 = OpLoad %ulong %12
%54 = OpBitcast %_ptr_CrossWorkgroup_ulong %48
OpStore %54 %49
OpReturn
OpFunctionEnd

35
ptx/src/test/spirv_run/stateful_ld_st_ntid_sub.ptx Normal file
View file

@ -0,0 +1,35 @@
.version 6.5
.target sm_30
.address_size 64
.visible .entry stateful_ld_st_ntid_sub(
.param .u64 input,
.param .u64 output
)
{
.reg .b64 in_addr1;
.reg .b64 in_addr2;
.reg .b64 in_addr3;
.reg .b64 out_addr1;
.reg .b64 out_addr2;
.reg .b64 out_addr3;
.reg .u32 tid_32;
.reg .u64 tid_64;
.reg .u64 temp;
ld.param.u64 in_addr1, [input];
ld.param.u64 out_addr1, [output];
cvta.to.global.u64 in_addr2, in_addr1;
cvta.to.global.u64 out_addr2, out_addr1;
mov.u32 tid_32, %tid.x;
cvt.u64.u32 tid_64, tid_32;
sub.s64 in_addr3, in_addr2, tid_64;
sub.s64 out_addr3, out_addr2, tid_64;
ld.global.u64 temp, [in_addr3+-0];
st.global.u64 [out_addr3+-0], temp;
ret;
}

105
ptx/src/test/spirv_run/stateful_ld_st_ntid_sub.spvtxt Normal file
View file

@ -0,0 +1,105 @@
OpCapability GenericPointer
OpCapability Linkage
OpCapability Addresses
OpCapability Kernel
OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%65 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "stateful_ld_st_ntid_sub" %gl_LocalInvocationID
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%v4uint = OpTypeVector %uint 4
%_ptr_Input_v4uint = OpTypePointer Input %v4uint
%gl_LocalInvocationID = OpVariable %_ptr_Input_v4uint Input
%uchar = OpTypeInt 8 0
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%72 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
%_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
%_ptr_Function_uint = OpTypePointer Function %uint
%ulong = OpTypeInt 64 0
%_ptr_Function_ulong = OpTypePointer Function %ulong
%ulong_0 = OpConstant %ulong 0
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%ulong_0_0 = OpConstant %ulong 0
%1 = OpFunction %void None %72
%30 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%31 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%63 = OpLabel
%2 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%3 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%14 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%15 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%16 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%17 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%18 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%19 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%10 = OpVariable %_ptr_Function_uint Function
%11 = OpVariable %_ptr_Function_ulong Function
%12 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %30
OpStore %3 %31
%21 = OpBitcast %_ptr_Function_ulong %2
%57 = OpLoad %ulong %21
%20 = OpCopyObject %ulong %57
%32 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %20
OpStore %14 %32
%23 = OpBitcast %_ptr_Function_ulong %3
%58 = OpLoad %ulong %23
%22 = OpCopyObject %ulong %58
%33 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %22
OpStore %17 %33
%34 = OpLoad %_ptr_CrossWorkgroup_uchar %14
%25 = OpConvertPtrToU %ulong %34
%24 = OpCopyObject %ulong %25
%35 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %24
OpStore %15 %35
%36 = OpLoad %_ptr_CrossWorkgroup_uchar %17
%27 = OpConvertPtrToU %ulong %36
%26 = OpCopyObject %ulong %27
%37 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %26
OpStore %18 %37
%39 = OpLoad %v4uint %gl_LocalInvocationID
%52 = OpCompositeExtract %uint %39 0
%38 = OpCopyObject %uint %52
OpStore %10 %38
%41 = OpLoad %uint %10
%77 = OpBitcast %uint %41
%40 = OpUConvert %ulong %77
OpStore %11 %40
%42 = OpLoad %ulong %11
%59 = OpCopyObject %ulong %42
%28 = OpSNegate %ulong %59
%44 = OpLoad %_ptr_CrossWorkgroup_uchar %15
%78 = OpBitcast %_ptr_CrossWorkgroup_uchar %44
%79 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %78 %28
%43 = OpBitcast %_ptr_CrossWorkgroup_uchar %79
OpStore %16 %43
%45 = OpLoad %ulong %11
%60 = OpCopyObject %ulong %45
%29 = OpSNegate %ulong %60
%47 = OpLoad %_ptr_CrossWorkgroup_uchar %18
%80 = OpBitcast %_ptr_CrossWorkgroup_uchar %47
%81 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %80 %29
%46 = OpBitcast %_ptr_CrossWorkgroup_uchar %81
OpStore %19 %46
%49 = OpLoad %_ptr_CrossWorkgroup_uchar %16
%61 = OpBitcast %_ptr_CrossWorkgroup_ulong %49
%83 = OpBitcast %_ptr_CrossWorkgroup_uchar %61
%84 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %83 %ulong_0
%54 = OpBitcast %_ptr_CrossWorkgroup_ulong %84
%48 = OpLoad %ulong %54
OpStore %12 %48
%50 = OpLoad %_ptr_CrossWorkgroup_uchar %19
%51 = OpLoad %ulong %12
%62 = OpBitcast %_ptr_CrossWorkgroup_ulong %50
%85 = OpBitcast %_ptr_CrossWorkgroup_uchar %62
%86 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %85 %ulong_0_0
%56 = OpBitcast %_ptr_CrossWorkgroup_ulong %86
OpStore %56 %51
OpReturn
OpFunctionEnd

25
ptx/src/test/spirv_run/stateful_ld_st_simple.ptx Normal file
View file

@ -0,0 +1,25 @@
.version 6.5
.target sm_30
.address_size 64
.visible .entry stateful_ld_st_simple(
.param .u64 input,
.param .u64 output
)
{
.reg .u64 in_addr;
.reg .u64 out_addr;
.reg .u64 in_addr2;
.reg .u64 out_addr2;
.reg .u64 temp;
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
cvta.to.global.u64 in_addr2, in_addr;
cvta.to.global.u64 out_addr2, out_addr;
ld.global.u64 temp, [in_addr2];
st.global.u64 [out_addr2], temp;
ret;
}

65
ptx/src/test/spirv_run/stateful_ld_st_simple.spvtxt Normal file
View file

@ -0,0 +1,65 @@
OpCapability GenericPointer
OpCapability Linkage
OpCapability Addresses
OpCapability Kernel
OpCapability Int8
OpCapability Int16
OpCapability Int64
OpCapability Float16
OpCapability Float64
%41 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "stateful_ld_st_simple"
%void = OpTypeVoid
%uchar = OpTypeInt 8 0
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%45 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
%_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
%ulong = OpTypeInt 64 0
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%1 = OpFunction %void None %45
%21 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%22 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%39 = OpLabel
%2 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%3 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%9 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%10 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%11 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%12 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%8 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %21
OpStore %3 %22
%14 = OpBitcast %_ptr_Function_ulong %2
%13 = OpLoad %ulong %14
%23 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %13
OpStore %9 %23
%16 = OpBitcast %_ptr_Function_ulong %3
%15 = OpLoad %ulong %16
%24 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %15
OpStore %10 %24
%25 = OpLoad %_ptr_CrossWorkgroup_uchar %9
%18 = OpConvertPtrToU %ulong %25
%34 = OpCopyObject %ulong %18
%33 = OpCopyObject %ulong %34
%17 = OpCopyObject %ulong %33
%26 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %17
OpStore %11 %26
%27 = OpLoad %_ptr_CrossWorkgroup_uchar %10
%20 = OpConvertPtrToU %ulong %27
%36 = OpCopyObject %ulong %20
%35 = OpCopyObject %ulong %36
%19 = OpCopyObject %ulong %35
%28 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %19
OpStore %12 %28
%30 = OpLoad %_ptr_CrossWorkgroup_uchar %11
%37 = OpBitcast %_ptr_CrossWorkgroup_ulong %30
%29 = OpLoad %ulong %37
OpStore %8 %29
%31 = OpLoad %_ptr_CrossWorkgroup_uchar %12
%32 = OpLoad %ulong %8
%38 = OpBitcast %_ptr_CrossWorkgroup_ulong %31
OpStore %38 %32
OpReturn
OpFunctionEnd

1121
ptx/src/translate.rs
View file

File diff suppressed because it is too large Load diff