Finish fixing shared memory pass

14 files changed, 329 insertions, 155 deletions

diff --git a/ptx/src/test/spirv_run/atom_add.spvtxt b/ptx/src/test/spirv_run/atom_add.spvtxt
index b4de00a..3609247 100644
--- a/ptx/src/test/spirv_run/atom_add.spvtxt
+++ b/ptx/src/test/spirv_run/atom_add.spvtxt
@@ -7,29 +7,33 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %38 = OpExtInstImport "OpenCL.std"
+               OpCapability DenormFlushToZero
+         %42 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %1 "atom_add" %4
-               OpDecorate %4 Alignment 4
+               OpEntryPoint Kernel %1 "atom_add" %38
+               OpExecutionMode %1 ContractionOff
        %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
-          %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
+         %38 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
       %ulong = OpTypeInt 64 0
-         %46 = OpTypeFunction %void %ulong %ulong
+         %50 = 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 %46
+          %1 = OpFunction %void None %50
           %9 = OpFunctionParameter %ulong
          %10 = OpFunctionParameter %ulong
          %36 = OpLabel
@@ -51,19 +55,22 @@
                OpStore %7 %13
          %16 = OpLoad %ulong %5
          %30 = OpConvertUToPtr %_ptr_Generic_uint %16
-         %51 = OpBitcast %_ptr_Generic_uchar %30
-         %52 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %51 %ulong_4
-         %26 = OpBitcast %_ptr_Generic_uint %52
+         %55 = OpBitcast %_ptr_Generic_uchar %30
+         %56 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %55 %ulong_4
+         %26 = OpBitcast %_ptr_Generic_uint %56
          %15 = OpLoad %uint %26 Aligned 4
                OpStore %8 %15
          %17 = OpLoad %uint %7
-         %31 = OpBitcast %_ptr_Workgroup_uint %4
+         %39 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %38
+         %31 = OpBitcast %_ptr_Workgroup_uint %39
                OpStore %31 %17 Aligned 4
          %19 = OpLoad %uint %8
-         %32 = OpBitcast %_ptr_Workgroup_uint %4
+         %40 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %38
+         %32 = OpBitcast %_ptr_Workgroup_uint %40
          %18 = OpAtomicIAdd %uint %32 %uint_1 %uint_0 %19
                OpStore %7 %18
-         %33 = OpBitcast %_ptr_Workgroup_uint %4
+         %41 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %38
+         %33 = OpBitcast %_ptr_Workgroup_uint %41
          %20 = OpLoad %uint %33 Aligned 4
                OpStore %8 %20
          %21 = OpLoad %ulong %6
@@ -73,9 +80,9 @@
          %23 = OpLoad %ulong %6
          %24 = OpLoad %uint %8
          %35 = OpConvertUToPtr %_ptr_Generic_uint %23
-         %56 = OpBitcast %_ptr_Generic_uchar %35
-         %57 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %56 %ulong_4_0
-         %28 = OpBitcast %_ptr_Generic_uint %57
+         %63 = OpBitcast %_ptr_Generic_uchar %35
+         %64 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %63 %ulong_4_0
+         %28 = OpBitcast %_ptr_Generic_uint %64
                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 7d25632..9533d83 100644
--- a/ptx/src/test/spirv_run/atom_add_float.spvtxt
+++ b/ptx/src/test/spirv_run/atom_add_float.spvtxt
@@ -7,34 +7,38 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %42 = OpExtInstImport "OpenCL.std"
+               OpCapability DenormFlushToZero
+         %46 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %1 "atom_add_float" %4
+               OpEntryPoint Kernel %1 "atom_add_float" %42
+               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
-         %46 = OpTypeFunction %float %_ptr_Workgroup_float %float
+         %50 = 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
-          %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
+         %42 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
       %ulong = OpTypeInt 64 0
-         %53 = OpTypeFunction %void %ulong %ulong
+         %57 = 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 %46
+         %37 = OpFunction %float None %50
          %39 = OpFunctionParameter %_ptr_Workgroup_float
          %40 = OpFunctionParameter %float
                OpFunctionEnd
