aboutsummaryrefslogtreecommitdiffhomepage
path: root/ptx/src/test/spirv_run
diff options
context:
space:
mode:
Diffstat (limited to 'ptx/src/test/spirv_run')
-rw-r--r--ptx/src/test/spirv_run/atom_add.spvtxt37
-rw-r--r--ptx/src/test/spirv_run/atom_add_float.spvtxt41
-rw-r--r--ptx/src/test/spirv_run/extern_shared.spvtxt24
-rw-r--r--ptx/src/test/spirv_run/extern_shared_call.spvtxt40
-rw-r--r--ptx/src/test/spirv_run/ld_st_implicit.spvtxt39
-rw-r--r--ptx/src/test/spirv_run/mod.rs3
-rw-r--r--ptx/src/test/spirv_run/shared_ptr_32.spvtxt23
-rw-r--r--ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt16
-rw-r--r--ptx/src/test/spirv_run/shared_unify_extern.ptx22
-rw-r--r--ptx/src/test/spirv_run/shared_unify_extern.spvtxt153
-rw-r--r--ptx/src/test/spirv_run/shared_unify_private.ptx32
-rw-r--r--ptx/src/test/spirv_run/shared_unify_private.spvtxt84
-rw-r--r--ptx/src/test/spirv_run/shared_variable.spvtxt21
13 files changed, 232 insertions, 303 deletions
diff --git a/ptx/src/test/spirv_run/atom_add.spvtxt b/ptx/src/test/spirv_run/atom_add.spvtxt
index 3609247..987fdef 100644
--- a/ptx/src/test/spirv_run/atom_add.spvtxt
+++ b/ptx/src/test/spirv_run/atom_add.spvtxt
@@ -8,32 +8,32 @@
OpCapability Float16
OpCapability Float64
OpCapability DenormFlushToZero
- %42 = OpExtInstImport "OpenCL.std"
+ OpExtension "SPV_KHR_float_controls"
+ OpExtension "SPV_KHR_no_integer_wrap_decoration"
+ %38 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %1 "atom_add" %38
+ OpEntryPoint Kernel %1 "atom_add" %4
OpExecutionMode %1 ContractionOff
+ OpDecorate %4 Alignment 4
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%uint_1024 = OpConstant %uint 1024
%_arr_uchar_uint_1024 = OpTypeArray %uchar %uint_1024
%_ptr_Workgroup__arr_uchar_uint_1024 = OpTypePointer Workgroup %_arr_uchar_uint_1024
- %38 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
+ %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
%ulong = OpTypeInt 64 0
- %50 = OpTypeFunction %void %ulong %ulong
+ %46 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Generic_uint = OpTypePointer Generic %uint
%ulong_4 = OpConstant %ulong 4
%_ptr_Generic_uchar = OpTypePointer Generic %uchar
-%uint_1024_0 = OpConstant %uint 1024
%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
-%uint_1024_1 = OpConstant %uint 1024
%uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0
-%uint_1024_2 = OpConstant %uint 1024
%ulong_4_0 = OpConstant %ulong 4
- %1 = OpFunction %void None %50
+ %1 = OpFunction %void None %46
%9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong
%36 = OpLabel
@@ -55,22 +55,19 @@
OpStore %7 %13
%16 = OpLoad %ulong %5
%30 = OpConvertUToPtr %_ptr_Generic_uint %16
- %55 = OpBitcast %_ptr_Generic_uchar %30
- %56 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %55 %ulong_4
- %26 = OpBitcast %_ptr_Generic_uint %56
+ %51 = OpBitcast %_ptr_Generic_uchar %30
+ %52 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %51 %ulong_4
+ %26 = OpBitcast %_ptr_Generic_uint %52
%15 = OpLoad %uint %26 Aligned 4
OpStore %8 %15
%17 = OpLoad %uint %7
- %39 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %38
- %31 = OpBitcast %_ptr_Workgroup_uint %39
+ %31 = OpBitcast %_ptr_Workgroup_uint %4
OpStore %31 %17 Aligned 4
%19 = OpLoad %uint %8
- %40 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %38
- %32 = OpBitcast %_ptr_Workgroup_uint %40
+ %32 = OpBitcast %_ptr_Workgroup_uint %4
%18 = OpAtomicIAdd %uint %32 %uint_1 %uint_0 %19
OpStore %7 %18
- %41 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %38
- %33 = OpBitcast %_ptr_Workgroup_uint %41
+ %33 = OpBitcast %_ptr_Workgroup_uint %4
%20 = OpLoad %uint %33 Aligned 4
OpStore %8 %20
%21 = OpLoad %ulong %6
@@ -80,9 +77,9 @@
%23 = OpLoad %ulong %6
%24 = OpLoad %uint %8
%35 = OpConvertUToPtr %_ptr_Generic_uint %23
- %63 = OpBitcast %_ptr_Generic_uchar %35
- %64 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %63 %ulong_4_0
- %28 = OpBitcast %_ptr_Generic_uint %64
+ %56 = OpBitcast %_ptr_Generic_uchar %35
+ %57 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %56 %ulong_4_0
+ %28 = OpBitcast %_ptr_Generic_uint %57
OpStore %28 %24 Aligned 4
OpReturn
OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/atom_add_float.spvtxt b/ptx/src/test/spirv_run/atom_add_float.spvtxt
index 9533d83..067c347 100644
--- a/ptx/src/test/spirv_run/atom_add_float.spvtxt
+++ b/ptx/src/test/spirv_run/atom_add_float.spvtxt
@@ -8,37 +8,37 @@
OpCapability Float16
OpCapability Float64
OpCapability DenormFlushToZero
- %46 = OpExtInstImport "OpenCL.std"
+ OpExtension "SPV_KHR_float_controls"
+ OpExtension "SPV_KHR_no_integer_wrap_decoration"
+ %42 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %1 "atom_add_float" %42
+ OpEntryPoint Kernel %1 "atom_add_float" %4
OpExecutionMode %1 ContractionOff
OpDecorate %37 LinkageAttributes "__zluda_ptx_impl__atom_relaxed_gpu_shared_add_f32" Import
+ OpDecorate %4 Alignment 4
%void = OpTypeVoid
%float = OpTypeFloat 32
%_ptr_Workgroup_float = OpTypePointer Workgroup %float
- %50 = OpTypeFunction %float %_ptr_Workgroup_float %float
+ %46 = OpTypeFunction %float %_ptr_Workgroup_float %float
%uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%uint_1024 = OpConstant %uint 1024
%_arr_uchar_uint_1024 = OpTypeArray %uchar %uint_1024
%_ptr_Workgroup__arr_uchar_uint_1024 = OpTypePointer Workgroup %_arr_uchar_uint_1024
- %42 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
+ %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
%ulong = OpTypeInt 64 0
- %57 = OpTypeFunction %void %ulong %ulong
+ %53 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Function_float = OpTypePointer Function %float
%_ptr_Generic_float = OpTypePointer Generic %float
%ulong_4 = OpConstant %ulong 4
%_ptr_Generic_uchar = OpTypePointer Generic %uchar
-%uint_1024_0 = OpConstant %uint 1024
-%uint_1024_1 = OpConstant %uint 1024
-%uint_1024_2 = OpConstant %uint 1024
%ulong_4_0 = OpConstant %ulong 4
- %37 = OpFunction %float None %50
+ %37 = OpFunction %float None %46
%39 = OpFunctionParameter %_ptr_Workgroup_float
%40 = OpFunctionParameter %float
OpFunctionEnd
- %1 = OpFunction %void None %57
+ %1 = OpFunction %void None %53
%9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong
%36 = OpLabel
@@ -60,22 +60,19 @@
OpStore %7 %13
%16 = OpLoad %ulong %5
%30 = OpConvertUToPtr %_ptr_Generic_float %16
- %62 = OpBitcast %_ptr_Generic_uchar %30
- %63 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %62 %ulong_4
- %26 = OpBitcast %_ptr_Generic_float %63
+ %58 = OpBitcast %_ptr_Generic_uchar %30
+ %59 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %58 %ulong_4
+ %26 = OpBitcast %_ptr_Generic_float %59
%15 = OpLoad %float %26 Aligned 4
OpStore %8 %15
%17 = OpLoad %float %7
- %43 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %42
- %31 = OpBitcast %_ptr_Workgroup_float %43
+ %31 = OpBitcast %_ptr_Workgroup_float %4
OpStore %31 %17 Aligned 4
%19 = OpLoad %float %8
- %44 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %42
- %32 = OpBitcast %_ptr_Workgroup_float %44
+ %32 = OpBitcast %_ptr_Workgroup_float %4
%18 = OpFunctionCall %float %37 %32 %19
OpStore %7 %18
- %45 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %42
- %33 = OpBitcast %_ptr_Workgroup_float %45
+ %33 = OpBitcast %_ptr_Workgroup_float %4
%20 = OpLoad %float %33 Aligned 4
OpStore %8 %20
%21 = OpLoad %ulong %6
@@ -85,9 +82,9 @@
%23 = OpLoad %ulong %6
%24 = OpLoad %float %8
%35 = OpConvertUToPtr %_ptr_Generic_float %23
- %67 = OpBitcast %_ptr_Generic_uchar %35
- %68 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %67 %ulong_4_0
- %28 = OpBitcast %_ptr_Generic_float %68
+ %60 = OpBitcast %_ptr_Generic_uchar %35
+ %61 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %60 %ulong_4_0
+ %28 = OpBitcast %_ptr_Generic_float %61
OpStore %28 %24 Aligned 4
OpReturn
OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/extern_shared.spvtxt b/ptx/src/test/spirv_run/extern_shared.spvtxt
index 82d86ae..025cd81 100644
--- a/ptx/src/test/spirv_run/extern_shared.spvtxt
+++ b/ptx/src/test/spirv_run/extern_shared.spvtxt
@@ -8,22 +8,22 @@
OpCapability Float16
OpCapability Float64
OpCapability DenormFlushToZero
- %27 = OpExtInstImport "OpenCL.std"
+ OpExtension "SPV_KHR_float_controls"
+ OpExtension "SPV_KHR_no_integer_wrap_decoration"
+ %24 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %2 "extern_shared" %24
+ OpEntryPoint Kernel %2 "extern_shared" %1
OpExecutionMode %2 ContractionOff
%void = OpTypeVoid
- %uchar = OpTypeInt 8 0
-%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
- %24 = OpVariable %_ptr_Workgroup_uchar Workgroup
+ %uint = OpTypeInt 32 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %1 = OpVariable %_ptr_Workgroup_uint Workgroup
%ulong = OpTypeInt 64 0
- %32 = OpTypeFunction %void %ulong %ulong
+ %29 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
- %uint = OpTypeInt 32 0
-%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
- %2 = OpFunction %void None %32
+ %2 = OpFunction %void None %29
%8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong
%22 = OpLabel
@@ -43,11 +43,9 @@
%12 = OpLoad %ulong %18 Aligned 8
OpStore %7 %12
%14 = OpLoad %ulong %7
- %25 = OpBitcast %_ptr_Workgroup_uint %24
- %19 = OpBitcast %_ptr_Workgroup_ulong %25
+ %19 = OpBitcast %_ptr_Workgroup_ulong %1
OpStore %19 %14 Aligned 8
- %26 = OpBitcast %_ptr_Workgroup_uint %24
- %20 = OpBitcast %_ptr_Workgroup_ulong %26
+ %20 = OpBitcast %_ptr_Workgroup_ulong %1
%15 = OpLoad %ulong %20 Aligned 8
OpStore %7 %15
%16 = OpLoad %ulong %6
diff --git a/ptx/src/test/spirv_run/extern_shared_call.spvtxt b/ptx/src/test/spirv_run/extern_shared_call.spvtxt
index 3cc78cb..bf1dccd 100644
--- a/ptx/src/test/spirv_run/extern_shared_call.spvtxt
+++ b/ptx/src/test/spirv_run/extern_shared_call.spvtxt
@@ -8,41 +8,40 @@
OpCapability Float16
OpCapability Float64
OpCapability DenormFlushToZero
- %41 = OpExtInstImport "OpenCL.std"
+ OpExtension "SPV_KHR_float_controls"
+ OpExtension "SPV_KHR_no_integer_wrap_decoration"
+ %35 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %12 "extern_shared_call" %37
+ OpEntryPoint Kernel %12 "extern_shared_call" %1
OpExecutionMode %12 ContractionOff
+ OpDecorate %1 Alignment 4
%void = OpTypeVoid
- %uchar = OpTypeInt 8 0
-%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
- %45 = OpTypeFunction %void %_ptr_Workgroup_uchar
- %ulong = OpTypeInt 64 0
-%_ptr_Function_ulong = OpTypePointer Function %ulong
%uint = OpTypeInt 32 0
%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %1 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %39 = OpTypeFunction %void %_ptr_Workgroup_uint
+ %ulong = OpTypeInt 64 0
+%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%ulong_2 = OpConstant %ulong 2
- %37 = OpVariable %_ptr_Workgroup_uchar Workgroup
- %51 = OpTypeFunction %void %ulong %ulong
+ %43 = OpTypeFunction %void %ulong %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
- %2 = OpFunction %void None %45
- %34 = OpFunctionParameter %_ptr_Workgroup_uchar
+ %2 = OpFunction %void None %39
+ %34 = OpFunctionParameter %_ptr_Workgroup_uint
%11 = OpLabel
%3 = OpVariable %_ptr_Function_ulong Function
- %35 = OpBitcast %_ptr_Workgroup_uint %34
- %9 = OpBitcast %_ptr_Workgroup_ulong %35
+ %9 = OpBitcast %_ptr_Workgroup_ulong %34
%4 = OpLoad %ulong %9 Aligned 8
OpStore %3 %4
%6 = OpLoad %ulong %3
%5 = OpIAdd %ulong %6 %ulong_2
OpStore %3 %5
%7 = OpLoad %ulong %3
- %36 = OpBitcast %_ptr_Workgroup_uint %34
- %10 = OpBitcast %_ptr_Workgroup_ulong %36
+ %10 = OpBitcast %_ptr_Workgroup_ulong %34
OpStore %10 %7 Aligned 8
OpReturn
OpFunctionEnd
- %12 = OpFunction %void None %51
+ %12 = OpFunction %void None %43
%18 = OpFunctionParameter %ulong
%19 = OpFunctionParameter %ulong
%32 = OpLabel
@@ -62,13 +61,10 @@
%22 = OpLoad %ulong %28 Aligned 8
OpStore %17 %22
%24 = OpLoad %ulong %17
- %38 = OpBitcast %_ptr_Workgroup_uint %37
- %29 = OpBitcast %_ptr_Workgroup_ulong %38
+ %29 = OpBitcast %_ptr_Workgroup_ulong %1
OpStore %29 %24 Aligned 8
- %39 = OpBitcast %_ptr_Workgroup_uchar %37
- %53 = OpFunctionCall %void %2 %39
- %40 = OpBitcast %_ptr_Workgroup_uint %37
- %30 = OpBitcast %_ptr_Workgroup_ulong %40
+ %45 = OpFunctionCall %void %2 %1
+ %30 = OpBitcast %_ptr_Workgroup_ulong %1
%25 = OpLoad %ulong %30 Aligned 8
OpStore %17 %25
%26 = OpLoad %ulong %16
diff --git a/ptx/src/test/spirv_run/ld_st_implicit.spvtxt b/ptx/src/test/spirv_run/ld_st_implicit.spvtxt
index 29f46f9..9c0e508 100644
--- a/ptx/src/test/spirv_run/ld_st_implicit.spvtxt
+++ b/ptx/src/test/spirv_run/ld_st_implicit.spvtxt
@@ -7,20 +7,25 @@
OpCapability Int64
OpCapability Float16
OpCapability Float64
- %21 = OpExtInstImport "OpenCL.std"
+ OpCapability DenormFlushToZero
+ OpExtension "SPV_KHR_float_controls"
+ OpExtension "SPV_KHR_no_integer_wrap_decoration"
+ %23 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "ld_st_implicit"
+ OpExecutionMode %1 ContractionOff
%void = OpTypeVoid
%ulong = OpTypeInt 64 0
- %24 = OpTypeFunction %void %ulong %ulong
+ %26 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
+%ulong_81985529216486895 = OpConstant %ulong 81985529216486895
%float = OpTypeFloat 32
%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
%uint = OpTypeInt 32 0
- %1 = OpFunction %void None %24
+ %1 = OpFunction %void None %26
%7 = OpFunctionParameter %ulong
%8 = OpFunctionParameter %ulong
- %19 = OpLabel
+ %21 = OpLabel
%2 = OpVariable %_ptr_Function_ulong Function
%3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function
@@ -32,18 +37,20 @@
OpStore %4 %9
%10 = OpLoad %ulong %3 Aligned 8
OpStore %5 %10
- %12 = OpLoad %ulong %4
- %16 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %12
- %15 = OpLoad %float %16 Aligned 4
- %29 = OpBitcast %uint %15
- %11 = OpUConvert %ulong %29
+ %11 = OpCopyObject %ulong %ulong_81985529216486895
OpStore %6 %11
- %13 = OpLoad %ulong %5
- %14 = OpLoad %ulong %6
- %17 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %13
- %30 = OpBitcast %ulong %14
- %31 = OpUConvert %uint %30
- %18 = OpBitcast %float %31
- OpStore %17 %18 Aligned 4
+ %13 = OpLoad %ulong %4
+ %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %13
+ %17 = OpLoad %float %18 Aligned 4
+ %31 = OpBitcast %uint %17
+ %12 = OpUConvert %ulong %31
+ OpStore %6 %12
+ %14 = OpLoad %ulong %5
+ %15 = OpLoad %ulong %6
+ %19 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %14
+ %32 = OpBitcast %ulong %15
+ %33 = OpUConvert %uint %32
+ %20 = OpBitcast %float %33
+ OpStore %19 %20 Aligned 4
OpReturn
OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 512b6cf..dfc252d 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -218,8 +218,7 @@ test_ptx!(cvt_f64_f32, [0.125f32], [0.125f64]);
test_ptx!(prmt, [0x70c507d6u32, 0x6fbd4b5cu32], [0x6fbdd65cu32]);
test_ptx!(activemask, [0u32], [1u32]);
test_ptx!(membar, [152731u32], [152731u32]);
-test_ptx!(shared_unify_extern, [7681u64], [15362u64]);
-test_ptx!(shared_unify_private, [67153u64], [134306u64]);
+test_ptx!(shared_unify_extern, [7681u64, 7682u64], [15363u64]);
test_ptx!(assertfail);
test_ptx!(func_ptr);
diff --git a/ptx/src/test/spirv_run/shared_ptr_32.spvtxt b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt
index 020c15b..787a71c 100644
--- a/ptx/src/test/spirv_run/shared_ptr_32.spvtxt
+++ b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt
@@ -8,27 +8,29 @@
OpCapability Float16
OpCapability Float64
OpCapability DenormFlushToZero
- %34 = OpExtInstImport "OpenCL.std"
+ OpExtension "SPV_KHR_float_controls"
+ OpExtension "SPV_KHR_no_integer_wrap_decoration"
+ %32 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %1 "shared_ptr_32" %32
+ OpEntryPoint Kernel %1 "shared_ptr_32" %4
OpExecutionMode %1 ContractionOff
+ OpDecorate %4 Alignment 4
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%uint_128 = OpConstant %uint 128
%_arr_uchar_uint_128 = OpTypeArray %uchar %uint_128
%_ptr_Workgroup__arr_uchar_uint_128 = OpTypePointer Workgroup %_arr_uchar_uint_128
- %32 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
+ %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
%ulong = OpTypeInt 64 0
- %42 = OpTypeFunction %void %ulong %ulong
+ %40 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Function_uint = OpTypePointer Function %uint
- %uint_128_0 = OpConstant %uint 128
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%ulong_0 = OpConstant %ulong 0
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
- %1 = OpFunction %void None %42
+ %1 = OpFunction %void None %40
%10 = OpFunctionParameter %ulong
%11 = OpFunctionParameter %ulong
%30 = OpLabel
@@ -45,8 +47,7 @@
OpStore %5 %12
%13 = OpLoad %ulong %3 Aligned 8
OpStore %6 %13
- %33 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_128 %32
- %25 = OpConvertPtrToU %uint %33
+ %25 = OpConvertPtrToU %uint %4
%14 = OpCopyObject %uint %25
OpStore %7 %14
%16 = OpLoad %ulong %5
@@ -59,9 +60,9 @@
OpStore %27 %18 Aligned 8
%20 = OpLoad %uint %7
%28 = OpConvertUToPtr %_ptr_Workgroup_ulong %20
- %49 = OpBitcast %_ptr_Workgroup_uchar %28
- %50 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %49 %ulong_0
- %24 = OpBitcast %_ptr_Workgroup_ulong %50
+ %46 = OpBitcast %_ptr_Workgroup_uchar %28
+ %47 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %46 %ulong_0
+ %24 = OpBitcast %_ptr_Workgroup_ulong %47
%19 = OpLoad %ulong %24 Aligned 8
OpStore %9 %19
%21 = OpLoad %ulong %6
diff --git a/ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt b/ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt
index 90e04f3..14926ef 100644
--- a/ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt
+++ b/ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt
@@ -8,20 +8,23 @@
OpCapability Float16
OpCapability Float64
OpCapability DenormFlushToZero
- %32 = OpExtInstImport "OpenCL.std"
+ OpExtension "SPV_KHR_float_controls"
+ OpExtension "SPV_KHR_no_integer_wrap_decoration"
+ %30 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %2 "shared_ptr_take_address" %30
+ OpEntryPoint Kernel %2 "shared_ptr_take_address" %1
OpExecutionMode %2 ContractionOff
+ OpDecorate %1 Alignment 4
%void = OpTypeVoid
%uchar = OpTypeInt 8 0
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
- %30 = OpVariable %_ptr_Workgroup_uchar Workgroup
+ %1 = OpVariable %_ptr_Workgroup_uchar Workgroup
%ulong = OpTypeInt 64 0
- %37 = OpTypeFunction %void %ulong %ulong
+ %35 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
- %2 = OpFunction %void None %37
+ %2 = OpFunction %void None %35
%10 = OpFunctionParameter %ulong
%11 = OpFunctionParameter %ulong
%28 = OpLabel
@@ -38,8 +41,7 @@
OpStore %5 %12
%13 = OpLoad %ulong %4 Aligned 8
OpStore %6 %13
- %31 = OpBitcast %_ptr_Workgroup_uchar %30
- %23 = OpConvertPtrToU %ulong %31
+ %23 = OpConvertPtrToU %ulong %1
%14 = OpCopyObject %ulong %23
OpStore %7 %14
%16 = OpLoad %ulong %5
diff --git a/ptx/src/test/spirv_run/shared_unify_extern.ptx b/ptx/src/test/spirv_run/shared_unify_extern.ptx
index 8b406b2..0416e2c 100644
--- a/ptx/src/test/spirv_run/shared_unify_extern.ptx
+++ b/ptx/src/test/spirv_run/shared_unify_extern.ptx
@@ -5,10 +5,20 @@
.extern .shared .b32 shared_ex[];
.shared .b32 shared_mod[4];
+.func (.reg .b64 out) add()
+{
+ .reg .u64 temp1;
+ .reg .u64 temp2;
+ ld.shared.u64 temp1, [shared_mod];
+ ld.shared.u64 temp2, [shared_ex];
+ add.u64 out, temp2, temp1;
+ ret;
+}
-.func (.reg .b64 out) load_from_shared()
+.func (.reg .b64 out) set_shared_temp1(.reg .b64 temp1)
{
- ld.shared.u64 out, [shared_mod];
+ st.shared.u64 [shared_ex], temp1;
+ call (out), add;
ret;
}
@@ -25,10 +35,10 @@
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
- ld.u64 temp1, [in_addr];
- st.shared.u64 [shared_ex], temp1;
- call (temp2), load_from_shared;
- add.u64 temp2, temp2, temp1;
+ ld.global.u64 temp1, [in_addr];
+ ld.global.u64 temp2, [in_addr+8];
+ st.shared.u64 [shared_mod], temp2;
+ call (temp2), set_shared_temp1, (temp1);
st.u64 [out_addr], temp2;
ret;
}
diff --git a/ptx/src/test/spirv_run/shared_unify_extern.spvtxt b/ptx/src/test/spirv_run/shared_unify_extern.spvtxt
index 2dd2056..90fc156 100644
--- a/ptx/src/test/spirv_run/shared_unify_extern.spvtxt
+++ b/ptx/src/test/spirv_run/shared_unify_extern.spvtxt
@@ -8,72 +8,111 @@
OpCapability Float16
OpCapability Float64
OpCapability DenormFlushToZero
- %41 = OpExtInstImport "OpenCL.std"
+ OpExtension "SPV_KHR_float_controls"
+ OpExtension "SPV_KHR_no_integer_wrap_decoration"
+ %61 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %10 "shared_unify_extern" %38
- OpExecutionMode %10 ContractionOff
+ OpEntryPoint Kernel %27 "shared_unify_extern" %1 %2
+ OpExecutionMode %27 ContractionOff
%void = OpTypeVoid
- %ulong = OpTypeInt 64 0
- %uchar = OpTypeInt 8 0
-%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
- %46 = OpTypeFunction %ulong %_ptr_Workgroup_uchar
-%_ptr_Function_ulong = OpTypePointer Function %ulong
%uint = OpTypeInt 32 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %1 = OpVariable %_ptr_Workgroup_uint Workgroup
%uint_4 = OpConstant %uint 4
%_arr_uint_uint_4 = OpTypeArray %uint %uint_4
%_ptr_Workgroup__arr_uint_uint_4 = OpTypePointer Workgroup %_arr_uint_uint_4
+ %2 = OpVariable %_ptr_Workgroup__arr_uint_uint_4 Workgroup
+ %ulong = OpTypeInt 64 0
+ %uint_4_0 = OpConstant %uint 4
+ %70 = OpTypeFunction %ulong %_ptr_Workgroup_uint %_ptr_Workgroup__arr_uint_uint_4
+ %uint_4_1 = OpConstant %uint 4
+%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
- %38 = OpVariable %_ptr_Workgroup_uchar Workgroup
- %53 = OpTypeFunction %void %ulong %ulong
+ %uint_4_2 = OpConstant %uint 4
+ %75 = OpTypeFunction %ulong %ulong %_ptr_Workgroup_uint %_ptr_Workgroup__arr_uint_uint_4
+ %uint_4_3 = OpConstant %uint 4
+ %77 = OpTypeFunction %void %ulong %ulong
+%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
+ %ulong_8 = OpConstant %ulong 8
+ %uchar = OpTypeInt 8 0
+%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%_ptr_Generic_ulong = OpTypePointer Generic %ulong
-%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
- %3 = OpFunction %ulong None %46
- %36 = OpFunctionParameter %_ptr_Workgroup_uchar
- %9 = OpLabel
+ %3 = OpFunction %ulong None %70
+ %57 = OpFunctionParameter %_ptr_Workgroup_uint
+ %58 = OpFunctionParameter %_ptr_Workgroup__arr_uint_uint_4
+ %16 = OpLabel
%4 = OpVariable %_ptr_Function_ulong Function
- %37 = OpBitcast %_ptr_Workgroup__arr_uint_uint_4 %36
- %8 = OpBitcast %_ptr_Workgroup_ulong %37
- %7 = OpLoad %ulong %8 Aligned 8
- %5 = OpCopyObject %ulong %7
- OpStore %4 %5
- %6 = OpLoad %ulong %4
- OpReturnValue %6
+ %5 = OpVariable %_ptr_Function_ulong Function
+ %6 = OpVariable %_ptr_Function_ulong Function
+ %13 = OpBitcast %_ptr_Workgroup_ulong %58
+ %7 = OpLoad %ulong %13 Aligned 8
+ OpStore %5 %7
+ %14 = OpBitcast %_ptr_Workgroup_ulong %57
+ %8 = OpLoad %ulong %14 Aligned 8
+ OpStore %6 %8
+ %10 = OpLoad %ulong %6
+ %11 = OpLoad %ulong %5
+ %15 = OpIAdd %ulong %10 %11
+ %9 = OpCopyObject %ulong %15
+ OpStore %4 %9
+ %12 = OpLoad %ulong %4
+ OpReturnValue %12
+ OpFunctionEnd
+ %17 = OpFunction %ulong None %75
+ %20 = OpFunctionParameter %ulong
+ %59 = OpFunctionParameter %_ptr_Workgroup_uint
+ %60 = OpFunctionParameter %_ptr_Workgroup__arr_uint_uint_4
+ %26 = OpLabel
+ %19 = OpVariable %_ptr_Function_ulong Function
+ %18 = OpVariable %_ptr_Function_ulong Function
+ OpStore %19 %20
+ %21 = OpLoad %ulong %19
+ %24 = OpBitcast %_ptr_Workgroup_ulong %59
+ %25 = OpCopyObject %ulong %21
+ OpStore %24 %25 Aligned 8
+ %22 = OpFunctionCall %ulong %3 %59 %60
+ OpStore %18 %22
+ %23 = OpLoad %ulong %18
+ OpReturnValue %23
OpFunctionEnd
- %10 = OpFunction %void None %53
- %17 = OpFunctionParameter %ulong
- %18 = OpFunctionParameter %ulong
- %34 = OpLabel
- %11 = OpVariable %_ptr_Function_ulong Function
- %12 = OpVariable %_ptr_Function_ulong Function
- %13 = OpVariable %_ptr_Function_ulong Function
- %14 = OpVariable %_ptr_Function_ulong Function
- %15 = OpVariable %_ptr_Function_ulong Function
- %16 = OpVariable %_ptr_Function_ulong Function
- OpStore %11 %17
- OpStore %12 %18
- %19 = OpLoad %ulong %11 Aligned 8
- OpStore %13 %19
- %20 = OpLoad %ulong %12 Aligned 8
- OpStore %14 %20
- %22 = OpLoad %ulong %13
- %30 = OpConvertUToPtr %_ptr_Generic_ulong %22
- %21 = OpLoad %ulong %30 Aligned 8
- OpStore %15 %21
- %23 = OpLoad %ulong %15
- %39 = OpBitcast %_ptr_Workgroup_uint %38
- %31 = OpBitcast %_ptr_Workgroup_ulong %39
- OpStore %31 %23 Aligned 8
- %40 = OpBitcast %_ptr_Workgroup_uchar %38
- %32 = OpFunctionCall %ulong %3 %40
- %24 = OpCopyObject %ulong %32
- OpStore %16 %24
- %26 = OpLoad %ulong %16
- %27 = OpLoad %ulong %15
- %25 = OpIAdd %ulong %26 %27
- OpStore %16 %25
- %28 = OpLoad %ulong %14
- %29 = OpLoad %ulong %16
- %33 = OpConvertUToPtr %_ptr_Generic_ulong %28
- OpStore %33 %29 Aligned 8
+ %27 = OpFunction %void None %77
+ %34 = OpFunctionParameter %ulong
+ %35 = OpFunctionParameter %ulong
+ %55 = OpLabel
+ %28 = OpVariable %_ptr_Function_ulong Function
+ %29 = OpVariable %_ptr_Function_ulong Function
+ %30 = OpVariable %_ptr_Function_ulong Function
+ %31 = OpVariable %_ptr_Function_ulong Function
+ %32 = OpVariable %_ptr_Function_ulong Function
+ %33 = OpVariable %_ptr_Function_ulong Function
+ OpStore %28 %34
+ OpStore %29 %35
+ %36 = OpLoad %ulong %28 Aligned 8
+ OpStore %30 %36
+ %37 = OpLoad %ulong %29 Aligned 8
+ OpStore %31 %37
+ %39 = OpLoad %ulong %30
+ %49 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %39
+ %38 = OpLoad %ulong %49 Aligned 8
+ OpStore %32 %38
+ %41 = OpLoad %ulong %30
+ %50 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %41
+ %81 = OpBitcast %_ptr_CrossWorkgroup_uchar %50
+ %82 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %81 %ulong_8
+ %48 = OpBitcast %_ptr_CrossWorkgroup_ulong %82
+ %40 = OpLoad %ulong %48 Aligned 8
+ OpStore %33 %40
+ %42 = OpLoad %ulong %33
+ %51 = OpBitcast %_ptr_Workgroup_ulong %2
+ OpStore %51 %42 Aligned 8
+ %44 = OpLoad %ulong %32
+ %53 = OpCopyObject %ulong %44
+ %52 = OpFunctionCall %ulong %17 %53 %1 %2
+ %43 = OpCopyObject %ulong %52
+ OpStore %33 %43
+ %45 = OpLoad %ulong %31
+ %46 = OpLoad %ulong %33
+ %54 = OpConvertUToPtr %_ptr_Generic_ulong %45
+ OpStore %54 %46 Aligned 8
OpReturn
OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/shared_unify_private.ptx b/ptx/src/test/spirv_run/shared_unify_private.ptx
deleted file mode 100644
index fd31357..0000000
--- a/ptx/src/test/spirv_run/shared_unify_private.ptx
+++ /dev/null
@@ -1,32 +0,0 @@
-.version 6.5
-.target sm_30
-.address_size 64
-
-.func (.reg .b64 out) load_from_shared()
-{
- .shared .b32 shared_mod[4];
- ld.shared.u64 out, [shared_mod];
- ret;
-}
-
-.visible .entry shared_unify_private(
- .param .u64 input,
- .param .u64 output
-)
-{
- .shared .b32 shared_ex[2];
- .reg .u64 in_addr;
- .reg .u64 out_addr;
- .reg .u64 temp1;
- .reg .u64 temp2;
-
- ld.param.u64 in_addr, [input];
- ld.param.u64 out_addr, [output];
-
- ld.u64 temp1, [in_addr];
- st.shared.u64 [shared_ex], temp1;
- call (temp2), load_from_shared;
- add.u64 temp2, temp2, temp1;
- st.u64 [out_addr], temp2;
- ret;
-}
diff --git a/ptx/src/test/spirv_run/shared_unify_private.spvtxt b/ptx/src/test/spirv_run/shared_unify_private.spvtxt
deleted file mode 100644
index 69bf018..0000000
--- a/ptx/src/test/spirv_run/shared_unify_private.spvtxt
+++ /dev/null
@@ -1,84 +0,0 @@
- OpCapability GenericPointer
- OpCapability Linkage
- OpCapability Addresses
- OpCapability Kernel
- OpCapability Int8
- OpCapability Int16
- OpCapability Int64
- OpCapability Float16
- OpCapability Float64
- OpCapability DenormFlushToZero
- %41 = OpExtInstImport "OpenCL.std"
- OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %9 "shared_unify_private" %38
- OpExecutionMode %9 ContractionOff
- %void = OpTypeVoid
- %ulong = OpTypeInt 64 0
- %uchar = OpTypeInt 8 0
-%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
- %46 = OpTypeFunction %ulong %_ptr_Workgroup_uchar
-%_ptr_Function_ulong = OpTypePointer Function %ulong
- %uint = OpTypeInt 32 0
- %uint_4 = OpConstant %uint 4
-%_arr_uint_uint_4 = OpTypeArray %uint %uint_4
-%_ptr_Workgroup__arr_uint_uint_4 = OpTypePointer Workgroup %_arr_uint_uint_4
-%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
- %uint_16 = OpConstant %uint 16
-%_arr_uchar_uint_16 = OpTypeArray %uchar %uint_16
-%_ptr_Workgroup__arr_uchar_uint_16 = OpTypePointer Workgroup %_arr_uchar_uint_16
- %38 = OpVariable %_ptr_Workgroup__arr_uchar_uint_16 Workgroup
- %56 = OpTypeFunction %void %ulong %ulong
-%_ptr_Generic_ulong = OpTypePointer Generic %ulong
- %uint_2 = OpConstant %uint 2
-%_arr_uint_uint_2 = OpTypeArray %uint %uint_2
-%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2
- %1 = OpFunction %ulong None %46
- %36 = OpFunctionParameter %_ptr_Workgroup_uchar
- %8 = OpLabel
- %2 = OpVariable %_ptr_Function_ulong Function
- %37 = OpBitcast %_ptr_Workgroup__arr_uint_uint_4 %36
- %7 = OpBitcast %_ptr_Workgroup_ulong %37
- %6 = OpLoad %ulong %7 Aligned 8
- %4 = OpCopyObject %ulong %6
- OpStore %2 %4
- %5 = OpLoad %ulong %2
- OpReturnValue %5
- OpFunctionEnd
- %9 = OpFunction %void None %56
- %17 = OpFunctionParameter %ulong
- %18 = OpFunctionParameter %ulong
- %34 = OpLabel
- %10 = OpVariable %_ptr_Function_ulong Function
- %11 = OpVariable %_ptr_Function_ulong Function
- %13 = OpVariable %_ptr_Function_ulong Function
- %14 = OpVariable %_ptr_Function_ulong Function
- %15 = OpVariable %_ptr_Function_ulong Function
- %16 = OpVariable %_ptr_Function_ulong Function
- OpStore %10 %17
- OpStore %11 %18
- %19 = OpLoad %ulong %10 Aligned 8
- OpStore %13 %19
- %20 = OpLoad %ulong %11 Aligned 8
- OpStore %14 %20
- %22 = OpLoad %ulong %13
- %30 = OpConvertUToPtr %_ptr_Generic_ulong %22
- %21 = OpLoad %ulong %30 Aligned 8
- OpStore %15 %21
- %23 = OpLoad %ulong %15
- %39 = OpBitcast %_ptr_Workgroup__arr_uint_uint_2 %38
- %31 = OpBitcast %_ptr_Workgroup_ulong %39
- OpStore %31 %23 Aligned 8
- %40 = OpBitcast %_ptr_Workgroup_uchar %38
- %32 = OpFunctionCall %ulong %1 %40
- %24 = OpCopyObject %ulong %32
- OpStore %16 %24
- %26 = OpLoad %ulong %16
- %27 = OpLoad %ulong %15
- %25 = OpIAdd %ulong %26 %27
- OpStore %16 %25
- %28 = OpLoad %ulong %14
- %29 = OpLoad %ulong %16
- %33 = OpConvertUToPtr %_ptr_Generic_ulong %28
- OpStore %33 %29 Aligned 8
- OpReturn
- OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/shared_variable.spvtxt b/ptx/src/test/spirv_run/shared_variable.spvtxt
index 7a97c0e..fbbfe4a 100644
--- a/ptx/src/test/spirv_run/shared_variable.spvtxt
+++ b/ptx/src/test/spirv_run/shared_variable.spvtxt
@@ -8,25 +8,26 @@
OpCapability Float16
OpCapability Float64
OpCapability DenormFlushToZero
- %28 = OpExtInstImport "OpenCL.std"
+ OpExtension "SPV_KHR_float_controls"
+ OpExtension "SPV_KHR_no_integer_wrap_decoration"
+ %25 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %1 "shared_variable" %25
+ OpEntryPoint Kernel %1 "shared_variable" %4
OpExecutionMode %1 ContractionOff
+ OpDecorate %4 Alignment 4
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%uint_128 = OpConstant %uint 128
%_arr_uchar_uint_128 = OpTypeArray %uchar %uint_128
%_ptr_Workgroup__arr_uchar_uint_128 = OpTypePointer Workgroup %_arr_uchar_uint_128
- %25 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
+ %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
%ulong = OpTypeInt 64 0
- %36 = OpTypeFunction %void %ulong %ulong
+ %33 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
- %uint_128_0 = OpConstant %uint 128
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
- %uint_128_1 = OpConstant %uint 128
- %1 = OpFunction %void None %36
+ %1 = OpFunction %void None %33
%9 = OpFunctionParameter %ulong
%10 = OpFunctionParameter %ulong
%23 = OpLabel
@@ -47,11 +48,9 @@
%13 = OpLoad %ulong %19 Aligned 8
OpStore %7 %13
%15 = OpLoad %ulong %7
- %26 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_128 %25
- %20 = OpBitcast %_ptr_Workgroup_ulong %26
+ %20 = OpBitcast %_ptr_Workgroup_ulong %4
OpStore %20 %15 Aligned 8
- %27 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_128 %25
- %21 = OpBitcast %_ptr_Workgroup_ulong %27
+ %21 = OpBitcast %_ptr_Workgroup_ulong %4
%16 = OpLoad %ulong %21 Aligned 8
OpStore %8 %16
%17 = OpLoad %ulong %6