Redo shared memory transformation

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