Applied-Software/ZLUDA
5
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-07-17 09:16:23 +03:00
Code Issues Packages Projects Releases Wiki Activity

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

Browse Source
This commit is contained in:
Andrzej Janik
2021-09-09 00:17:39 +00:00
parent a27d1e119f
commit da9cf4d583
16 changed files with 270 additions and 252 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
ptx/src/test/spirv_run/cos.spvtxt
View File

@ -10,6 +10,7 @@
%21 = OpExtInstImport "OpenCL.std" %21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "cos" OpEntryPoint Kernel %1 "cos"
OpExecutionMode %1 ContractionOff
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%24 = OpTypeFunction %void %ulong %ulong %24 = OpTypeFunction %void %ulong %ulong
@ -37,7 +38,7 @@
%11 = OpLoad %float %17 Aligned 4 %11 = OpLoad %float %17 Aligned 4
OpStore %6 %11 OpStore %6 %11
%14 = OpLoad %float %6 %14 = OpLoad %float %6
%13 = OpExtInst %float %21 native_cos %14 %13 = OpExtInst %float %21 cos %14
OpStore %6 %13 OpStore %6 %13
%15 = OpLoad %ulong %5 %15 = OpLoad %ulong %5
%16 = OpLoad %float %6 %16 = OpLoad %float %6

3
ptx/src/test/spirv_run/ex2.spvtxt
View File

@ -10,6 +10,7 @@
%21 = OpExtInstImport "OpenCL.std" %21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "ex2" OpEntryPoint Kernel %1 "ex2"
OpExecutionMode %1 ContractionOff
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%24 = OpTypeFunction %void %ulong %ulong %24 = OpTypeFunction %void %ulong %ulong
@ -37,7 +38,7 @@
%11 = OpLoad %float %17 Aligned 4 %11 = OpLoad %float %17 Aligned 4
OpStore %6 %11 OpStore %6 %11
%14 = OpLoad %float %6 %14 = OpLoad %float %6
%13 = OpExtInst %float %21 native_exp2 %14 %13 = OpExtInst %float %21 exp2 %14
OpStore %6 %13 OpStore %6 %13
%15 = OpLoad %ulong %5 %15 = OpLoad %ulong %5
%16 = OpLoad %float %6 %16 = OpLoad %float %6

17
ptx/src/test/spirv_run/extern_shared.spvtxt
View File

@ -7,24 +7,23 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%27 = OpExtInstImport "OpenCL.std" %24 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "extern_shared" %1 OpEntryPoint Kernel %2 "extern_shared" %1
OpExecutionMode %2 ContractionOff
OpDecorate %1 LinkageAttributes "shared_mem" Import
%void = OpTypeVoid %void = OpTypeVoid
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
%1 = OpVariable %_ptr_Workgroup_uint Workgroup %1 = OpVariable %_ptr_Workgroup_uint Workgroup
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%uchar = OpTypeInt 8 0 %29 = OpTypeFunction %void %ulong %ulong
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
%34 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%2 = OpFunction %void None %34 %2 = OpFunction %void None %29
%8 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong
%9 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong
%24 = OpFunctionParameter %_ptr_Workgroup_uchar
%22 = OpLabel %22 = OpLabel
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -42,11 +41,9 @@
%12 = OpLoad %ulong %18 Aligned 8 %12 = OpLoad %ulong %18 Aligned 8
OpStore %7 %12 OpStore %7 %12
%14 = OpLoad %ulong %7 %14 = OpLoad %ulong %7
%25 = OpBitcast %_ptr_Workgroup_uint %24 %19 = OpBitcast %_ptr_Workgroup_ulong %1
%19 = OpBitcast %_ptr_Workgroup_ulong %25
OpStore %19 %14 Aligned 8 OpStore %19 %14 Aligned 8
%26 = OpBitcast %_ptr_Workgroup_uint %24 %20 = OpBitcast %_ptr_Workgroup_ulong %1
%20 = OpBitcast %_ptr_Workgroup_ulong %26
%15 = OpLoad %ulong %20 Aligned 8 %15 = OpLoad %ulong %20 Aligned 8
OpStore %7 %15 OpStore %7 %15
%16 = OpLoad %ulong %6 %16 = OpLoad %ulong %6

30
ptx/src/test/spirv_run/extern_shared_call.spvtxt
View File