-          %1 = OpFunction %void None %53
+          %1 = OpFunction %void None %57
           %9 = OpFunctionParameter %ulong
          %10 = OpFunctionParameter %ulong
          %36 = OpLabel
@@ -56,19 +60,22 @@
                OpStore %7 %13
          %16 = OpLoad %ulong %5
          %30 = OpConvertUToPtr %_ptr_Generic_float %16
-         %58 = OpBitcast %_ptr_Generic_uchar %30
-         %59 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %58 %ulong_4
-         %26 = OpBitcast %_ptr_Generic_float %59
+         %62 = OpBitcast %_ptr_Generic_uchar %30
+         %63 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %62 %ulong_4
+         %26 = OpBitcast %_ptr_Generic_float %63
          %15 = OpLoad %float %26 Aligned 4
                OpStore %8 %15
          %17 = OpLoad %float %7
-         %31 = OpBitcast %_ptr_Workgroup_float %4
+         %43 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %42
+         %31 = OpBitcast %_ptr_Workgroup_float %43
                OpStore %31 %17 Aligned 4
          %19 = OpLoad %float %8
-         %32 = OpBitcast %_ptr_Workgroup_float %4
+         %44 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %42
+         %32 = OpBitcast %_ptr_Workgroup_float %44
          %18 = OpFunctionCall %float %37 %32 %19
                OpStore %7 %18
-         %33 = OpBitcast %_ptr_Workgroup_float %4
+         %45 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_1024 %42
+         %33 = OpBitcast %_ptr_Workgroup_float %45
          %20 = OpLoad %float %33 Aligned 4
                OpStore %8 %20
          %21 = OpLoad %ulong %6
@@ -78,9 +85,9 @@
          %23 = OpLoad %ulong %6
          %24 = OpLoad %float %8
          %35 = OpConvertUToPtr %_ptr_Generic_float %23
-         %60 = OpBitcast %_ptr_Generic_uchar %35
-         %61 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %60 %ulong_4_0
-         %28 = OpBitcast %_ptr_Generic_float %61
+         %67 = OpBitcast %_ptr_Generic_uchar %35
+         %68 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %67 %ulong_4_0
+         %28 = OpBitcast %_ptr_Generic_float %68
                OpStore %28 %24 Aligned 4
                OpReturn
                OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/call.spvtxt b/ptx/src/test/spirv_run/call.spvtxt
index 6929b1e..c29984e 100644
--- a/ptx/src/test/spirv_run/call.spvtxt
+++ b/ptx/src/test/spirv_run/call.spvtxt
@@ -7,9 +7,13 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
+               OpCapability DenormFlushToZero
          %37 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %4 "call"
+               OpExecutionMode %4 ContractionOff
+               OpDecorate %4 LinkageAttributes "call" Export
+               OpDecorate %1 LinkageAttributes "incr" Export
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
          %40 = OpTypeFunction %void %ulong %ulong
diff --git a/ptx/src/test/spirv_run/extern_shared.spvtxt b/ptx/src/test/spirv_run/extern_shared.spvtxt
index ed1c489..82d86ae 100644
--- a/ptx/src/test/spirv_run/extern_shared.spvtxt
+++ b/ptx/src/test/spirv_run/extern_shared.spvtxt
@@ -7,21 +7,23 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %24 = OpExtInstImport "OpenCL.std"
+               OpCapability DenormFlushToZero
+         %27 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %2 "extern_shared" %1
+               OpEntryPoint Kernel %2 "extern_shared" %24
                OpExecutionMode %2 ContractionOff
-               OpDecorate %1 LinkageAttributes "shared_mem" Import
        %void = OpTypeVoid
-       %uint = OpTypeInt 32 0
-%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
-          %1 = OpVariable %_ptr_Workgroup_uint Workgroup
+      %uchar = OpTypeInt 8 0
+%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
+         %24 = OpVariable %_ptr_Workgroup_uchar Workgroup
       %ulong = OpTypeInt 64 0
