Update tests, disable OpenCL-style shared mem conversion, emit linking information

16 files changed, 270 insertions, 252 deletions

diff --git a/ptx/src/test/spirv_run/cos.spvtxt b/ptx/src/test/spirv_run/cos.spvtxt
index 8d6a0ca..a79cdbe 100644
--- a/ptx/src/test/spirv_run/cos.spvtxt
+++ b/ptx/src/test/spirv_run/cos.spvtxt
@@ -10,6 +10,7 @@
          %21 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %1 "cos"
+               OpExecutionMode %1 ContractionOff
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
          %24 = OpTypeFunction %void %ulong %ulong
@@ -37,7 +38,7 @@
          %11 = OpLoad %float %17 Aligned 4
                OpStore %6 %11
          %14 = OpLoad %float %6
-         %13 = OpExtInst %float %21 native_cos %14
+         %13 = OpExtInst %float %21 cos %14
                OpStore %6 %13
          %15 = OpLoad %ulong %5
          %16 = OpLoad %float %6
diff --git a/ptx/src/test/spirv_run/ex2.spvtxt b/ptx/src/test/spirv_run/ex2.spvtxt
index 3d7b58d..29e5e86 100644
--- a/ptx/src/test/spirv_run/ex2.spvtxt
+++ b/ptx/src/test/spirv_run/ex2.spvtxt
@@ -10,6 +10,7 @@
          %21 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %1 "ex2"
+               OpExecutionMode %1 ContractionOff
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
          %24 = OpTypeFunction %void %ulong %ulong
@@ -37,7 +38,7 @@
          %11 = OpLoad %float %17 Aligned 4
                OpStore %6 %11
          %14 = OpLoad %float %6
-         %13 = OpExtInst %float %21 native_exp2 %14
+         %13 = OpExtInst %float %21 exp2 %14
                OpStore %6 %13
          %15 = OpLoad %ulong %5
          %16 = OpLoad %float %6
diff --git a/ptx/src/test/spirv_run/extern_shared.spvtxt b/ptx/src/test/spirv_run/extern_shared.spvtxt
index 13587d5..ed1c489 100644
--- a/ptx/src/test/spirv_run/extern_shared.spvtxt
+++ b/ptx/src/test/spirv_run/extern_shared.spvtxt
@@ -7,24 +7,23 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %27 = OpExtInstImport "OpenCL.std"
+         %24 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %2 "extern_shared" %1
+               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
       %ulong = OpTypeInt 64 0
-      %uchar = OpTypeInt 8 0
-%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
-         %34 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar
+         %29 = 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 %34
+          %2 = OpFunction %void None %29
           %8 = OpFunctionParameter %ulong
           %9 = OpFunctionParameter %ulong
-         %24 = OpFunctionParameter %_ptr_Workgroup_uchar
          %22 = OpLabel
           %3 = OpVariable %_ptr_Function_ulong Function
           %4 = OpVariable %_ptr_Function_ulong Function
@@ -42,11 +41,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 5af7168..941eb39 100644
--- a/ptx/src/test/spirv_run/extern_shared_call.spvtxt
+++ b/ptx/src/test/spirv_run/extern_shared_call.spvtxt
@@ -7,44 +7,40 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %40 = OpExtInstImport "OpenCL.std"
+         %34 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %12 "extern_shared_call" %1
+               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
-      %uchar = OpTypeInt 8 0
-%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
-         %46 = OpTypeFunction %void %_ptr_Workgroup_uchar
+         %38 = OpTypeFunction %void
       %ulong = OpTypeInt 64 0
 %_ptr_Function_ulong = OpTypePointer Function %ulong
 %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
     %ulong_2 = OpConstant %ulong 2
-         %50 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar
+         %42 = OpTypeFunction %void %ulong %ulong
 %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
-          %2 = OpFunction %void None %46
-         %34 = OpFunctionParameter %_ptr_Workgroup_uchar
+          %2 = OpFunction %void None %38
          %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 %1
           %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 %1
                OpStore %10 %7 Aligned 8
                OpReturn
                OpFunctionEnd
-         %12 = OpFunction %void None %50
+         %12 = OpFunction %void None %42
          %18 = OpFunctionParameter %ulong
          %19 = OpFunctionParameter %ulong
-         %37 = OpFunctionParameter %_ptr_Workgroup_uchar
          %32 = OpLabel
          %13 = OpVariable %_ptr_Function_ulong Function
          %14 = OpVariable %_ptr_Function_ulong Function
@@ -62,12 +58,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
-         %52 = OpFunctionCall %void %2 %37
-         %39 = OpBitcast %_ptr_Workgroup_uint %37
-         %30 = OpBitcast %_ptr_Workgroup_ulong %39
+         %44 = OpFunctionCall %void %2
+         %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/lg2.spvtxt b/ptx/src/test/spirv_run/lg2.spvtxt
index c30eeff..a8175cf 100644
--- a/ptx/src/test/spirv_run/lg2.spvtxt
+++ b/ptx/src/test/spirv_run/lg2.spvtxt
@@ -10,6 +10,7 @@
          %21 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %1 "lg2"
+               OpExecutionMode %1 ContractionOff
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
          %24 = OpTypeFunction %void %ulong %ulong
@@ -37,7 +38,7 @@
          %11 = OpLoad %float %17 Aligned 4
                OpStore %6 %11
          %14 = OpLoad %float %6
-         %13 = OpExtInst %float %21 native_log2 %14
+         %13 = OpExtInst %float %21 log2 %14
                OpStore %6 %13
          %15 = OpLoad %ulong %5
          %16 = OpLoad %float %6
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 798fff2..1bb6ab7 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -32,6 +32,7 @@ use std::io;
 use std::io::Read;

 use std::io::Write;

 use std::mem;