@ -7,44 +7,40 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%40 = OpExtInstImport "OpenCL.std" %34 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %12 "extern_shared_call" %1 OpEntryPoint Kernel %12 "extern_shared_call" %1
OpExecutionMode %12 ContractionOff
OpDecorate %1 Alignment 4 OpDecorate %1 Alignment 4
OpDecorate %1 LinkageAttributes "shared_mem" Import
%void = OpTypeVoid %void = OpTypeVoid
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
%1 = OpVariable %_ptr_Workgroup_uint Workgroup %1 = OpVariable %_ptr_Workgroup_uint Workgroup
%uchar = OpTypeInt 8 0 %38 = OpTypeFunction %void
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
%46 = OpTypeFunction %void %_ptr_Workgroup_uchar
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%ulong_2 = OpConstant %ulong 2 %ulong_2 = OpConstant %ulong 2
%50 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar %42 = OpTypeFunction %void %ulong %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%2 = OpFunction %void None %46 %2 = OpFunction %void None %38
%34 = OpFunctionParameter %_ptr_Workgroup_uchar
%11 = OpLabel %11 = OpLabel
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%35 = OpBitcast %_ptr_Workgroup_uint %34 %9 = OpBitcast %_ptr_Workgroup_ulong %1
%9 = OpBitcast %_ptr_Workgroup_ulong %35
%4 = OpLoad %ulong %9 Aligned 8 %4 = OpLoad %ulong %9 Aligned 8
OpStore %3 %4 OpStore %3 %4
%6 = OpLoad %ulong %3 %6 = OpLoad %ulong %3
%5 = OpIAdd %ulong %6 %ulong_2 %5 = OpIAdd %ulong %6 %ulong_2
OpStore %3 %5 OpStore %3 %5
%7 = OpLoad %ulong %3 %7 = OpLoad %ulong %3
%36 = OpBitcast %_ptr_Workgroup_uint %34 %10 = OpBitcast %_ptr_Workgroup_ulong %1
%10 = OpBitcast %_ptr_Workgroup_ulong %36
OpStore %10 %7 Aligned 8 OpStore %10 %7 Aligned 8
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%12 = OpFunction %void None %50 %12 = OpFunction %void None %42
%18 = OpFunctionParameter %ulong %18 = OpFunctionParameter %ulong
%19 = OpFunctionParameter %ulong %19 = OpFunctionParameter %ulong
%37 = OpFunctionParameter %_ptr_Workgroup_uchar
%32 = OpLabel %32 = OpLabel
%13 = OpVariable %_ptr_Function_ulong Function %13 = OpVariable %_ptr_Function_ulong Function
%14 = OpVariable %_ptr_Function_ulong Function %14 = OpVariable %_ptr_Function_ulong Function
@ -62,12 +58,10 @@
%22 = OpLoad %ulong %28 Aligned 8 %22 = OpLoad %ulong %28 Aligned 8
OpStore %17 %22 OpStore %17 %22
%24 = OpLoad %ulong %17 %24 = OpLoad %ulong %17
%38 = OpBitcast %_ptr_Workgroup_uint %37 %29 = OpBitcast %_ptr_Workgroup_ulong %1
%29 = OpBitcast %_ptr_Workgroup_ulong %38
OpStore %29 %24 Aligned 8 OpStore %29 %24 Aligned 8
%52 = OpFunctionCall %void %2 %37 %44 = OpFunctionCall %void %2
%39 = OpBitcast %_ptr_Workgroup_uint %37 %30 = OpBitcast %_ptr_Workgroup_ulong %1
%30 = OpBitcast %_ptr_Workgroup_ulong %39
%25 = OpLoad %ulong %30 Aligned 8 %25 = OpLoad %ulong %30 Aligned 8
OpStore %17 %25 OpStore %17 %25
%26 = OpLoad %ulong %16 %26 = OpLoad %ulong %16

3
ptx/src/test/spirv_run/lg2.spvtxt
View File

@ -10,6 +10,7 @@
%21 = OpExtInstImport "OpenCL.std" %21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "lg2" OpEntryPoint Kernel %1 "lg2"
OpExecutionMode %1 ContractionOff
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%24 = OpTypeFunction %void %ulong %ulong %24 = OpTypeFunction %void %ulong %ulong
@ -37,7 +38,7 @@
%11 = OpLoad %float %17 Aligned 4 %11 = OpLoad %float %17 Aligned 4
OpStore %6 %11 OpStore %6 %11
%14 = OpLoad %float %6 %14 = OpLoad %float %6
%13 = OpExtInst %float %21 native_log2 %14 %13 = OpExtInst %float %21 log2 %14
OpStore %6 %13 OpStore %6 %13
%15 = OpLoad %ulong %5 %15 = OpLoad %ulong %5
%16 = OpLoad %float %6 %16 = OpLoad %float %6

22
ptx/src/test/spirv_run/mod.rs
View File