-         %29 = OpTypeFunction %void %ulong %ulong
+         %32 = 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 %29
+          %2 = OpFunction %void None %32
           %8 = OpFunctionParameter %ulong
           %9 = OpFunctionParameter %ulong
          %22 = OpLabel
@@ -41,9 +43,11 @@
          %12 = OpLoad %ulong %18 Aligned 8
                OpStore %7 %12
          %14 = OpLoad %ulong %7
-         %19 = OpBitcast %_ptr_Workgroup_ulong %1
+         %25 = OpBitcast %_ptr_Workgroup_uint %24
+         %19 = OpBitcast %_ptr_Workgroup_ulong %25
                OpStore %19 %14 Aligned 8
-         %20 = OpBitcast %_ptr_Workgroup_ulong %1
+         %26 = OpBitcast %_ptr_Workgroup_uint %24
+         %20 = OpBitcast %_ptr_Workgroup_ulong %26
          %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 941eb39..3cc78cb 100644
--- a/ptx/src/test/spirv_run/extern_shared_call.spvtxt
+++ b/ptx/src/test/spirv_run/extern_shared_call.spvtxt
@@ -7,38 +7,42 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %34 = OpExtInstImport "OpenCL.std"
+               OpCapability DenormFlushToZero
+         %41 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %12 "extern_shared_call" %1
+               OpEntryPoint Kernel %12 "extern_shared_call" %37
                OpExecutionMode %12 ContractionOff
-               OpDecorate %1 Alignment 4
-               OpDecorate %1 LinkageAttributes "shared_mem" Import
        %void = OpTypeVoid
-       %uint = OpTypeInt 32 0
-%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
-          %1 = OpVariable %_ptr_Workgroup_uint Workgroup
-         %38 = OpTypeFunction %void
+      %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
 %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
     %ulong_2 = OpConstant %ulong 2
-         %42 = OpTypeFunction %void %ulong %ulong
+         %37 = OpVariable %_ptr_Workgroup_uchar Workgroup
+         %51 = OpTypeFunction %void %ulong %ulong
 %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
-          %2 = OpFunction %void None %38
+          %2 = OpFunction %void None %45
+         %34 = OpFunctionParameter %_ptr_Workgroup_uchar
          %11 = OpLabel
           %3 = OpVariable %_ptr_Function_ulong Function
-          %9 = OpBitcast %_ptr_Workgroup_ulong %1
+         %35 = OpBitcast %_ptr_Workgroup_uint %34
+          %9 = OpBitcast %_ptr_Workgroup_ulong %35
           %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
-         %10 = OpBitcast %_ptr_Workgroup_ulong %1
+         %36 = OpBitcast %_ptr_Workgroup_uint %34
+         %10 = OpBitcast %_ptr_Workgroup_ulong %36
                OpStore %10 %7 Aligned 8
                OpReturn
                OpFunctionEnd
-         %12 = OpFunction %void None %42
+         %12 = OpFunction %void None %51
          %18 = OpFunctionParameter %ulong
          %19 = OpFunctionParameter %ulong
          %32 = OpLabel
@@ -58,10 +62,13 @@
          %22 = OpLoad %ulong %28 Aligned 8
                OpStore %17 %22
          %24 = OpLoad %ulong %17
-         %29 = OpBitcast %_ptr_Workgroup_ulong %1
+         %38 = OpBitcast %_ptr_Workgroup_uint %37
+         %29 = OpBitcast %_ptr_Workgroup_ulong %38
                OpStore %29 %24 Aligned 8
-         %44 = OpFunctionCall %void %2
-         %30 = OpBitcast %_ptr_Workgroup_ulong %1
+         %39 = OpBitcast %_ptr_Workgroup_uchar %37
+         %53 = OpFunctionCall %void %2 %39
+         %40 = OpBitcast %_ptr_Workgroup_uint %37
+         %30 = OpBitcast %_ptr_Workgroup_ulong %40
          %25 = OpLoad %ulong %30 Aligned 8
                OpStore %17 %25
          %26 = OpLoad %ulong %16
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index b7fd386..6c073f3 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -222,6 +222,7 @@ 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!(func_ptr);

 test_ptx!(lanemask_lt);