+use std::path::Path;

 use std::process::Command;

 use std::slice;

 use std::{borrow::Cow, collections::HashMap, env, fs, path::PathBuf, ptr, str};

@@ -292,7 +293,7 @@ fn run_spirv<Input: From<u8> + Copy + Debug, Output: From<u8> + Copy + Debug + D
         hip_call! { hipMemcpyWithStream(inp_b, input.as_ptr() as _, input.len() * mem::size_of::<Input>(), hipMemcpyKind::hipMemcpyHostToDevice, stream) };

         hip_call! { hipMemset(out_b, 0, output.len() * mem::size_of::<Output>()) };

         let mut args = [&inp_b, &out_b];

-        hip_call! { hipModuleLaunchKernel(kernel, 1,1,1,1,1,1, 0, stream, args.as_mut_ptr() as _, ptr::null_mut()) };

+        hip_call! { hipModuleLaunchKernel(kernel, 1,1,1,1,1,1, 1024, stream, args.as_mut_ptr() as _, ptr::null_mut()) };

         hip_call! { hipMemcpyAsync(result.as_mut_ptr() as _, out_b, output.len() * mem::size_of::<Output>(), hipMemcpyKind::hipMemcpyDeviceToHost, stream) };

         hip_call! { hipStreamSynchronize(stream) };

     }