@ -32,6 +32,7 @@ use std::io;
use std::io::Read; use std::io::Read;
use std::io::Write; use std::io::Write;
use std::mem; use std::mem;
use std::path::Path;
use std::process::Command; use std::process::Command;
use std::slice; use std::slice;
use std::{borrow::Cow, collections::HashMap, env, fs, path::PathBuf, ptr, str}; 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! { 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>()) }; hip_call! { hipMemset(out_b, 0, output.len() * mem::size_of::<Output>()) };
let mut args = [&inp_b, &out_b]; 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! { hipMemcpyAsync(result.as_mut_ptr() as _, out_b, output.len() * mem::size_of::<Output>(), hipMemcpyKind::hipMemcpyDeviceToHost, stream) };
hip_call! { hipStreamSynchronize(stream) }; hip_call! { hipStreamSynchronize(stream) };
} }
@ -600,6 +601,9 @@ fn compile_amd(
.arg(spirv.path()) .arg(spirv.path())
.status()?; .status()?;
assert!(to_llvm_cmd.success()); assert!(to_llvm_cmd.success());
if cfg!(debug_assertions) {
persist_file(llvm.path())?;
}
let linked_binary = NamedTempFile::new_in(&dir)?; let linked_binary = NamedTempFile::new_in(&dir)?;
let mut llvm_link = PathBuf::from(AMDGPU); let mut llvm_link = PathBuf::from(AMDGPU);
llvm_link.push("llvm"); llvm_link.push("llvm");
@ -617,6 +621,9 @@ fn compile_amd(
} }
let status = linker_cmd.status()?; let status = linker_cmd.status()?;
assert!(status.success()); assert!(status.success());
if cfg!(debug_assertions) {
persist_file(linked_binary.path())?;
}
let mut ptx_lib_bitcode = NamedTempFile::new_in(&dir)?; let mut ptx_lib_bitcode = NamedTempFile::new_in(&dir)?;
let compiled_binary = NamedTempFile::new_in(&dir)?; let compiled_binary = NamedTempFile::new_in(&dir)?;
let mut clang_exe = PathBuf::from(AMDGPU); let mut clang_exe = PathBuf::from(AMDGPU);
@ -651,11 +658,18 @@ fn compile_amd(
let compiled_bin_path = compiled_binary.path(); let compiled_bin_path = compiled_binary.path();
let mut compiled_binary = File::open(compiled_bin_path)?; let mut compiled_binary = File::open(compiled_bin_path)?;
compiled_binary.read_to_end(&mut result)?; 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"); let mut persistent = PathBuf::from("/tmp/zluda");
std::fs::create_dir_all(&persistent)?; std::fs::create_dir_all(&persistent)?;
persistent.push(compiled_bin_path.file_name().unwrap()); persistent.push(path.file_name().unwrap());
std::fs::copy(compiled_bin_path, persistent)?; std::fs::copy(path, persistent)?;
Ok(result) Ok(())
} }
fn get_bitcode_paths(device_name: &str) -> impl Iterator<Item = PathBuf> { fn get_bitcode_paths(device_name: &str) -> impl Iterator<Item = PathBuf> {

5
ptx/src/test/spirv_run/reg_local.spvtxt
View File

@ -10,6 +10,7 @@
%34 = OpExtInstImport "OpenCL.std" %34 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "reg_local" OpEntryPoint Kernel %1 "reg_local"
OpExecutionMode %1 ContractionOff
OpDecorate %4 Alignment 8 OpDecorate %4 Alignment 8
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
@ -50,10 +51,10 @@
OpStore %7 %12 OpStore %7 %12
%14 = OpLoad %ulong %7 %14 = OpLoad %ulong %7
%19 = OpIAdd %ulong %14 %ulong_1 %19 = OpIAdd %ulong %14 %ulong_1
%26 = OpBitcast %_ptr_Generic_ulong %4 %26 = OpPtrCastToGeneric %_ptr_Generic_ulong %4
%27 = OpCopyObject %ulong %19 %27 = OpCopyObject %ulong %19
OpStore %26 %27 Aligned 8 OpStore %26 %27 Aligned 8
%28 = OpBitcast %_ptr_Generic_ulong %4 %28 = OpPtrCastToGeneric %_ptr_Generic_ulong %4
%47 = OpBitcast %_ptr_Generic_uchar %28 %47 = OpBitcast %_ptr_Generic_uchar %28
%48 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %47 %ulong_0 %48 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %47 %ulong_0
%21 = OpBitcast %_ptr_Generic_ulong %48 %21 = OpBitcast %_ptr_Generic_ulong %48

3
ptx/src/test/spirv_run/rsqrt.spvtxt
View File

@ -10,6 +10,7 @@
%21 = OpExtInstImport "OpenCL.std" %21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "rsqrt" OpEntryPoint Kernel %1 "rsqrt"
OpExecutionMode %1 ContractionOff
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%24 = OpTypeFunction %void %ulong %ulong %24 = OpTypeFunction %void %ulong %ulong
@ -37,7 +38,7 @@
%11 = OpLoad %double %17 Aligned 8 %11 = OpLoad %double %17 Aligned 8
OpStore %6 %11 OpStore %6 %11
%14 = OpLoad %double %6 %14 = OpLoad %double %6
%13 = OpExtInst %double %21 native_rsqrt %14 %13 = OpExtInst %double %21 rsqrt %14
OpStore %6 %13 OpStore %6 %13
%15 = OpLoad %ulong %5 %15 = OpLoad %ulong %5
%16 = OpLoad %double %6 %16 = OpLoad %double %6

11
ptx/src/test/spirv_run/shared_ptr_take_address.spvtxt
View File

@ -7,23 +7,24 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%31 = OpExtInstImport "OpenCL.std" %30 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "shared_ptr_take_address" %1 OpEntryPoint Kernel %2 "shared_ptr_take_address" %1
OpExecutionMode %2 ContractionOff
OpDecorate %1 Alignment 4 OpDecorate %1 Alignment 4
OpDecorate %1 LinkageAttributes "shared_mem" Import
%void = OpTypeVoid %void = OpTypeVoid
%uchar = OpTypeInt 8 0 %uchar = OpTypeInt 8 0
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar %_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
%1 = OpVariable %_ptr_Workgroup_uchar Workgroup %1 = OpVariable %_ptr_Workgroup_uchar Workgroup
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%36 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar %35 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
%2 = OpFunction %void None %36 %2 = OpFunction %void None %35
%10 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong
%11 = OpFunctionParameter %ulong %11 = OpFunctionParameter %ulong
%30 = OpFunctionParameter %_ptr_Workgroup_uchar
%28 = OpLabel %28 = OpLabel
%3 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function
%4 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function
@ -38,7 +39,7 @@
OpStore %5 %12 OpStore %5 %12
%13 = OpLoad %ulong %4 Aligned 8 %13 = OpLoad %ulong %4 Aligned 8
OpStore %6 %13 OpStore %6 %13
%23 = OpConvertPtrToU %ulong %30 %23 = OpConvertPtrToU %ulong %1
%14 = OpCopyObject %ulong %23 %14 = OpCopyObject %ulong %23
OpStore %7 %14 OpStore %7 %14
%16 = OpLoad %ulong %5 %16 = OpLoad %ulong %5

3
ptx/src/test/spirv_run/sin.spvtxt
View File

@ -10,6 +10,7 @@
%21 = OpExtInstImport "OpenCL.std" %21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "sin" OpEntryPoint Kernel %1 "sin"
OpExecutionMode %1 ContractionOff
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%24 = OpTypeFunction %void %ulong %ulong %24 = OpTypeFunction %void %ulong %ulong
@ -37,7 +38,7 @@
%11 = OpLoad %float %17 Aligned 4 %11 = OpLoad %float %17 Aligned 4
OpStore %6 %11 OpStore %6 %11
%14 = OpLoad %float %6 %14 = OpLoad %float %6
%13 = OpExtInst %float %21 native_sin %14 %13 = OpExtInst %float %21 sin %14
OpStore %6 %13 OpStore %6 %13
%15 = OpLoad %ulong %5 %15 = OpLoad %ulong %5
%16 = OpLoad %float %6 %16 = OpLoad %float %6

3
ptx/src/test/spirv_run/sqrt.spvtxt
View File

@ -10,6 +10,7 @@
%21 = OpExtInstImport "OpenCL.std" %21 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "sqrt" OpEntryPoint Kernel %1 "sqrt"
OpExecutionMode %1 ContractionOff
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%24 = OpTypeFunction %void %ulong %ulong %24 = OpTypeFunction %void %ulong %ulong
@ -37,7 +38,7 @@
%11 = OpLoad %float %17 Aligned 4 %11 = OpLoad %float %17 Aligned 4
OpStore %6 %11 OpStore %6 %11
%14 = OpLoad %float %6 %14 = OpLoad %float %6
%13 = OpExtInst %float %21 native_sqrt %14 %13 = OpExtInst %float %21 sqrt %14
OpStore %6 %13 OpStore %6 %13
%15 = OpLoad %ulong %5 %15 = OpLoad %ulong %5
%16 = OpLoad %float %6 %16 = OpLoad %float %6

70
ptx/src/test/spirv_run/stateful_ld_st_ntid.spvtxt
View File

@ -7,27 +7,30 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%54 = OpExtInstImport "OpenCL.std" %57 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "stateful_ld_st_ntid" %gl_LocalInvocationID OpEntryPoint Kernel %1 "stateful_ld_st_ntid"
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId OpExecutionMode %1 ContractionOff
OpDecorate %44 LinkageAttributes "_Z12get_local_idj" Import
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%v3ulong = OpTypeVector %ulong 3 %uint = OpTypeInt 32 0
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong %61 = OpTypeFunction %ulong %uint
%gl_LocalInvocationID = OpVariable %_ptr_Input_v3ulong Input
%uchar = OpTypeInt 8 0 %uchar = OpTypeInt 8 0
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar %_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 %_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
%uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint_0 = OpConstant %uint 0
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_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 %20 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%21 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %21 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%52 = OpLabel %55 = OpLabel
%12 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %12 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%13 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %13 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%10 = 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 %8 = OpVariable %_ptr_Function_ulong Function
OpStore %12 %20 OpStore %12 %20
OpStore %13 %21 OpStore %13 %21
%45 = OpBitcast %_ptr_Function_ulong %12 %48 = OpBitcast %_ptr_Function_ulong %12
%44 = OpLoad %ulong %45 Aligned 8 %47 = OpLoad %ulong %48 Aligned 8
%14 = OpCopyObject %ulong %44 %14 = OpCopyObject %ulong %47
%22 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %14 %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %14
OpStore %10 %22 OpStore %10 %22
%47 = OpBitcast %_ptr_Function_ulong %13 %50 = OpBitcast %_ptr_Function_ulong %13
%46 = OpLoad %ulong %47 Aligned 8 %49 = OpLoad %ulong %50 Aligned 8
%15 = OpCopyObject %ulong %46 %15 = OpCopyObject %ulong %49
%23 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %15 %23 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %15
OpStore %11 %23 OpStore %11 %23
%24 = OpLoad %_ptr_CrossWorkgroup_uchar %10 %24 = OpLoad %_ptr_CrossWorkgroup_uchar %10
@ -57,37 +60,36 @@
%18 = OpCopyObject %ulong %19 %18 = OpCopyObject %ulong %19
%27 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %18 %27 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %18
OpStore %11 %27 OpStore %11 %27
%66 = OpLoad %v3ulong %gl_LocalInvocationID %43 = OpFunctionCall %ulong %44 %uint_0
%43 = OpCompositeExtract %ulong %66 0 %68 = OpBitcast %ulong %43
%67 = OpBitcast %ulong %43 %29 = OpUConvert %uint %68
%29 = OpUConvert %uint %67
%28 = OpCopyObject %uint %29 %28 = OpCopyObject %uint %29
OpStore %6 %28 OpStore %6 %28
%31 = OpLoad %uint %6 %31 = OpLoad %uint %6
%68 = OpBitcast %uint %31 %69 = OpBitcast %uint %31
%30 = OpUConvert %ulong %68 %30 = OpUConvert %ulong %69
OpStore %7 %30 OpStore %7 %30
%33 = OpLoad %_ptr_CrossWorkgroup_uchar %10 %33 = OpLoad %_ptr_CrossWorkgroup_uchar %10
%34 = OpLoad %ulong %7 %34 = OpLoad %ulong %7
%48 = OpCopyObject %ulong %34 %51 = OpCopyObject %ulong %34
%69 = OpBitcast %_ptr_CrossWorkgroup_uchar %33 %70 = OpBitcast %_ptr_CrossWorkgroup_uchar %33
%70 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %69 %48 %71 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %70 %51
%32 = OpBitcast %_ptr_CrossWorkgroup_uchar %70 %32 = OpBitcast %_ptr_CrossWorkgroup_uchar %71
OpStore %10 %32 OpStore %10 %32
%36 = OpLoad %_ptr_CrossWorkgroup_uchar %11 %36 = OpLoad %_ptr_CrossWorkgroup_uchar %11
%37 = OpLoad %ulong %7 %37 = OpLoad %ulong %7
%49 = OpCopyObject %ulong %37 %52 = OpCopyObject %ulong %37
%71 = OpBitcast %_ptr_CrossWorkgroup_uchar %36 %72 = OpBitcast %_ptr_CrossWorkgroup_uchar %36
%72 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %71 %49 %73 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %72 %52
%35 = OpBitcast %_ptr_CrossWorkgroup_uchar %72 %35 = OpBitcast %_ptr_CrossWorkgroup_uchar %73
OpStore %11 %35 OpStore %11 %35
%39 = OpLoad %_ptr_CrossWorkgroup_uchar %10 %39 = OpLoad %_ptr_CrossWorkgroup_uchar %10
%50 = OpBitcast %_ptr_CrossWorkgroup_ulong %39 %53 = OpBitcast %_ptr_CrossWorkgroup_ulong %39
%38 = OpLoad %ulong %50 Aligned 8 %38 = OpLoad %ulong %53 Aligned 8
OpStore %8 %38 OpStore %8 %38
%40 = OpLoad %_ptr_CrossWorkgroup_uchar %11 %40 = OpLoad %_ptr_CrossWorkgroup_uchar %11
%41 = OpLoad %ulong %8 %41 = OpLoad %ulong %8
%51 = OpBitcast %_ptr_CrossWorkgroup_ulong %40 %54 = OpBitcast %_ptr_CrossWorkgroup_ulong %40
OpStore %51 %41 Aligned 8 OpStore %54 %41 Aligned 8
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

70
ptx/src/test/spirv_run/stateful_ld_st_ntid_chain.spvtxt
View File

@ -7,27 +7,30 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%62 = OpExtInstImport "OpenCL.std" %65 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "stateful_ld_st_ntid_chain" %gl_LocalInvocationID OpEntryPoint Kernel %1 "stateful_ld_st_ntid_chain"
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId OpExecutionMode %1 ContractionOff
OpDecorate %52 LinkageAttributes "_Z12get_local_idj" Import
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%v3ulong = OpTypeVector %ulong 3 %uint = OpTypeInt 32 0
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong %69 = OpTypeFunction %ulong %uint
%gl_LocalInvocationID = OpVariable %_ptr_Input_v3ulong Input
%uchar = OpTypeInt 8 0 %uchar = OpTypeInt 8 0
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar %_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 %_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
%uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint_0 = OpConstant %uint 0
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_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 %28 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%29 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %29 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%60 = OpLabel %63 = OpLabel
%20 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %20 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%21 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %21 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%14 = 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 %12 = OpVariable %_ptr_Function_ulong Function
OpStore %20 %28 OpStore %20 %28
OpStore %21 %29 OpStore %21 %29
%53 = OpBitcast %_ptr_Function_ulong %20 %56 = OpBitcast %_ptr_Function_ulong %20
%52 = OpLoad %ulong %53 Aligned 8 %55 = OpLoad %ulong %56 Aligned 8
%22 = OpCopyObject %ulong %52 %22 = OpCopyObject %ulong %55
%30 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %22 %30 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %22
OpStore %14 %30 OpStore %14 %30
%55 = OpBitcast %_ptr_Function_ulong %21 %58 = OpBitcast %_ptr_Function_ulong %21
%54 = OpLoad %ulong %55 Aligned 8 %57 = OpLoad %ulong %58 Aligned 8
%23 = OpCopyObject %ulong %54 %23 = OpCopyObject %ulong %57
%31 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %23 %31 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %23
OpStore %17 %31 OpStore %17 %31
%32 = OpLoad %_ptr_CrossWorkgroup_uchar %14 %32 = OpLoad %_ptr_CrossWorkgroup_uchar %14
@ -61,37 +64,36 @@
%26 = OpCopyObject %ulong %27 %26 = OpCopyObject %ulong %27
%35 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %26 %35 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %26
OpStore %18 %35 OpStore %18 %35
%74 = OpLoad %v3ulong %gl_LocalInvocationID %51 = OpFunctionCall %ulong %52 %uint_0
%51 = OpCompositeExtract %ulong %74 0 %76 = OpBitcast %ulong %51
%75 = OpBitcast %ulong %51 %37 = OpUConvert %uint %76
%37 = OpUConvert %uint %75
%36 = OpCopyObject %uint %37 %36 = OpCopyObject %uint %37
OpStore %10 %36 OpStore %10 %36
%39 = OpLoad %uint %10 %39 = OpLoad %uint %10
%76 = OpBitcast %uint %39 %77 = OpBitcast %uint %39
%38 = OpUConvert %ulong %76 %38 = OpUConvert %ulong %77
OpStore %11 %38 OpStore %11 %38
%41 = OpLoad %_ptr_CrossWorkgroup_uchar %15 %41 = OpLoad %_ptr_CrossWorkgroup_uchar %15
%42 = OpLoad %ulong %11 %42 = OpLoad %ulong %11
%56 = OpCopyObject %ulong %42 %59 = OpCopyObject %ulong %42
%77 = OpBitcast %_ptr_CrossWorkgroup_uchar %41 %78 = OpBitcast %_ptr_CrossWorkgroup_uchar %41
%78 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %77 %56 %79 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %78 %59
%40 = OpBitcast %_ptr_CrossWorkgroup_uchar %78 %40 = OpBitcast %_ptr_CrossWorkgroup_uchar %79
OpStore %16 %40 OpStore %16 %40
%44 = OpLoad %_ptr_CrossWorkgroup_uchar %18 %44 = OpLoad %_ptr_CrossWorkgroup_uchar %18
%45 = OpLoad %ulong %11 %45 = OpLoad %ulong %11
%57 = OpCopyObject %ulong %45 %60 = OpCopyObject %ulong %45
%79 = OpBitcast %_ptr_CrossWorkgroup_uchar %44 %80 = OpBitcast %_ptr_CrossWorkgroup_uchar %44
%80 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %79 %57 %81 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %80 %60
%43 = OpBitcast %_ptr_CrossWorkgroup_uchar %80 %43 = OpBitcast %_ptr_CrossWorkgroup_uchar %81
OpStore %19 %43 OpStore %19 %43
%47 = OpLoad %_ptr_CrossWorkgroup_uchar %16 %47 = OpLoad %_ptr_CrossWorkgroup_uchar %16
%58 = OpBitcast %_ptr_CrossWorkgroup_ulong %47 %61 = OpBitcast %_ptr_CrossWorkgroup_ulong %47
%46 = OpLoad %ulong %58 Aligned 8 %46 = OpLoad %ulong %61 Aligned 8
OpStore %12 %46 OpStore %12 %46
%48 = OpLoad %_ptr_CrossWorkgroup_uchar %19 %48 = OpLoad %_ptr_CrossWorkgroup_uchar %19
%49 = OpLoad %ulong %12 %49 = OpLoad %ulong %12
%59 = OpBitcast %_ptr_CrossWorkgroup_ulong %48 %62 = OpBitcast %_ptr_CrossWorkgroup_ulong %48
OpStore %59 %49 Aligned 8 OpStore %62 %49 Aligned 8
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

98
ptx/src/test/spirv_run/stateful_ld_st_ntid_sub.spvtxt
View File

@ -7,31 +7,34 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%66 = OpExtInstImport "OpenCL.std" %71 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %1 "stateful_ld_st_ntid_sub" %gl_LocalInvocationID OpEntryPoint Kernel %1 "stateful_ld_st_ntid_sub"
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId OpExecutionMode %1 ContractionOff
OpDecorate %54 LinkageAttributes "_Z12get_local_idj" Import
%void = OpTypeVoid %void = OpTypeVoid
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%v3ulong = OpTypeVector %ulong 3 %uint = OpTypeInt 32 0
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong %75 = OpTypeFunction %ulong %uint
%gl_LocalInvocationID = OpVariable %_ptr_Input_v3ulong Input
%uchar = OpTypeInt 8 0 %uchar = OpTypeInt 8 0
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar %_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 %_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
%uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%uint_0 = OpConstant %uint 0
%ulong_0 = OpConstant %ulong 0 %ulong_0 = OpConstant %ulong 0
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
%ulong_0_0 = OpConstant %ulong 0 %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 %30 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%31 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %31 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%64 = OpLabel %69 = OpLabel
%2 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %20 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%3 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %21 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%14 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %14 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%15 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %15 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%16 = 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 %10 = OpVariable %_ptr_Function_uint Function
%11 = OpVariable %_ptr_Function_ulong Function %11 = OpVariable %_ptr_Function_ulong Function
%12 = OpVariable %_ptr_Function_ulong Function %12 = OpVariable %_ptr_Function_ulong Function
OpStore %2 %30 OpStore %20 %30
OpStore %3 %31 OpStore %21 %31
%21 = OpBitcast %_ptr_Function_ulong %2 %62 = OpBitcast %_ptr_Function_ulong %20
%58 = OpLoad %ulong %21 Aligned 8 %61 = OpLoad %ulong %62 Aligned 8
%20 = OpCopyObject %ulong %58 %22 = OpCopyObject %ulong %61
%32 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %20 %32 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %22
OpStore %14 %32 OpStore %14 %32
%23 = OpBitcast %_ptr_Function_ulong %3 %64 = OpBitcast %_ptr_Function_ulong %21
%59 = OpLoad %ulong %23 Aligned 8 %63 = OpLoad %ulong %64 Aligned 8
%22 = OpCopyObject %ulong %59 %23 = OpCopyObject %ulong %63
%33 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %22 %33 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %23
OpStore %17 %33 OpStore %17 %33
%34 = OpLoad %_ptr_CrossWorkgroup_uchar %14 %34 = OpLoad %_ptr_CrossWorkgroup_uchar %14
%25 = OpConvertPtrToU %ulong %34 %25 = OpConvertPtrToU %ulong %34
@ -63,45 +66,44 @@
%26 = OpCopyObject %ulong %27 %26 = OpCopyObject %ulong %27
%37 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %26 %37 = OpConvertUToPtr %_ptr_CrossWorkgroup_uchar %26
OpStore %18 %37 OpStore %18 %37
%78 = OpLoad %v3ulong %gl_LocalInvocationID %53 = OpFunctionCall %ulong %54 %uint_0
%53 = OpCompositeExtract %ulong %78 0 %82 = OpBitcast %ulong %53
%79 = OpBitcast %ulong %53 %39 = OpUConvert %uint %82
%39 = OpUConvert %uint %79
%38 = OpCopyObject %uint %39 %38 = OpCopyObject %uint %39
OpStore %10 %38 OpStore %10 %38
%41 = OpLoad %uint %10 %41 = OpLoad %uint %10
%80 = OpBitcast %uint %41 %83 = OpBitcast %uint %41
%40 = OpUConvert %ulong %80 %40 = OpUConvert %ulong %83
OpStore %11 %40 OpStore %11 %40
%42 = OpLoad %ulong %11 %42 = OpLoad %ulong %11
%60 = OpCopyObject %ulong %42 %65 = OpCopyObject %ulong %42
%28 = OpSNegate %ulong %60 %28 = OpSNegate %ulong %65
%44 = OpLoad %_ptr_CrossWorkgroup_uchar %15 %44 = OpLoad %_ptr_CrossWorkgroup_uchar %15
%81 = OpBitcast %_ptr_CrossWorkgroup_uchar %44 %84 = OpBitcast %_ptr_CrossWorkgroup_uchar %44
%82 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %81 %28 %85 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %84 %28
%43 = OpBitcast %_ptr_CrossWorkgroup_uchar %82 %43 = OpBitcast %_ptr_CrossWorkgroup_uchar %85
OpStore %16 %43 OpStore %16 %43
%45 = OpLoad %ulong %11 %45 = OpLoad %ulong %11
%61 = OpCopyObject %ulong %45 %66 = OpCopyObject %ulong %45
%29 = OpSNegate %ulong %61 %29 = OpSNegate %ulong %66
%47 = OpLoad %_ptr_CrossWorkgroup_uchar %18 %47 = OpLoad %_ptr_CrossWorkgroup_uchar %18
%83 = OpBitcast %_ptr_CrossWorkgroup_uchar %47 %86 = OpBitcast %_ptr_CrossWorkgroup_uchar %47
%84 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %83 %29 %87 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %86 %29
%46 = OpBitcast %_ptr_CrossWorkgroup_uchar %84 %46 = OpBitcast %_ptr_CrossWorkgroup_uchar %87
OpStore %19 %46 OpStore %19 %46
%49 = OpLoad %_ptr_CrossWorkgroup_uchar %16 %49 = OpLoad %_ptr_CrossWorkgroup_uchar %16
%62 = OpBitcast %_ptr_CrossWorkgroup_ulong %49 %67 = OpBitcast %_ptr_CrossWorkgroup_ulong %49
%86 = OpBitcast %_ptr_CrossWorkgroup_uchar %62 %89 = OpBitcast %_ptr_CrossWorkgroup_uchar %67
%87 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %86 %ulong_0 %90 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %89 %ulong_0
%55 = OpBitcast %_ptr_CrossWorkgroup_ulong %87 %58 = OpBitcast %_ptr_CrossWorkgroup_ulong %90
%48 = OpLoad %ulong %55 Aligned 8 %48 = OpLoad %ulong %58 Aligned 8
OpStore %12 %48 OpStore %12 %48
%50 = OpLoad %_ptr_CrossWorkgroup_uchar %19 %50 = OpLoad %_ptr_CrossWorkgroup_uchar %19
%51 = OpLoad %ulong %12 %51 = OpLoad %ulong %12
%63 = OpBitcast %_ptr_CrossWorkgroup_ulong %50 %68 = OpBitcast %_ptr_CrossWorkgroup_ulong %50
%88 = OpBitcast %_ptr_CrossWorkgroup_uchar %63 %91 = OpBitcast %_ptr_CrossWorkgroup_uchar %68
%89 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %88 %ulong_0_0 %92 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %91 %ulong_0_0
%57 = OpBitcast %_ptr_CrossWorkgroup_ulong %89 %60 = OpBitcast %_ptr_CrossWorkgroup_ulong %92
OpStore %57 %51 Aligned 8 OpStore %60 %51 Aligned 8
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

121
ptx/src/test/spirv_run/vector4.spvtxt
View File

@ -7,93 +7,50 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%51 = OpExtInstImport "OpenCL.std" %24 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %25 "vector" OpEntryPoint Kernel %1 "vector4"
OpExecutionMode %1 ContractionOff
%void = OpTypeVoid %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 %ulong = OpTypeInt 64 0
%67 = OpTypeFunction %void %ulong %ulong %27 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_v2uint = OpTypePointer Generic %v2uint %uint = OpTypeInt 32 0
%1 = OpFunction %v2uint None %55 %v4uint = OpTypeVector %uint 4
%7 = OpFunctionParameter %v2uint %_ptr_Function_v4uint = OpTypePointer Function %v4uint
%24 = OpLabel %_ptr_Function_uint = OpTypePointer Function %uint
%3 = OpVariable %_ptr_Function_v2uint Function %_ptr_Generic_v4uint = OpTypePointer Generic %v4uint
%2 = OpVariable %_ptr_Function_v2uint Function %uint_3 = OpConstant %uint 3
%4 = OpVariable %_ptr_Function_v2uint Function %_ptr_Generic_uint = OpTypePointer Generic %uint
%5 = OpVariable %_ptr_Function_uint Function %1 = OpFunction %void None %27
%6 = OpVariable %_ptr_Function_uint Function %8 = OpFunctionParameter %ulong
OpStore %3 %7 %9 = OpFunctionParameter %ulong
%59 = OpInBoundsAccessChain %_ptr_Function_uint %3 %uint_0 %22 = OpLabel
%9 = OpLoad %uint %59 %2 = OpVariable %_ptr_Function_ulong Function
%8 = OpCopyObject %uint %9 %3 = OpVariable %_ptr_Function_ulong Function
OpStore %5 %8 %4 = OpVariable %_ptr_Function_ulong Function
%61 = OpInBoundsAccessChain %_ptr_Function_uint %3 %uint_1 %5 = OpVariable %_ptr_Function_ulong Function
%11 = OpLoad %uint %61 %6 = OpVariable %_ptr_Function_v4uint Function
%10 = OpCopyObject %uint %11 %7 = OpVariable %_ptr_Function_uint Function
OpStore %6 %10 OpStore %2 %8
%13 = OpLoad %uint %5 OpStore %3 %9
%14 = OpLoad %uint %6 %10 = OpLoad %ulong %2 Aligned 8
%12 = OpIAdd %uint %13 %14 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 OpStore %6 %12
%16 = OpLoad %uint %6 %35 = OpInBoundsAccessChain %_ptr_Function_uint %6 %uint_3
%15 = OpCopyObject %uint %16 %15 = OpLoad %uint %35
%62 = OpInBoundsAccessChain %_ptr_Function_uint %4 %uint_0 %20 = OpCopyObject %uint %15
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
%19 = OpCopyObject %uint %20 %19 = OpCopyObject %uint %20
%65 = OpInBoundsAccessChain %_ptr_Function_uint %4 %uint_0 %14 = OpCopyObject %uint %19
OpStore %65 %19 OpStore %7 %14
%22 = OpLoad %v2uint %4 %16 = OpLoad %ulong %5
%21 = OpCopyObject %v2uint %22 %17 = OpLoad %uint %7
OpStore %2 %21 %21 = OpConvertUToPtr %_ptr_Generic_uint %16
%23 = OpLoad %v2uint %2 OpStore %21 %17 Aligned 4
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
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

60
ptx/src/translate.rs
View File

@ -431,7 +431,7 @@ pub fn to_spirv_module<'a>(ast: ast::Module<'a>) -> Result<Module, TranslateErro
}) })
.collect::<Result<Vec<_>, _>>()?; .collect::<Result<Vec<_>, _>>()?;
let must_link_ptx_impl = ptx_impl_imports.len() > 0; let must_link_ptx_impl = ptx_impl_imports.len() > 0;
let directives = ptx_impl_imports let mut directives = ptx_impl_imports
.into_iter() .into_iter()
.map(|(_, v)| v) .map(|(_, v)| v)
.chain(directives.into_iter()) .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(); let mut builder = dr::Builder::new();
builder.reserve_ids(id_defs.current_id()); builder.reserve_ids(id_defs.current_id());
let call_map = get_kernels_call_map(&directives); 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); normalize_variable_decls(&mut directives);
let denorm_information = compute_denorm_information(&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 // 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(); let empty_body = Vec::new();
for d in directives.iter() { for d in directives.iter() {
match d { match d {
Directive::Variable(_, var) => { Directive::Variable(linking, var) => {
emit_variable(builder, map, &var)?; emit_variable(builder, map, id_defs, *linking, &var)?;
} }
Directive::Method(f) => { Directive::Method(f) => {
let f_body = match &f.body { let f_body = match &f.body {
@ -547,7 +547,7 @@ fn emit_directives<'input>(
} }
}; };
for var in f.globals.iter() { 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 func_decl = (*f.func_decl).borrow();
let fn_id = emit_function_header( 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()?; builder.end_function()?;
if let ( if let (
ast::MethodDeclaration { 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, builder: &mut dr::Builder,
map: &mut TypeWordMap, map: &mut TypeWordMap,
id_defs: &GlobalStringIdResolver<'input>,
opencl: spirv::Word, opencl: spirv::Word,
func: &[ExpandedStatement], func: &[ExpandedStatement],
) -> Result<(), TranslateError> { ) -> Result<(), TranslateError> {
@ -2541,7 +2542,7 @@ fn emit_function_body_ops(
builder.function_call(result_type, result_id, call.name, arg_list)?; builder.function_call(result_type, result_id, call.name, arg_list)?;
} }
Statement::Variable(var) => { Statement::Variable(var) => {
emit_variable(builder, map, var)?; emit_variable(builder, map, id_defs, ast::LinkingDirective::NONE, var)?;
} }
Statement::Constant(cnst) => { Statement::Constant(cnst) => {
let typ_id = map.get_or_add_scalar(builder, cnst.typ); 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 result
} }
fn emit_variable( fn emit_variable<'input>(
builder: &mut dr::Builder, builder: &mut dr::Builder,
map: &mut TypeWordMap, map: &mut TypeWordMap,
id_defs: &GlobalStringIdResolver<'input>,
linking: ast::LinkingDirective,
var: &ast::Variable<spirv::Word>, var: &ast::Variable<spirv::Word>,
) -> Result<(), TranslateError> { ) -> Result<(), TranslateError> {
let (must_init, st_class) = match var.state_space { let (must_init, st_class) = match var.state_space {
@ -3323,9 +3326,45 @@ fn emit_variable(
[dr::Operand::LiteralInt32(align)].iter().cloned(), [dr::Operand::LiteralInt32(align)].iter().cloned(),
); );
} }
emit_linking_decoration(builder, id_defs, var.name, linking);
Ok(()) 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( fn emit_mad_uint(
builder: &mut dr::Builder, builder: &mut dr::Builder,
map: &mut TypeWordMap, map: &mut TypeWordMap,
@ -4902,6 +4941,7 @@ impl<'input> FnSigMapper<'input> {
struct GlobalStringIdResolver<'input> { struct GlobalStringIdResolver<'input> {
current_id: spirv::Word, current_id: spirv::Word,
variables: HashMap<Cow<'input, str>, 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)>>, variables_type_check: HashMap<u32, Option<(ast::Type, ast::StateSpace, bool)>>,
special_registers: SpecialRegistersMap, special_registers: SpecialRegistersMap,
fns: HashMap<spirv::Word, FnSigMapper<'input>>, fns: HashMap<spirv::Word, FnSigMapper<'input>>,
@ -4912,6 +4952,7 @@ impl<'input> GlobalStringIdResolver<'input> {
Self { Self {
current_id: start_id, current_id: start_id,
variables: HashMap::new(), variables: HashMap::new(),
reverse_variables: HashMap::new(),
variables_type_check: HashMap::new(), variables_type_check: HashMap::new(),
special_registers: SpecialRegistersMap::new(), special_registers: SpecialRegistersMap::new(),
fns: HashMap::new(), fns: HashMap::new(),
@ -4942,6 +4983,7 @@ impl<'input> GlobalStringIdResolver<'input> {
hash_map::Entry::Vacant(e) => { hash_map::Entry::Vacant(e) => {
let numeric_id = self.current_id; let numeric_id = self.current_id;
e.insert(numeric_id); e.insert(numeric_id);
self.reverse_variables.insert(numeric_id, id);
self.current_id += 1; self.current_id += 1;
numeric_id numeric_id
} }