diff --git a/ptx/src/test/spirv_run/reg_local.spvtxt b/ptx/src/test/spirv_run/reg_local.spvtxt
index 4a69450..ddb6a9e 100644
--- a/ptx/src/test/spirv_run/reg_local.spvtxt
+++ b/ptx/src/test/spirv_run/reg_local.spvtxt
@@ -7,6 +7,7 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
+               OpCapability DenormFlushToZero
          %34 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %1 "reg_local"
@@ -51,22 +52,24 @@
                OpStore %7 %12
          %14 = OpLoad %ulong %7
          %19 = OpIAdd %ulong %14 %ulong_1
-         %26 = OpPtrCastToGeneric %_ptr_Generic_ulong %4
+         %46 = OpBitcast %_ptr_Function_ulong %4
+         %26 = OpPtrCastToGeneric %_ptr_Generic_ulong %46
          %27 = OpCopyObject %ulong %19
                OpStore %26 %27 Aligned 8
-         %28 = OpPtrCastToGeneric %_ptr_Generic_ulong %4
-         %47 = OpBitcast %_ptr_Generic_uchar %28
-         %48 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %47 %ulong_0
-         %21 = OpBitcast %_ptr_Generic_ulong %48
+         %47 = OpBitcast %_ptr_Function_ulong %4
+         %28 = OpPtrCastToGeneric %_ptr_Generic_ulong %47
+         %49 = OpBitcast %_ptr_Generic_uchar %28
+         %50 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %49 %ulong_0
+         %21 = OpBitcast %_ptr_Generic_ulong %50
          %29 = OpLoad %ulong %21 Aligned 8
          %15 = OpCopyObject %ulong %29
                OpStore %7 %15
          %16 = OpLoad %ulong %6
          %17 = OpLoad %ulong %7
          %30 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %16
-         %50 = OpBitcast %_ptr_CrossWorkgroup_uchar %30
-         %51 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %50 %ulong_0_0
-         %23 = OpBitcast %_ptr_CrossWorkgroup_ulong %51
+         %52 = OpBitcast %_ptr_CrossWorkgroup_uchar %30
+         %53 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %52 %ulong_0_0
+         %23 = OpBitcast %_ptr_CrossWorkgroup_ulong %53
          %31 = OpCopyObject %ulong %17
                OpStore %23 %31 Aligned 8
                OpReturn
diff --git a/ptx/src/test/spirv_run/shared_ptr_32.spvtxt b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt
index 1b2e3dd..020c15b 100644
--- a/ptx/src/test/spirv_run/shared_ptr_32.spvtxt
+++ b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt
@@ -7,26 +7,28 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %32 = OpExtInstImport "OpenCL.std"
+               OpCapability DenormFlushToZero
+         %34 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %1 "shared_ptr_32" %4
-               OpDecorate %4 Alignment 4
+               OpEntryPoint Kernel %1 "shared_ptr_32" %32
+               OpExecutionMode %1 ContractionOff
        %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
-          %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
+         %32 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
       %ulong = OpTypeInt 64 0
-         %40 = OpTypeFunction %void %ulong %ulong
+         %42 = 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 %40
+          %1 = OpFunction %void None %42
          %10 = OpFunctionParameter %ulong
          %11 = OpFunctionParameter %ulong
          %30 = OpLabel
@@ -43,7 +45,8 @@
                OpStore %5 %12
          %13 = OpLoad %ulong %3 Aligned 8
                OpStore %6 %13
-         %25 = OpConvertPtrToU %uint %4
+         %33 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_128 %32
+         %25 = OpConvertPtrToU %uint %33
          %14 = OpCopyObject %uint %25
                OpStore %7 %14
          %16 = OpLoad %ulong %5