@@ -600,6 +601,9 @@ fn compile_amd(
         .arg(spirv.path())

         .status()?;

     assert!(to_llvm_cmd.success());

+    if cfg!(debug_assertions) {

+        persist_file(llvm.path())?;

+    }

     let linked_binary = NamedTempFile::new_in(&dir)?;

     let mut llvm_link = PathBuf::from(AMDGPU);

     llvm_link.push("llvm");

@@ -617,6 +621,9 @@ fn compile_amd(
     }

     let status = linker_cmd.status()?;

     assert!(status.success());

+    if cfg!(debug_assertions) {

+        persist_file(linked_binary.path())?;

+    }

     let mut ptx_lib_bitcode = NamedTempFile::new_in(&dir)?;

     let compiled_binary = NamedTempFile::new_in(&dir)?;

     let mut clang_exe = PathBuf::from(AMDGPU);

@@ -651,11 +658,18 @@ fn compile_amd(
     let compiled_bin_path = compiled_binary.path();

     let mut compiled_binary = File::open(compiled_bin_path)?;

     compiled_binary.read_to_end(&mut result)?;

+    if cfg!(debug_assertions) {

+        persist_file(compiled_bin_path)?;

+    }

+    Ok(result)

+}

+

+fn persist_file(path: &Path) -> io::Result<()> {

     let mut persistent = PathBuf::from("/tmp/zluda");

     std::fs::create_dir_all(&persistent)?;

-    persistent.push(compiled_bin_path.file_name().unwrap());

-    std::fs::copy(compiled_bin_path, persistent)?;

-    Ok(result)

+    persistent.push(path.file_name().unwrap());

+    std::fs::copy(path, persistent)?;

+    Ok(())

 }

 

 fn get_bitcode_paths(device_name: &str) -> impl Iterator<Item = PathBuf> {

diff --git a/ptx/src/test/spirv_run/reg_local.spvtxt b/ptx/src/test/spirv_run/reg_local.spvtxt
index a0b957a..4a69450 100644
--- a/ptx/src/test/spirv_run/reg_local.spvtxt
+++ b/ptx/src/test/spirv_run/reg_local.spvtxt
@@ -10,6 +10,7 @@
          %34 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %1 "reg_local"
+               OpExecutionMode %1 ContractionOff
                OpDecorate %4 Alignment 8
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
@@ -50,10 +51,10 @@
                OpStore %7 %12
          %14 = OpLoad %ulong %7
          %19 = OpIAdd %ulong %14 %ulong_1
-         %26 = OpBitcast %_ptr_Generic_ulong %4
+         %26 = OpPtrCastToGeneric %_ptr_Generic_ulong %4
          %27 = OpCopyObject %ulong %19
                OpStore %26 %27 Aligned 8
-         %28 = OpBitcast %_ptr_Generic_ulong %4
+         %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
diff --git a/ptx/src/test/spirv_run/rsqrt.spvtxt b/ptx/src/test/spirv_run/rsqrt.spvtxt
index fc1a7e1..6c87113 100644
--- a/ptx/src/test/spirv_run/rsqrt.spvtxt
+++ b/ptx/src/test/spirv_run/rsqrt.spvtxt
@@ -10,6 +10,7 @@
          %21 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %1 "rsqrt"
+               OpExecutionMode %1 ContractionOff
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
          %24 = OpTypeFunction %void %ulong %ulong
@@ -37,7 +38,7 @@
          %11 = OpLoad %double %17 Aligned 8
                OpStore %6 %11
          %14 = OpLoad %double %6
-         %13 = OpExtInst %double %21 native_rsqrt %14
+         %13 = OpExtInst %double %21 rsqrt %14
                OpStore %6 %13
          %15 = OpLoad %ulong %5
          %16 = OpLoad %double %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 fd4f893..3ebe810 100644
--- a/ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt
+++ b/ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt
@@ -7,23 +7,24 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %31 = OpExtInstImport "OpenCL.std"
+         %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
-         %36 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar
+         %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 %36
+          %2 = OpFunction %void None %35
          %10 = OpFunctionParameter %ulong
          %11 = OpFunctionParameter %ulong
-         %30 = OpFunctionParameter %_ptr_Workgroup_uchar
          %28 = OpLabel
           %3 = OpVariable %_ptr_Function_ulong Function
           %4 = OpVariable %_ptr_Function_ulong Function
@@ -38,7 +39,7 @@
                OpStore %5 %12
          %13 = OpLoad %ulong %4 Aligned 8
                OpStore %6 %13
-         %23 = OpConvertPtrToU %ulong %30
+         %23 = OpConvertPtrToU %ulong %1
          %14 = OpCopyObject %ulong %23
                OpStore %7 %14
          %16 = OpLoad %ulong %5
diff --git a/ptx/src/test/spirv_run/sin.spvtxt b/ptx/src/test/spirv_run/sin.spvtxt
index 02eba40..6dd3e53 100644
--- a/ptx/src/test/spirv_run/sin.spvtxt
+++ b/ptx/src/test/spirv_run/sin.spvtxt
@@ -10,6 +10,7 @@
          %21 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %1 "sin"
+               OpExecutionMode %1 ContractionOff
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
          %24 = OpTypeFunction %void %ulong %ulong
@@ -37,7 +38,7 @@
          %11 = OpLoad %float %17 Aligned 4
                OpStore %6 %11
          %14 = OpLoad %float %6
-         %13 = OpExtInst %float %21 native_sin %14
+         %13 = OpExtInst %float %21 sin %14
                OpStore %6 %13
          %15 = OpLoad %ulong %5
          %16 = OpLoad %float %6
diff --git a/ptx/src/test/spirv_run/sqrt.spvtxt b/ptx/src/test/spirv_run/sqrt.spvtxt
index 17f223d..1c65aa3 100644
--- a/ptx/src/test/spirv_run/sqrt.spvtxt
+++ b/ptx/src/test/spirv_run/sqrt.spvtxt
@@ -10,6 +10,7 @@
          %21 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %1 "sqrt"
+               OpExecutionMode %1 ContractionOff
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
          %24 = OpTypeFunction %void %ulong %ulong
@@ -37,7 +38,7 @@
          %11 = OpLoad %float %17 Aligned 4
                OpStore %6 %11
          %14 = OpLoad %float %6
-         %13 = OpExtInst %float %21 native_sqrt %14
+         %13 = OpExtInst %float %21 sqrt %14
                OpStore %6 %13
          %15 = OpLoad %ulong %5
          %16 = OpLoad %float %6
diff --git a/ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt b/ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt
index cf0d86e..b99fb50 100644
--- a/ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt
+++ b/ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt
@@ -7,27 +7,30 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %54 = OpExtInstImport "OpenCL.std"
+         %57 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %1 "stateful_ld_st_ntid" %gl_LocalInvocationID
-               OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
+               OpEntryPoint Kernel %1 "stateful_ld_st_ntid"
+               OpExecutionMode %1 ContractionOff
+               OpDecorate %44 LinkageAttributes "_Z12get_local_idj" Import
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
-    %v3ulong = OpTypeVector %ulong 3
-%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
-%gl_LocalInvocationID = OpVariable %_ptr_Input_v3ulong Input
+       %uint = OpTypeInt 32 0
+         %61 = OpTypeFunction %ulong %uint
       %uchar = OpTypeInt 8 0
 %_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
-         %61 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
+         %64 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
 %_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
-       %uint = OpTypeInt 32 0
 %_ptr_Function_uint = OpTypePointer Function %uint
 %_ptr_Function_ulong = OpTypePointer Function %ulong
+     %uint_0 = OpConstant %uint 0
 %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
-          %1 = OpFunction %void None %61
+         %44 = OpFunction %ulong None %61
+         %46 = OpFunctionParameter %uint
+               OpFunctionEnd
+          %1 = OpFunction %void None %64
          %20 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
          %21 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
-         %52 = OpLabel
+         %55 = OpLabel
          %12 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
          %13 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
          %10 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
@@ -37,14 +40,14 @@
           %8 = OpVariable %_ptr_Function_ulong Function
                OpStore %12 %20
                OpStore %13 %21
-         %45 = OpBitcast %_ptr_Function_ulong %12
-         %44 = OpLoad %ulong %45 Aligned 8
-         %14 = OpCopyObject %ulong %44
+         %48 = OpBitcast %_ptr_Function_ulong %12
+         %47 = OpLoad %ulong %48 Aligned 8
+         %14 = OpCopyObject %ulong %47
          %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %14
                OpStore %10 %22
-         %47 = OpBitcast %_ptr_Function_ulong %13
-         %46 = OpLoad %ulong %47 Aligned 8
-         %15 = OpCopyObject %ulong %46
+         %50 = OpBitcast %_ptr_Function_ulong %13
+         %49 = OpLoad %ulong %50 Aligned 8
+         %15 = OpCopyObject %ulong %49
          %23 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %15
                OpStore %11 %23
          %24 = OpLoad %_ptr_CrossWorkgroup_uchar %10
@@ -57,37 +60,36 @@
          %18 = OpCopyObject %ulong %19
          %27 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %18
                OpStore %11 %27
-         %66 = OpLoad %v3ulong %gl_LocalInvocationID
-         %43 = OpCompositeExtract %ulong %66 0
-         %67 = OpBitcast %ulong %43
-         %29 = OpUConvert %uint %67
+         %43 = OpFunctionCall %ulong %44 %uint_0
+         %68 = OpBitcast %ulong %43
+         %29 = OpUConvert %uint %68
          %28 = OpCopyObject %uint %29
                OpStore %6 %28
          %31 = OpLoad %uint %6
-         %68 = OpBitcast %uint %31
-         %30 = OpUConvert %ulong %68
+         %69 = OpBitcast %uint %31
+         %30 = OpUConvert %ulong %69
                OpStore %7 %30
          %33 = OpLoad %_ptr_CrossWorkgroup_uchar %10
          %34 = OpLoad %ulong %7
-         %48 = OpCopyObject %ulong %34
-         %69 = OpBitcast %_ptr_CrossWorkgroup_uchar %33
-         %70 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %69 %48
-         %32 = OpBitcast %_ptr_CrossWorkgroup_uchar %70
+         %51 = OpCopyObject %ulong %34
+         %70 = OpBitcast %_ptr_CrossWorkgroup_uchar %33
+         %71 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %70 %51
+         %32 = OpBitcast %_ptr_CrossWorkgroup_uchar %71
                OpStore %10 %32
          %36 = OpLoad %_ptr_CrossWorkgroup_uchar %11
          %37 = OpLoad %ulong %7
-         %49 = OpCopyObject %ulong %37
-         %71 = OpBitcast %_ptr_CrossWorkgroup_uchar %36
-         %72 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %71 %49
-         %35 = OpBitcast %_ptr_CrossWorkgroup_uchar %72
+         %52 = OpCopyObject %ulong %37
+         %72 = OpBitcast %_ptr_CrossWorkgroup_uchar %36
+         %73 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %72 %52
+         %35 = OpBitcast %_ptr_CrossWorkgroup_uchar %73
                OpStore %11 %35
          %39 = OpLoad %_ptr_CrossWorkgroup_uchar %10
-         %50 = OpBitcast %_ptr_CrossWorkgroup_ulong %39
-         %38 = OpLoad %ulong %50 Aligned 8
+         %53 = OpBitcast %_ptr_CrossWorkgroup_ulong %39
+         %38 = OpLoad %ulong %53 Aligned 8
                OpStore %8 %38
          %40 = OpLoad %_ptr_CrossWorkgroup_uchar %11
          %41 = OpLoad %ulong %8
-         %51 = OpBitcast %_ptr_CrossWorkgroup_ulong %40
-               OpStore %51 %41 Aligned 8
+         %54 = OpBitcast %_ptr_CrossWorkgroup_ulong %40
+               OpStore %54 %41 Aligned 8
                OpReturn
                OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/stateful_ld_st_ntid_chain.spvtxt b/ptx/src/test/spirv_run/stateful_ld_st_ntid_chain.spvtxt
index 97bf000..0239632 100644
--- a/ptx/src/test/spirv_run/stateful_ld_st_ntid_chain.spvtxt
+++ b/ptx/src/test/spirv_run/stateful_ld_st_ntid_chain.spvtxt
@@ -7,27 +7,30 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %62 = OpExtInstImport "OpenCL.std"
+         %65 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %1 "stateful_ld_st_ntid_chain" %gl_LocalInvocationID
-               OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
+               OpEntryPoint Kernel %1 "stateful_ld_st_ntid_chain"
+               OpExecutionMode %1 ContractionOff
+               OpDecorate %52 LinkageAttributes "_Z12get_local_idj" Import
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
-    %v3ulong = OpTypeVector %ulong 3
-%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
-%gl_LocalInvocationID = OpVariable %_ptr_Input_v3ulong Input
+       %uint = OpTypeInt 32 0
+         %69 = OpTypeFunction %ulong %uint
       %uchar = OpTypeInt 8 0
 %_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
-         %69 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
+         %72 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
 %_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
-       %uint = OpTypeInt 32 0
 %_ptr_Function_uint = OpTypePointer Function %uint
 %_ptr_Function_ulong = OpTypePointer Function %ulong
+     %uint_0 = OpConstant %uint 0
 %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
-          %1 = OpFunction %void None %69
+         %52 = OpFunction %ulong None %69
+         %54 = OpFunctionParameter %uint
+               OpFunctionEnd
+          %1 = OpFunction %void None %72
          %28 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
          %29 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
-         %60 = OpLabel
+         %63 = OpLabel
          %20 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
          %21 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
          %14 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
@@ -41,14 +44,14 @@
          %12 = OpVariable %_ptr_Function_ulong Function
                OpStore %20 %28
                OpStore %21 %29
-         %53 = OpBitcast %_ptr_Function_ulong %20
-         %52 = OpLoad %ulong %53 Aligned 8
-         %22 = OpCopyObject %ulong %52
+         %56 = OpBitcast %_ptr_Function_ulong %20
+         %55 = OpLoad %ulong %56 Aligned 8
+         %22 = OpCopyObject %ulong %55
          %30 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %22
                OpStore %14 %30
-         %55 = OpBitcast %_ptr_Function_ulong %21
-         %54 = OpLoad %ulong %55 Aligned 8
-         %23 = OpCopyObject %ulong %54
+         %58 = OpBitcast %_ptr_Function_ulong %21
+         %57 = OpLoad %ulong %58 Aligned 8
+         %23 = OpCopyObject %ulong %57
          %31 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %23
                OpStore %17 %31
          %32 = OpLoad %_ptr_CrossWorkgroup_uchar %14
@@ -61,37 +64,36 @@
          %26 = OpCopyObject %ulong %27
          %35 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %26
                OpStore %18 %35
-         %74 = OpLoad %v3ulong %gl_LocalInvocationID
-         %51 = OpCompositeExtract %ulong %74 0
-         %75 = OpBitcast %ulong %51
-         %37 = OpUConvert %uint %75
+         %51 = OpFunctionCall %ulong %52 %uint_0
+         %76 = OpBitcast %ulong %51
+         %37 = OpUConvert %uint %76
          %36 = OpCopyObject %uint %37
                OpStore %10 %36
          %39 = OpLoad %uint %10
-         %76 = OpBitcast %uint %39
-         %38 = OpUConvert %ulong %76
+         %77 = OpBitcast %uint %39
+         %38 = OpUConvert %ulong %77
                OpStore %11 %38
          %41 = OpLoad %_ptr_CrossWorkgroup_uchar %15
          %42 = OpLoad %ulong %11
-         %56 = OpCopyObject %ulong %42
-         %77 = OpBitcast %_ptr_CrossWorkgroup_uchar %41
-         %78 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %77 %56
-         %40 = OpBitcast %_ptr_CrossWorkgroup_uchar %78
+         %59 = OpCopyObject %ulong %42
+         %78 = OpBitcast %_ptr_CrossWorkgroup_uchar %41
+         %79 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %78 %59
+         %40 = OpBitcast %_ptr_CrossWorkgroup_uchar %79
                OpStore %16 %40
          %44 = OpLoad %_ptr_CrossWorkgroup_uchar %18
          %45 = OpLoad %ulong %11
-         %57 = OpCopyObject %ulong %45
-         %79 = OpBitcast %_ptr_CrossWorkgroup_uchar %44
-         %80 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %79 %57
-         %43 = OpBitcast %_ptr_CrossWorkgroup_uchar %80
+         %60 = OpCopyObject %ulong %45
+         %80 = OpBitcast %_ptr_CrossWorkgroup_uchar %44
+         %81 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %80 %60
+         %43 = OpBitcast %_ptr_CrossWorkgroup_uchar %81
                OpStore %19 %43
          %47 = OpLoad %_ptr_CrossWorkgroup_uchar %16
-         %58 = OpBitcast %_ptr_CrossWorkgroup_ulong %47
-         %46 = OpLoad %ulong %58 Aligned 8
+         %61 = OpBitcast %_ptr_CrossWorkgroup_ulong %47
+         %46 = OpLoad %ulong %61 Aligned 8
                OpStore %12 %46
          %48 = OpLoad %_ptr_CrossWorkgroup_uchar %19
          %49 = OpLoad %ulong %12
-         %59 = OpBitcast %_ptr_CrossWorkgroup_ulong %48
-               OpStore %59 %49 Aligned 8
+         %62 = OpBitcast %_ptr_CrossWorkgroup_ulong %48
+               OpStore %62 %49 Aligned 8
                OpReturn
                OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/stateful_ld_st_ntid_sub.spvtxt b/ptx/src/test/spirv_run/stateful_ld_st_ntid_sub.spvtxt
index 1d0fdfc..987e205 100644
--- a/ptx/src/test/spirv_run/stateful_ld_st_ntid_sub.spvtxt
+++ b/ptx/src/test/spirv_run/stateful_ld_st_ntid_sub.spvtxt
@@ -7,31 +7,34 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %66 = OpExtInstImport "OpenCL.std"
+         %71 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %1 "stateful_ld_st_ntid_sub" %gl_LocalInvocationID
-               OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
+               OpEntryPoint Kernel %1 "stateful_ld_st_ntid_sub"
+               OpExecutionMode %1 ContractionOff
+               OpDecorate %54 LinkageAttributes "_Z12get_local_idj" Import
        %void = OpTypeVoid
       %ulong = OpTypeInt 64 0
-    %v3ulong = OpTypeVector %ulong 3
-%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
-%gl_LocalInvocationID = OpVariable %_ptr_Input_v3ulong Input
+       %uint = OpTypeInt 32 0
+         %75 = OpTypeFunction %ulong %uint
       %uchar = OpTypeInt 8 0
 %_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
-         %73 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
+         %78 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
 %_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
-       %uint = OpTypeInt 32 0
 %_ptr_Function_uint = OpTypePointer Function %uint
 %_ptr_Function_ulong = OpTypePointer Function %ulong
+     %uint_0 = OpConstant %uint 0
     %ulong_0 = OpConstant %ulong 0
 %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
   %ulong_0_0 = OpConstant %ulong 0
-          %1 = OpFunction %void None %73
+         %54 = OpFunction %ulong None %75
+         %56 = OpFunctionParameter %uint
+               OpFunctionEnd
+          %1 = OpFunction %void None %78
          %30 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
          %31 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
-         %64 = OpLabel
-          %2 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
-          %3 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
+         %69 = OpLabel
+         %20 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
+         %21 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
          %14 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
          %15 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
          %16 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
@@ -41,17 +44,17 @@
          %10 = OpVariable %_ptr_Function_uint Function
          %11 = OpVariable %_ptr_Function_ulong Function
          %12 = OpVariable %_ptr_Function_ulong Function
-               OpStore %2 %30
-               OpStore %3 %31
-         %21 = OpBitcast %_ptr_Function_ulong %2
-         %58 = OpLoad %ulong %21 Aligned 8
-         %20 = OpCopyObject %ulong %58
-         %32 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %20
+               OpStore %20 %30
+               OpStore %21 %31
+         %62 = OpBitcast %_ptr_Function_ulong %20
+         %61 = OpLoad %ulong %62 Aligned 8
+         %22 = OpCopyObject %ulong %61
+         %32 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %22
                OpStore %14 %32
-         %23 = OpBitcast %_ptr_Function_ulong %3
-         %59 = OpLoad %ulong %23 Aligned 8
-         %22 = OpCopyObject %ulong %59
-         %33 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %22
+         %64 = OpBitcast %_ptr_Function_ulong %21
+         %63 = OpLoad %ulong %64 Aligned 8
+         %23 = OpCopyObject %ulong %63
+         %33 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %23
                OpStore %17 %33
          %34 = OpLoad %_ptr_CrossWorkgroup_uchar %14
          %25 = OpConvertPtrToU %ulong %34
@@ -63,45 +66,44 @@
          %26 = OpCopyObject %ulong %27
          %37 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %26
                OpStore %18 %37
-         %78 = OpLoad %v3ulong %gl_LocalInvocationID
-         %53 = OpCompositeExtract %ulong %78 0
-         %79 = OpBitcast %ulong %53
-         %39 = OpUConvert %uint %79
+         %53 = OpFunctionCall %ulong %54 %uint_0
+         %82 = OpBitcast %ulong %53
+         %39 = OpUConvert %uint %82
          %38 = OpCopyObject %uint %39
                OpStore %10 %38
          %41 = OpLoad %uint %10
-         %80 = OpBitcast %uint %41
-         %40 = OpUConvert %ulong %80
+         %83 = OpBitcast %uint %41
+         %40 = OpUConvert %ulong %83
                OpStore %11 %40
          %42 = OpLoad %ulong %11
-         %60 = OpCopyObject %ulong %42
-         %28 = OpSNegate %ulong %60
+         %65 = OpCopyObject %ulong %42
+         %28 = OpSNegate %ulong %65
          %44 = OpLoad %_ptr_CrossWorkgroup_uchar %15
-         %81 = OpBitcast %_ptr_CrossWorkgroup_uchar %44
-         %82 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %81 %28
-         %43 = OpBitcast %_ptr_CrossWorkgroup_uchar %82
+         %84 = OpBitcast %_ptr_CrossWorkgroup_uchar %44
+         %85 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %84 %28
+         %43 = OpBitcast %_ptr_CrossWorkgroup_uchar %85
                OpStore %16 %43
          %45 = OpLoad %ulong %11
-         %61 = OpCopyObject %ulong %45
-         %29 = OpSNegate %ulong %61
+         %66 = OpCopyObject %ulong %45
+         %29 = OpSNegate %ulong %66
          %47 = OpLoad %_ptr_CrossWorkgroup_uchar %18
-         %83 = OpBitcast %_ptr_CrossWorkgroup_uchar %47
-         %84 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %83 %29
-         %46 = OpBitcast %_ptr_CrossWorkgroup_uchar %84
+         %86 = OpBitcast %_ptr_CrossWorkgroup_uchar %47
+         %87 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %86 %29
+         %46 = OpBitcast %_ptr_CrossWorkgroup_uchar %87
                OpStore %19 %46
          %49 = OpLoad %_ptr_CrossWorkgroup_uchar %16
-         %62 = OpBitcast %_ptr_CrossWorkgroup_ulong %49
-         %86 = OpBitcast %_ptr_CrossWorkgroup_uchar %62
-         %87 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %86 %ulong_0
-         %55 = OpBitcast %_ptr_CrossWorkgroup_ulong %87
-         %48 = OpLoad %ulong %55 Aligned 8
+         %67 = OpBitcast %_ptr_CrossWorkgroup_ulong %49
+         %89 = OpBitcast %_ptr_CrossWorkgroup_uchar %67
+         %90 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %89 %ulong_0
+         %58 = OpBitcast %_ptr_CrossWorkgroup_ulong %90
+         %48 = OpLoad %ulong %58 Aligned 8
                OpStore %12 %48
          %50 = OpLoad %_ptr_CrossWorkgroup_uchar %19
          %51 = OpLoad %ulong %12
-         %63 = OpBitcast %_ptr_CrossWorkgroup_ulong %50
-         %88 = OpBitcast %_ptr_CrossWorkgroup_uchar %63
-         %89 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %88 %ulong_0_0
-         %57 = OpBitcast %_ptr_CrossWorkgroup_ulong %89
-               OpStore %57 %51 Aligned 8
+         %68 = OpBitcast %_ptr_CrossWorkgroup_ulong %50
+         %91 = OpBitcast %_ptr_CrossWorkgroup_uchar %68
+         %92 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %91 %ulong_0_0
+         %60 = OpBitcast %_ptr_CrossWorkgroup_ulong %92
+               OpStore %60 %51 Aligned 8
                OpReturn
                OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/vector4.spvtxt b/ptx/src/test/spirv_run/vector4.spvtxt
index 8253bf9..9b6349b 100644
--- a/ptx/src/test/spirv_run/vector4.spvtxt
+++ b/ptx/src/test/spirv_run/vector4.spvtxt
@@ -7,93 +7,50 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %51 = OpExtInstImport "OpenCL.std"
+         %24 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %25 "vector"
+               OpEntryPoint Kernel %1 "vector4"
+               OpExecutionMode %1 ContractionOff
        %void = OpTypeVoid
-       %uint = OpTypeInt 32 0
-     %v2uint = OpTypeVector %uint 2
-         %55 = OpTypeFunction %v2uint %v2uint
-%_ptr_Function_v2uint = OpTypePointer Function %v2uint
-%_ptr_Function_uint = OpTypePointer Function %uint
-     %uint_0 = OpConstant %uint 0
-     %uint_1 = OpConstant %uint 1
       %ulong = OpTypeInt 64 0
-         %67 = OpTypeFunction %void %ulong %ulong
+         %27 = OpTypeFunction %void %ulong %ulong
 %_ptr_Function_ulong = OpTypePointer Function %ulong
-%_ptr_Generic_v2uint = OpTypePointer Generic %v2uint
-          %1 = OpFunction %v2uint None %55
-          %7 = OpFunctionParameter %v2uint
-         %24 = OpLabel
-          %3 = OpVariable %_ptr_Function_v2uint Function
-          %2 = OpVariable %_ptr_Function_v2uint Function
-          %4 = OpVariable %_ptr_Function_v2uint Function
-          %5 = OpVariable %_ptr_Function_uint Function
-          %6 = OpVariable %_ptr_Function_uint Function
-               OpStore %3 %7
-         %59 = OpInBoundsAccessChain %_ptr_Function_uint %3 %uint_0
-          %9 = OpLoad %uint %59
-          %8 = OpCopyObject %uint %9
-               OpStore %5 %8
-         %61 = OpInBoundsAccessChain %_ptr_Function_uint %3 %uint_1
-         %11 = OpLoad %uint %61
-         %10 = OpCopyObject %uint %11
-               OpStore %6 %10
-         %13 = OpLoad %uint %5
-         %14 = OpLoad %uint %6
-         %12 = OpIAdd %uint %13 %14
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+%_ptr_Generic_v4uint = OpTypePointer Generic %v4uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_Generic_uint = OpTypePointer Generic %uint
+          %1 = OpFunction %void None %27
+          %8 = OpFunctionParameter %ulong
+          %9 = OpFunctionParameter %ulong
+         %22 = OpLabel
+          %2 = OpVariable %_ptr_Function_ulong Function
+          %3 = OpVariable %_ptr_Function_ulong Function
+          %4 = OpVariable %_ptr_Function_ulong Function
+          %5 = OpVariable %_ptr_Function_ulong Function
+          %6 = OpVariable %_ptr_Function_v4uint Function
+          %7 = OpVariable %_ptr_Function_uint Function
+               OpStore %2 %8
+               OpStore %3 %9
+         %10 = OpLoad %ulong %2 Aligned 8
+               OpStore %4 %10
+         %11 = OpLoad %ulong %3 Aligned 8
+               OpStore %5 %11
+         %13 = OpLoad %ulong %4
+         %18 = OpConvertUToPtr %_ptr_Generic_v4uint %13
+         %12 = OpLoad %v4uint %18 Aligned 16
                OpStore %6 %12
-         %16 = OpLoad %uint %6
-         %15 = OpCopyObject %uint %16
-         %62 = OpInBoundsAccessChain %_ptr_Function_uint %4 %uint_0
-               OpStore %62 %15
-         %18 = OpLoad %uint %6
-         %17 = OpCopyObject %uint %18
-         %63 = OpInBoundsAccessChain %_ptr_Function_uint %4 %uint_1
-               OpStore %63 %17
-         %64 = OpInBoundsAccessChain %_ptr_Function_uint %4 %uint_1
-         %20 = OpLoad %uint %64
+         %35 = OpInBoundsAccessChain %_ptr_Function_uint %6 %uint_3
+         %15 = OpLoad %uint %35
+         %20 = OpCopyObject %uint %15
          %19 = OpCopyObject %uint %20
-         %65 = OpInBoundsAccessChain %_ptr_Function_uint %4 %uint_0
-               OpStore %65 %19
-         %22 = OpLoad %v2uint %4
-         %21 = OpCopyObject %v2uint %22
-               OpStore %2 %21
-         %23 = OpLoad %v2uint %2
-               OpReturnValue %23
-               OpFunctionEnd
-         %25 = OpFunction %void None %67
-         %34 = OpFunctionParameter %ulong
-         %35 = OpFunctionParameter %ulong
-         %49 = OpLabel
-         %26 = OpVariable %_ptr_Function_ulong Function
-         %27 = OpVariable %_ptr_Function_ulong Function
-         %28 = OpVariable %_ptr_Function_ulong Function
-         %29 = OpVariable %_ptr_Function_ulong Function
-         %30 = OpVariable %_ptr_Function_v2uint Function
-         %31 = OpVariable %_ptr_Function_uint Function
-         %32 = OpVariable %_ptr_Function_uint Function
-         %33 = OpVariable %_ptr_Function_ulong Function
-               OpStore %26 %34
-               OpStore %27 %35
-         %36 = OpLoad %ulong %26 Aligned 8
-               OpStore %28 %36
-         %37 = OpLoad %ulong %27 Aligned 8
-               OpStore %29 %37
-         %39 = OpLoad %ulong %28
-         %46 = OpConvertUToPtr %_ptr_Generic_v2uint %39
-         %38 = OpLoad %v2uint %46 Aligned 8
-               OpStore %30 %38
-         %41 = OpLoad %v2uint %30
-         %40 = OpFunctionCall %v2uint %1 %41
-               OpStore %30 %40
-         %43 = OpLoad %v2uint %30
-         %47 = OpBitcast %ulong %43
-         %42 = OpCopyObject %ulong %47
-               OpStore %33 %42
-         %44 = OpLoad %ulong %29
-         %45 = OpLoad %v2uint %30
-         %48 = OpConvertUToPtr %_ptr_Generic_v2uint %44
-               OpStore %48 %45 Aligned 8
+         %14 = OpCopyObject %uint %19
+               OpStore %7 %14
+         %16 = OpLoad %ulong %5
+         %17 = OpLoad %uint %7
+         %21 = OpConvertUToPtr %_ptr_Generic_uint %16
+               OpStore %21 %17 Aligned 4
                OpReturn
                OpFunctionEnd
diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs
index 524196a..14e8473 100644
--- a/ptx/src/translate.rs
+++ b/ptx/src/translate.rs
@@ -431,7 +431,7 @@ pub fn to_spirv_module<'a>(ast: ast::Module<'a>) -> Result<Module, TranslateErro
         })

         .collect::<Result<Vec<_>, _>>()?;

     let must_link_ptx_impl = ptx_impl_imports.len() > 0;

