diff options
| -rw-r--r-- | ptx/src/test/spirv_run/extern_shared_call.ptx | 18 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/extern_shared_call.spvtxt | 124 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/mod.rs | 1 |
3 files changed, 88 insertions, 55 deletions
diff --git a/ptx/src/test/spirv_run/extern_shared_call.ptx b/ptx/src/test/spirv_run/extern_shared_call.ptx index 6626783..76b2ae7 100644 --- a/ptx/src/test/spirv_run/extern_shared_call.ptx +++ b/ptx/src/test/spirv_run/extern_shared_call.ptx @@ -4,28 +4,17 @@ .extern .shared .align 4 .b32 shared_mem[];
-.func (.param .u64 output) incr_shared_2_param(
- .param .u64 .ptr .shared shared_mem_addr
-)
-{
- .reg .u64 temp;
- ld.shared.u64 temp, [shared_mem_addr];
- add.u64 temp, temp, 2;
- st.param.u64 [output], temp;
- ret;
-}
-
-.func (.param .u64 output) incr_shared_2_global()
+.func incr_shared_2_global()
{
.reg .u64 temp;
ld.shared.u64 temp, [shared_mem];
add.u64 temp, temp, 2;
- st.param.u64 [output], temp;
+ st.shared.u64 [shared_mem], temp;
ret;
}
-.visible .entry extern_shared(
+.visible .entry extern_shared_call(
.param .u64 input,
.param .u64 output
)
@@ -39,6 +28,7 @@ ld.global.u64 temp, [in_addr];
st.shared.u64 [shared_mem], temp;
+ call incr_shared_2_global;
ld.shared.u64 temp, [shared_mem];
st.global.u64 [out_addr], temp;
ret;
diff --git a/ptx/src/test/spirv_run/extern_shared_call.spvtxt b/ptx/src/test/spirv_run/extern_shared_call.spvtxt index 84e7eac..d979193 100644 --- a/ptx/src/test/spirv_run/extern_shared_call.spvtxt +++ b/ptx/src/test/spirv_run/extern_shared_call.spvtxt @@ -2,52 +2,94 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %29 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %48 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL - OpEntryPoint Kernel %1 "cvta" + OpEntryPoint Kernel %14 "extern_shared_call" %1 + OpDecorate %1 Alignment 4 %void = OpTypeVoid + %uint = OpTypeInt 32 0 +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint +%_ptr_Workgroup__ptr_Workgroup_uint = OpTypePointer Workgroup %_ptr_Workgroup_uint + %1 = OpVariable %_ptr_Workgroup__ptr_Workgroup_uint Workgroup + %uchar = OpTypeInt 8 0 +%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar + %55 = OpTypeFunction %void %_ptr_Workgroup_uchar +%_ptr_Function__ptr_Workgroup_uchar = OpTypePointer Function %_ptr_Workgroup_uchar %ulong = OpTypeInt 64 0 - %32 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong - %float = OpTypeFloat 32 -%_ptr_Function_float = OpTypePointer Function %float -%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float - %1 = OpFunction %void None %32 - %7 = OpFunctionParameter %ulong - %8 = OpFunctionParameter %ulong - %27 = OpLabel - %2 = OpVariable %_ptr_Function_ulong Function +%_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 +%_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 %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 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 - OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 - %22 = OpCopyObject %ulong %14 - %21 = OpCopyObject %ulong %22 - %13 = OpCopyObject %ulong %21 - OpStore %4 %13 - %16 = OpLoad %ulong %5 - %24 = OpCopyObject %ulong %16 - %23 = OpCopyObject %ulong %24 - %15 = OpCopyObject %ulong %23 - OpStore %5 %15 - %18 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %18 - %17 = OpLoad %float %25 - OpStore %6 %17 - %19 = OpLoad %ulong %5 - %20 = OpLoad %float %6 - %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %19 - OpStore %26 %20 + OpStore %41 %40 + OpBranch %13 + %13 = OpLabel + %42 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %41 + %5 = OpLoad %_ptr_Workgroup_uint %42 + %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 + %9 = OpLoad %ulong %3 + %12 = OpBitcast %_ptr_Workgroup_ulong %8 + OpStore %12 %9 + OpReturn + OpFunctionEnd + %14 = OpFunction %void None %62 + %20 = OpFunctionParameter %ulong + %21 = OpFunctionParameter %ulong + %44 = OpFunctionParameter %_ptr_Workgroup_uchar + %63 = OpLabel + %45 = 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 %15 %20 + OpStore %16 %21 + %23 = OpLoad %ulong %15 + %22 = OpCopyObject %ulong %23 + 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 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 4c5f9b3..3a8acb1 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -79,6 +79,7 @@ test_ptx!(min, [555i32, 444i32], [444i32]); test_ptx!(max, [555i32, 444i32], [555i32]);
test_ptx!(global_array, [0xDEADu32], [1u32]);
test_ptx!(extern_shared, [127u64], [127u64]);
+test_ptx!(extern_shared_call, [121u64], [123u64]);
struct DisplayError<T: Debug> {
err: T,
|