@@ -56,9 +59,9 @@
                OpStore %27 %18 Aligned 8
          %20 = OpLoad %uint %7
          %28 = OpConvertUToPtr %_ptr_Workgroup_ulong %20
-         %46 = OpBitcast %_ptr_Workgroup_uchar %28
-         %47 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %46 %ulong_0
-         %24 = OpBitcast %_ptr_Workgroup_ulong %47
+         %49 = OpBitcast %_ptr_Workgroup_uchar %28
+         %50 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uchar %49 %ulong_0
+         %24 = OpBitcast %_ptr_Workgroup_ulong %50
          %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 3ebe810..90e04f3 100644
--- a/ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt
+++ b/ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt
@@ -7,22 +7,21 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %30 = OpExtInstImport "OpenCL.std"
+               OpCapability DenormFlushToZero
+         %32 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %2 "shared_ptr_take_address" %1
+               OpEntryPoint Kernel %2 "shared_ptr_take_address" %30
                OpExecutionMode %2 ContractionOff
-               OpDecorate %1 Alignment 4
-               OpDecorate %1 LinkageAttributes "shared_mem" Import
        %void = OpTypeVoid
       %uchar = OpTypeInt 8 0
 %_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
-          %1 = OpVariable %_ptr_Workgroup_uchar Workgroup
+         %30 = OpVariable %_ptr_Workgroup_uchar Workgroup
       %ulong = OpTypeInt 64 0
-         %35 = OpTypeFunction %void %ulong %ulong
+         %37 = 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 %35
+          %2 = OpFunction %void None %37
          %10 = OpFunctionParameter %ulong
          %11 = OpFunctionParameter %ulong
          %28 = OpLabel
@@ -39,7 +38,8 @@
                OpStore %5 %12
          %13 = OpLoad %ulong %4 Aligned 8
                OpStore %6 %13
-         %23 = OpConvertPtrToU %ulong %1
+         %31 = OpBitcast %_ptr_Workgroup_uchar %30
+         %23 = OpConvertPtrToU %ulong %31
          %14 = OpCopyObject %ulong %23
                OpStore %7 %14
          %16 = OpLoad %ulong %5
diff --git a/ptx/src/test/spirv_run/shared_unify_extern.spvtxt b/ptx/src/test/spirv_run/shared_unify_extern.spvtxt
index 9b62045..2dd2056 100644
--- a/ptx/src/test/spirv_run/shared_unify_extern.spvtxt
+++ b/ptx/src/test/spirv_run/shared_unify_extern.spvtxt
@@ -1,62 +1,79 @@
-               OpCapability GenericPointer

-               OpCapability Linkage

-               OpCapability Addresses

-               OpCapability Kernel

-               OpCapability Int8

-               OpCapability Int16

-               OpCapability Int64

-               OpCapability Float16

-               OpCapability Float64

-         %30 = OpExtInstImport "OpenCL.std"

-               OpMemoryModel Physical64 OpenCL

-               OpEntryPoint Kernel %2 "shared_ptr_take_address" %1

-               OpExecutionMode %2 ContractionOff

-               OpDecorate %1 Alignment 4

-               OpDecorate %1 LinkageAttributes "shared_mem" Import

-       %void = OpTypeVoid

-      %uchar = OpTypeInt 8 0

-%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar

-          %1 = OpVariable %_ptr_Workgroup_uchar Workgroup

-      %ulong = OpTypeInt 64 0

-         %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 %35

-         %10 = OpFunctionParameter %ulong

-         %11 = OpFunctionParameter %ulong

-         %28 = OpLabel

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

-          %9 = OpVariable %_ptr_Function_ulong Function

-               OpStore %3 %10

-               OpStore %4 %11

-         %12 = OpLoad %ulong %3 Aligned 8

-               OpStore %5 %12

-         %13 = OpLoad %ulong %4 Aligned 8

-               OpStore %6 %13

-         %23 = OpConvertPtrToU %ulong %1

-         %14 = OpCopyObject %ulong %23

-               OpStore %7 %14

-         %16 = OpLoad %ulong %5

-         %24 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %16