-    let directives = ptx_impl_imports

+    let mut directives = ptx_impl_imports

         .into_iter()

         .map(|(_, v)| v)

         .chain(directives.into_iter())

@@ -439,7 +439,7 @@ pub fn to_spirv_module<'a>(ast: ast::Module<'a>) -> Result<Module, TranslateErro
     let mut builder = dr::Builder::new();

     builder.reserve_ids(id_defs.current_id());

     let call_map = get_kernels_call_map(&directives);

-    let mut directives = convert_dynamic_shared_memory_usage(directives, &mut || builder.id());

+    //let mut directives = convert_dynamic_shared_memory_usage(directives, &mut || builder.id());

     normalize_variable_decls(&mut directives);

     let denorm_information = compute_denorm_information(&directives);

     // https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#_a_id_logicallayout_a_logical_layout_of_a_module

@@ -532,8 +532,8 @@ fn emit_directives<'input>(
     let empty_body = Vec::new();

     for d in directives.iter() {

         match d {

-            Directive::Variable(_, var) => {

-                emit_variable(builder, map, &var)?;

+            Directive::Variable(linking, var) => {

+                emit_variable(builder, map, id_defs, *linking, &var)?;

             }

             Directive::Method(f) => {

                 let f_body = match &f.body {

@@ -547,7 +547,7 @@ fn emit_directives<'input>(
                     }

                 };

                 for var in f.globals.iter() {

-                    emit_variable(builder, map, var)?;

+                    emit_variable(builder, map, id_defs, ast::LinkingDirective::NONE, var)?;

                 }

                 let func_decl = (*f.func_decl).borrow();

                 let fn_id = emit_function_header(

@@ -602,7 +602,7 @@ fn emit_directives<'input>(
                         }

                     }

                 }

-                emit_function_body_ops(builder, map, opencl_id, &f_body)?;

+                emit_function_body_ops(builder, map, id_defs, opencl_id, &f_body)?;

                 builder.end_function()?;

                 if let (

                     ast::MethodDeclaration {

@@ -2497,9 +2497,10 @@ fn get_function_type(
     )

 }

 

-fn emit_function_body_ops(

+fn emit_function_body_ops<'input>(

     builder: &mut dr::Builder,

     map: &mut TypeWordMap,

+    id_defs: &GlobalStringIdResolver<'input>,

     opencl: spirv::Word,

     func: &[ExpandedStatement],

 ) -> Result<(), TranslateError> {

@@ -2541,7 +2542,7 @@ fn emit_function_body_ops(
                 builder.function_call(result_type, result_id, call.name, arg_list)?;

             }

             Statement::Variable(var) => {

-                emit_variable(builder, map, var)?;

+                emit_variable(builder, map, id_defs, ast::LinkingDirective::NONE, var)?;

             }

             Statement::Constant(cnst) => {

                 let typ_id = map.get_or_add_scalar(builder, cnst.typ);

@@ -3287,9 +3288,11 @@ fn vec_repr<T: Copy>(t: T) -> Vec<u8> {
     result

 }

 

-fn emit_variable(

+fn emit_variable<'input>(

     builder: &mut dr::Builder,

     map: &mut TypeWordMap,

+    id_defs: &GlobalStringIdResolver<'input>,

+    linking: ast::LinkingDirective,

     var: &ast::Variable<spirv::Word>,

 ) -> Result<(), TranslateError> {

     let (must_init, st_class) = match var.state_space {

@@ -3323,9 +3326,45 @@ fn emit_variable(
             [dr::Operand::LiteralInt32(align)].iter().cloned(),

         );

     }

+    emit_linking_decoration(builder, id_defs, var.name, linking);

     Ok(())

 }

 

+fn emit_linking_decoration<'input>(

+    builder: &mut dr::Builder,

+    id_defs: &GlobalStringIdResolver<'input>,

+    name: spirv::Word,

+    linking: ast::LinkingDirective,

+) {

+    if linking.contains(ast::LinkingDirective::EXTERN) {

+        let external_name = id_defs.reverse_variables.get(&name).unwrap();

+        builder.decorate(

+            name,

+            spirv::Decoration::LinkageAttributes,

+            [

+                dr::Operand::LiteralString(external_name.to_string()),

+                dr::Operand::LinkageType(spirv::LinkageType::Import),

+            ]

+            .iter()

+            .cloned(),

+        );

+    }

+    if linking.contains(ast::LinkingDirective::VISIBLE) {

+        let external_name = id_defs.reverse_variables.get(&name).unwrap();

+        builder.decorate(

+            name,

+            spirv::Decoration::LinkageAttributes,

+            [

+                dr::Operand::LiteralString(external_name.to_string()),

+                dr::Operand::LinkageType(spirv::LinkageType::Export),

+            ]

+            .iter()

+            .cloned(),

+        );

+    }

+    // TODO: handle LinkingDirective::WEAK

+}

+

 fn emit_mad_uint(

     builder: &mut dr::Builder,

     map: &mut TypeWordMap,

@@ -4902,6 +4941,7 @@ impl<'input> FnSigMapper<'input> {
 struct GlobalStringIdResolver<'input> {

     current_id: spirv::Word,

     variables: HashMap<Cow<'input, str>, spirv::Word>,

+    reverse_variables: HashMap<spirv::Word, &'input str>,

     variables_type_check: HashMap<u32, Option<(ast::Type, ast::StateSpace, bool)>>,

     special_registers: SpecialRegistersMap,

     fns: HashMap<spirv::Word, FnSigMapper<'input>>,

@@ -4912,6 +4952,7 @@ impl<'input> GlobalStringIdResolver<'input> {
         Self {

             current_id: start_id,

             variables: HashMap::new(),

+            reverse_variables: HashMap::new(),

             variables_type_check: HashMap::new(),

             special_registers: SpecialRegistersMap::new(),

             fns: HashMap::new(),

@@ -4942,6 +4983,7 @@ impl<'input> GlobalStringIdResolver<'input> {
             hash_map::Entry::Vacant(e) => {

                 let numeric_id = self.current_id;

                 e.insert(numeric_id);

+                self.reverse_variables.insert(numeric_id, id);

                 self.current_id += 1;

                 numeric_id

             }