-         %15 = OpLoad %ulong %24 Aligned 8

-               OpStore %8 %15

-         %17 = OpLoad %ulong %7

-         %18 = OpLoad %ulong %8

-         %25 = OpConvertUToPtr %_ptr_Workgroup_ulong %17

-               OpStore %25 %18 Aligned 8

-         %20 = OpLoad %ulong %7

-         %26 = OpConvertUToPtr %_ptr_Workgroup_ulong %20

-         %19 = OpLoad %ulong %26 Aligned 8

-               OpStore %9 %19

-         %21 = OpLoad %ulong %6

-         %22 = OpLoad %ulong %9

-         %27 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %21

-               OpStore %27 %22 Aligned 8

-               OpReturn

-               OpFunctionEnd

+               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 %10 "shared_unify_extern" %38
+               OpExecutionMode %10 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
+         %38 = OpVariable %_ptr_Workgroup_uchar Workgroup
+         %53 = OpTypeFunction %void %ulong %ulong
+%_ptr_Generic_ulong = OpTypePointer Generic %ulong
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+          %3 = OpFunction %ulong None %46
+         %36 = OpFunctionParameter %_ptr_Workgroup_uchar
+          %9 = 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
+               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
+               OpReturn
+               OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/shared_unify_private.ptx b/ptx/src/test/spirv_run/shared_unify_private.ptx
new file mode 100644
index 0000000..fd31357
--- /dev/null
+++ b/ptx/src/test/spirv_run/shared_unify_private.ptx
@@ -0,0 +1,32 @@
+.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
new file mode 100644
index 0000000..69bf018
--- /dev/null
+++ b/ptx/src/test/spirv_run/shared_unify_private.spvtxt
@@ -0,0 +1,84 @@
+               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 49278a8..7a97c0e 100644
--- a/ptx/src/test/spirv_run/shared_variable.spvtxt
+++ b/ptx/src/test/spirv_run/shared_variable.spvtxt
@@ -7,23 +7,26 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %25 = OpExtInstImport "OpenCL.std"
+               OpCapability DenormFlushToZero
+         %28 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %1 "shared_variable" %4
-               OpDecorate %4 Alignment 4
+               OpEntryPoint Kernel %1 "shared_variable" %25
+               OpExecutionMode %1 ContractionOff
        %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
-          %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
+         %25 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup
       %ulong = OpTypeInt 64 0
-         %33 = OpTypeFunction %void %ulong %ulong
+         %36 = 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
-          %1 = OpFunction %void None %33
+ %uint_128_1 = OpConstant %uint 128
+          %1 = OpFunction %void None %36
           %9 = OpFunctionParameter %ulong
          %10 = OpFunctionParameter %ulong
          %23 = OpLabel
@@ -44,9 +47,11 @@
          %13 = OpLoad %ulong %19 Aligned 8
                OpStore %7 %13
          %15 = OpLoad %ulong %7
-         %20 = OpBitcast %_ptr_Workgroup_ulong %4
+         %26 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_128 %25
+         %20 = OpBitcast %_ptr_Workgroup_ulong %26
                OpStore %20 %15 Aligned 8
-         %21 = OpBitcast %_ptr_Workgroup_ulong %4
+         %27 = OpBitcast %_ptr_Workgroup__arr_uchar_uint_128 %25
+         %21 = OpBitcast %_ptr_Workgroup_ulong %27
          %16 = OpLoad %ulong %21 Aligned 8
                OpStore %8 %16
          %17 = OpLoad %ulong %6
diff --git a/ptx/src/test/spirv_run/verify.py b/ptx/src/test/spirv_run/verify.py
index dbfab00..4ef6465 100644
--- a/ptx/src/test/spirv_run/verify.py
+++ b/ptx/src/test/spirv_run/verify.py
@@ -1,7 +1,7 @@
 import os, sys, subprocess

 

 def main(path):

-    dirs = os.listdir(path)

+    dirs = sorted(os.listdir(path))

     for file in dirs:

         if not file.endswith(".spvtxt"):

             continue