diff options
| author | Andrzej Janik <[email protected]> | 2020-11-05 00:27:46 +0100 |
|---|---|---|
| committer | Andrzej Janik <[email protected]> | 2020-11-05 00:27:46 +0100 |
| commit | 96702d86c96ef2d14795a71af43015a8eacd0a94 (patch) | |
| tree | 7da0d15a396f31393668fe6ad2a8e999b6742fcc /ptx/src/test/spirv_run | |
| parent | e5a53ed5d30fad3d8ebae6d72ead1564d2b97275 (diff) | |
| download | ZLUDA-96702d86c96ef2d14795a71af43015a8eacd0a94.tar.gz ZLUDA-96702d86c96ef2d14795a71af43015a8eacd0a94.zip | |
Fix issues with .param/.local and implement sin, cos, ex2, lg2
Diffstat (limited to 'ptx/src/test/spirv_run')
57 files changed, 1912 insertions, 1760 deletions
diff --git a/ptx/src/test/spirv_run/add.spvtxt b/ptx/src/test/spirv_run/add.spvtxt index 6810fec..d9a5b9e 100644 --- a/ptx/src/test/spirv_run/add.spvtxt +++ b/ptx/src/test/spirv_run/add.spvtxt @@ -2,21 +2,24 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %25 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %23 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "add" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %28 = OpTypeFunction %void %ulong %ulong + %26 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_1 = OpConstant %ulong 1 - %1 = OpFunction %void None %28 + %1 = OpFunction %void None %26 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %23 = OpLabel + %21 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -25,22 +28,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpIAdd %ulong %17 %ulong_1 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %22 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpIAdd %ulong %15 %ulong_1 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %20 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/and.spvtxt b/ptx/src/test/spirv_run/and.spvtxt index 8358c28..57c36e6 100644 --- a/ptx/src/test/spirv_run/and.spvtxt +++ b/ptx/src/test/spirv_run/and.spvtxt @@ -1,66 +1,58 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 41 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%33 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "and" -%34 = OpTypeVoid -%35 = OpTypeInt 64 0 -%36 = OpTypeFunction %34 %35 %35 -%37 = OpTypePointer Function %35 -%38 = OpTypeInt 32 0 -%39 = OpTypePointer Function %38 -%40 = OpTypePointer Generic %38 -%23 = OpConstant %35 4 -%1 = OpFunction %34 None %36 -%8 = OpFunctionParameter %35 -%9 = OpFunctionParameter %35 -%31 = OpLabel -%2 = OpVariable %37 Function -%3 = OpVariable %37 Function -%4 = OpVariable %37 Function -%5 = OpVariable %37 Function -%6 = OpVariable %39 Function -%7 = OpVariable %39 Function -OpStore %2 %8 -OpStore %3 %9 -%11 = OpLoad %35 %2 -%10 = OpCopyObject %35 %11 -OpStore %4 %10 -%13 = OpLoad %35 %3 -%12 = OpCopyObject %35 %13 -OpStore %5 %12 -%15 = OpLoad %35 %4 -%25 = OpConvertUToPtr %40 %15 -%14 = OpLoad %38 %25 -OpStore %6 %14 -%17 = OpLoad %35 %4 -%24 = OpIAdd %35 %17 %23 -%26 = OpConvertUToPtr %40 %24 -%16 = OpLoad %38 %26 -OpStore %7 %16 -%19 = OpLoad %38 %6 -%20 = OpLoad %38 %7 -%28 = OpCopyObject %38 %19 -%29 = OpCopyObject %38 %20 -%27 = OpBitwiseAnd %38 %28 %29 -%18 = OpCopyObject %38 %27 -OpStore %6 %18 -%21 = OpLoad %35 %5 -%22 = OpLoad %38 %6 -%30 = OpConvertUToPtr %40 %21 -OpStore %30 %22 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %31 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "and" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %34 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %ulong_4 = OpConstant %ulong 4 + %1 = OpFunction %void None %34 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %29 = 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_uint Function + %7 = OpVariable %_ptr_Function_uint Function + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %4 %10 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %23 + OpStore %6 %12 + %15 = OpLoad %ulong %4 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_uint %22 + %14 = OpLoad %uint %24 + OpStore %7 %14 + %17 = OpLoad %uint %6 + %18 = OpLoad %uint %7 + %26 = OpCopyObject %uint %17 + %27 = OpCopyObject %uint %18 + %25 = OpBitwiseAnd %uint %26 %27 + %16 = OpCopyObject %uint %25 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %6 + %28 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %28 %20 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/atom_add.spvtxt b/ptx/src/test/spirv_run/atom_add.spvtxt index 2c83fe9..6a977e0 100644 --- a/ptx/src/test/spirv_run/atom_add.spvtxt +++ b/ptx/src/test/spirv_run/atom_add.spvtxt @@ -1,84 +1,76 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 55 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%40 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "atom_add" %4 -OpDecorate %4 Alignment 4 -%41 = OpTypeVoid -%42 = OpTypeInt 32 0 -%43 = OpTypeInt 8 0 -%44 = OpConstant %42 1024 -%45 = OpTypeArray %43 %44 -%46 = OpTypePointer Workgroup %45 -%4 = OpVariable %46 Workgroup -%47 = OpTypeInt 64 0 -%48 = OpTypeFunction %41 %47 %47 -%49 = OpTypePointer Function %47 -%50 = OpTypePointer Function %42 -%51 = OpTypePointer Generic %42 -%27 = OpConstant %47 4 -%52 = OpTypePointer Workgroup %42 -%53 = OpConstant %42 1 -%54 = OpConstant %42 0 -%29 = OpConstant %47 4 -%1 = OpFunction %41 None %48 -%9 = OpFunctionParameter %47 -%10 = OpFunctionParameter %47 -%38 = OpLabel -%2 = OpVariable %49 Function -%3 = OpVariable %49 Function -%5 = OpVariable %49 Function -%6 = OpVariable %49 Function -%7 = OpVariable %50 Function -%8 = OpVariable %50 Function -OpStore %2 %9 -OpStore %3 %10 -%12 = OpLoad %47 %2 -%11 = OpCopyObject %47 %12 -OpStore %5 %11 -%14 = OpLoad %47 %3 -%13 = OpCopyObject %47 %14 -OpStore %6 %13 -%16 = OpLoad %47 %5 -%31 = OpConvertUToPtr %51 %16 -%15 = OpLoad %42 %31 -OpStore %7 %15 -%18 = OpLoad %47 %5 -%28 = OpIAdd %47 %18 %27 -%32 = OpConvertUToPtr %51 %28 -%17 = OpLoad %42 %32 -OpStore %8 %17 -%19 = OpLoad %42 %7 -%33 = OpBitcast %52 %4 -OpStore %33 %19 -%21 = OpLoad %42 %8 -%34 = OpBitcast %52 %4 -%20 = OpAtomicIAdd %42 %34 %53 %54 %21 -OpStore %7 %20 -%35 = OpBitcast %52 %4 -%22 = OpLoad %42 %35 -OpStore %8 %22 -%23 = OpLoad %47 %6 -%24 = OpLoad %42 %7 -%36 = OpConvertUToPtr %51 %23 -OpStore %36 %24 -%25 = OpLoad %47 %6 -%26 = OpLoad %42 %8 -%30 = OpIAdd %47 %25 %29 -%37 = OpConvertUToPtr %51 %30 -OpStore %37 %26 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %38 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "atom_add" %4 + OpDecorate %4 Alignment 4 + %void = OpTypeVoid + %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 + %uint_1024 = OpConstant %uint 1024 +%_arr_uchar_uint_1024 = OpTypeArray %uchar %uint_1024 +%_ptr_Workgroup__arr_uchar_uint_1024 = OpTypePointer Workgroup %_arr_uchar_uint_1024 + %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup + %ulong = OpTypeInt 64 0 + %46 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %ulong_4 = OpConstant %ulong 4 +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %uint_1 = OpConstant %uint 1 + %uint_0 = OpConstant %uint 0 + %ulong_4_0 = OpConstant %ulong 4 + %1 = OpFunction %void None %46 + %9 = OpFunctionParameter %ulong + %10 = OpFunctionParameter %ulong + %36 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_ulong Function + %7 = OpVariable %_ptr_Function_uint Function + %8 = OpVariable %_ptr_Function_uint Function + OpStore %2 %9 + OpStore %3 %10 + %11 = OpLoad %ulong %2 + OpStore %5 %11 + %12 = OpLoad %ulong %3 + OpStore %6 %12 + %14 = OpLoad %ulong %5 + %29 = OpConvertUToPtr %_ptr_Generic_uint %14 + %13 = OpLoad %uint %29 + OpStore %7 %13 + %16 = OpLoad %ulong %5 + %26 = OpIAdd %ulong %16 %ulong_4 + %30 = OpConvertUToPtr %_ptr_Generic_uint %26 + %15 = OpLoad %uint %30 + OpStore %8 %15 + %17 = OpLoad %uint %7 + %31 = OpBitcast %_ptr_Workgroup_uint %4 + OpStore %31 %17 + %19 = OpLoad %uint %8 + %32 = OpBitcast %_ptr_Workgroup_uint %4 + %18 = OpAtomicIAdd %uint %32 %uint_1 %uint_0 %19 + OpStore %7 %18 + %33 = OpBitcast %_ptr_Workgroup_uint %4 + %20 = OpLoad %uint %33 + OpStore %8 %20 + %21 = OpLoad %ulong %6 + %22 = OpLoad %uint %7 + %34 = OpConvertUToPtr %_ptr_Generic_uint %21 + OpStore %34 %22 + %23 = OpLoad %ulong %6 + %24 = OpLoad %uint %8 + %28 = OpIAdd %ulong %23 %ulong_4_0 + %35 = OpConvertUToPtr %_ptr_Generic_uint %28 + OpStore %35 %24 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/atom_cas.spvtxt b/ptx/src/test/spirv_run/atom_cas.spvtxt index c5fb922..b28d3bc 100644 --- a/ptx/src/test/spirv_run/atom_cas.spvtxt +++ b/ptx/src/test/spirv_run/atom_cas.spvtxt @@ -1,77 +1,69 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 51 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%41 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "atom_cas" -%42 = OpTypeVoid -%43 = OpTypeInt 64 0 -%44 = OpTypeFunction %42 %43 %43 -%45 = OpTypePointer Function %43 -%46 = OpTypeInt 32 0 -%47 = OpTypePointer Function %46 -%48 = OpTypePointer Generic %46 -%25 = OpConstant %43 4 -%27 = OpConstant %46 100 -%49 = OpConstant %46 1 -%50 = OpConstant %46 0 -%28 = OpConstant %43 4 -%30 = OpConstant %43 4 -%1 = OpFunction %42 None %44 -%8 = OpFunctionParameter %43 -%9 = OpFunctionParameter %43 -%39 = OpLabel -%2 = OpVariable %45 Function -%3 = OpVariable %45 Function -%4 = OpVariable %45 Function -%5 = OpVariable %45 Function -%6 = OpVariable %47 Function -%7 = OpVariable %47 Function -OpStore %2 %8 -OpStore %3 %9 -%11 = OpLoad %43 %2 -%10 = OpCopyObject %43 %11 -OpStore %4 %10 -%13 = OpLoad %43 %3 -%12 = OpCopyObject %43 %13 -OpStore %5 %12 -%15 = OpLoad %43 %4 -%32 = OpConvertUToPtr %48 %15 -%14 = OpLoad %46 %32 -OpStore %6 %14 -%17 = OpLoad %43 %4 -%18 = OpLoad %46 %6 -%26 = OpIAdd %43 %17 %25 -%34 = OpConvertUToPtr %48 %26 -%35 = OpCopyObject %46 %18 -%33 = OpAtomicCompareExchange %46 %34 %49 %50 %50 %27 %35 -%16 = OpCopyObject %46 %33 -OpStore %6 %16 -%20 = OpLoad %43 %4 -%29 = OpIAdd %43 %20 %28 -%36 = OpConvertUToPtr %48 %29 -%19 = OpLoad %46 %36 -OpStore %7 %19 -%21 = OpLoad %43 %5 -%22 = OpLoad %46 %6 -%37 = OpConvertUToPtr %48 %21 -OpStore %37 %22 -%23 = OpLoad %43 %5 -%24 = OpLoad %46 %7 -%31 = OpIAdd %43 %23 %30 -%38 = OpConvertUToPtr %48 %31 -OpStore %38 %24 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %39 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "atom_cas" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %42 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %ulong_4 = OpConstant %ulong 4 + %uint_100 = OpConstant %uint 100 + %uint_1 = OpConstant %uint 1 + %uint_0 = OpConstant %uint 0 + %ulong_4_0 = OpConstant %ulong 4 + %ulong_4_1 = OpConstant %ulong 4 + %1 = OpFunction %void None %42 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %37 = 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_uint Function + %7 = OpVariable %_ptr_Function_uint Function + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %4 %10 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %30 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %30 + OpStore %6 %12 + %15 = OpLoad %ulong %4 + %16 = OpLoad %uint %6 + %24 = OpIAdd %ulong %15 %ulong_4 + %32 = OpConvertUToPtr %_ptr_Generic_uint %24 + %33 = OpCopyObject %uint %16 + %31 = OpAtomicCompareExchange %uint %32 %uint_1 %uint_0 %uint_0 %uint_100 %33 + %14 = OpCopyObject %uint %31 + OpStore %6 %14 + %18 = OpLoad %ulong %4 + %27 = OpIAdd %ulong %18 %ulong_4_0 + %34 = OpConvertUToPtr %_ptr_Generic_uint %27 + %17 = OpLoad %uint %34 + OpStore %7 %17 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %6 + %35 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %35 %20 + %21 = OpLoad %ulong %5 + %22 = OpLoad %uint %7 + %29 = OpIAdd %ulong %21 %ulong_4_1 + %36 = OpConvertUToPtr %_ptr_Generic_uint %29 + OpStore %36 %22 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/b64tof64.spvtxt b/ptx/src/test/spirv_run/b64tof64.spvtxt index 9146c90..e8cfcf4 100644 --- a/ptx/src/test/spirv_run/b64tof64.spvtxt +++ b/ptx/src/test/spirv_run/b64tof64.spvtxt @@ -2,23 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %26 = OpExtInstImport "OpenCL.std" + %24 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "b64tof64" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %29 = OpTypeFunction %void %ulong %ulong + %27 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %double = OpTypeFloat 64 %_ptr_Function_double = OpTypePointer Function %double %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %29 + %1 = OpFunction %void None %27 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %24 = OpLabel + %22 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_double Function @@ -27,24 +29,22 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %20 = OpBitcast %double %11 - %10 = OpCopyObject %double %20 + %18 = OpBitcast %_ptr_Function_double %2 + %10 = OpLoad %double %18 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %6 %12 - %15 = OpLoad %double %4 - %21 = OpBitcast %ulong %15 - %14 = OpCopyObject %ulong %21 - OpStore %5 %14 - %17 = OpLoad %ulong %5 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %17 - %16 = OpLoad %ulong %22 - OpStore %7 %16 - %18 = OpLoad %ulong %6 - %19 = OpLoad %ulong %7 - %23 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %23 %19 + %11 = OpLoad %ulong %3 + OpStore %6 %11 + %13 = OpLoad %double %4 + %19 = OpBitcast %ulong %13 + %12 = OpCopyObject %ulong %19 + OpStore %5 %12 + %15 = OpLoad %ulong %5 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %15 + %14 = OpLoad %ulong %20 + OpStore %7 %14 + %16 = OpLoad %ulong %6 + %17 = OpLoad %ulong %7 + %21 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %21 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/block.spvtxt b/ptx/src/test/spirv_run/block.spvtxt index 534167d..fe7e63a 100644 --- a/ptx/src/test/spirv_run/block.spvtxt +++ b/ptx/src/test/spirv_run/block.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %29 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %27 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "block" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %32 = OpTypeFunction %void %ulong %ulong + %30 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_1 = OpConstant %ulong 1 %ulong_1_0 = OpConstant %ulong 1 - %1 = OpFunction %void None %32 + %1 = OpFunction %void None %30 %9 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong - %27 = OpLabel + %25 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -27,25 +30,23 @@ %8 = OpVariable %_ptr_Function_ulong Function OpStore %2 %9 OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %2 OpStore %4 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %5 %13 - %16 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_ulong %16 - %15 = OpLoad %ulong %25 - OpStore %6 %15 - %18 = OpLoad %ulong %6 - %17 = OpIAdd %ulong %18 %ulong_1 - OpStore %7 %17 - %20 = OpLoad %ulong %8 - %19 = OpIAdd %ulong %20 %ulong_1_0 - OpStore %8 %19 - %21 = OpLoad %ulong %5 - %22 = OpLoad %ulong %7 - %26 = OpConvertUToPtr %_ptr_Generic_ulong %21 - OpStore %26 %22 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_ulong %14 + %13 = OpLoad %ulong %23 + OpStore %6 %13 + %16 = OpLoad %ulong %6 + %15 = OpIAdd %ulong %16 %ulong_1 + OpStore %7 %15 + %18 = OpLoad %ulong %8 + %17 = OpIAdd %ulong %18 %ulong_1_0 + OpStore %8 %17 + %19 = OpLoad %ulong %5 + %20 = OpLoad %ulong %7 + %24 = OpConvertUToPtr %_ptr_Generic_ulong %19 + OpStore %24 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/bra.spvtxt b/ptx/src/test/spirv_run/bra.spvtxt index f59fda5..b20e61a 100644 --- a/ptx/src/test/spirv_run/bra.spvtxt +++ b/ptx/src/test/spirv_run/bra.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %31 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %29 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "bra" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %34 = OpTypeFunction %void %ulong %ulong + %32 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_1 = OpConstant %ulong 1 %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %34 + %1 = OpFunction %void None %32 %11 = OpFunctionParameter %ulong %12 = OpFunctionParameter %ulong - %29 = OpLabel + %27 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function @@ -26,31 +29,29 @@ %10 = OpVariable %_ptr_Function_ulong Function OpStore %2 %11 OpStore %3 %12 - %14 = OpLoad %ulong %2 - %13 = OpCopyObject %ulong %14 + %13 = OpLoad %ulong %2 OpStore %7 %13 - %16 = OpLoad %ulong %3 - %15 = OpCopyObject %ulong %16 - OpStore %8 %15 - %18 = OpLoad %ulong %7 - %27 = OpConvertUToPtr %_ptr_Generic_ulong %18 - %17 = OpLoad %ulong %27 - OpStore %9 %17 + %14 = OpLoad %ulong %3 + OpStore %8 %14 + %16 = OpLoad %ulong %7 + %25 = OpConvertUToPtr %_ptr_Generic_ulong %16 + %15 = OpLoad %ulong %25 + OpStore %9 %15 OpBranch %4 %4 = OpLabel + %18 = OpLoad %ulong %9 + %17 = OpIAdd %ulong %18 %ulong_1 + OpStore %10 %17 + OpBranch %6 + %35 = OpLabel %20 = OpLoad %ulong %9 - %19 = OpIAdd %ulong %20 %ulong_1 + %19 = OpIAdd %ulong %20 %ulong_2 OpStore %10 %19 OpBranch %6 - %37 = OpLabel - %22 = OpLoad %ulong %9 - %21 = OpIAdd %ulong %22 %ulong_2 - OpStore %10 %21 - OpBranch %6 %6 = OpLabel - %23 = OpLoad %ulong %8 - %24 = OpLoad %ulong %10 - %28 = OpConvertUToPtr %_ptr_Generic_ulong %23 - OpStore %28 %24 + %21 = OpLoad %ulong %8 + %22 = OpLoad %ulong %10 + %26 = OpConvertUToPtr %_ptr_Generic_ulong %21 + OpStore %26 %22 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/call.spvtxt b/ptx/src/test/spirv_run/call.spvtxt index ca4685a..31f5307 100644 --- a/ptx/src/test/spirv_run/call.spvtxt +++ b/ptx/src/test/spirv_run/call.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %47 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %37 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %4 "call" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %50 = OpTypeFunction %void %ulong %ulong + %40 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong - %53 = OpTypeFunction %ulong %ulong + %44 = OpTypeFunction %void %_ptr_Function_ulong %_ptr_Function_ulong %ulong_1 = OpConstant %ulong 1 - %4 = OpFunction %void None %50 + %4 = OpFunction %void None %40 %12 = OpFunctionParameter %ulong %13 = OpFunctionParameter %ulong - %32 = OpLabel + %26 = OpLabel %5 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function @@ -27,49 +30,38 @@ %11 = OpVariable %_ptr_Function_ulong Function OpStore %5 %12 OpStore %6 %13 - %15 = OpLoad %ulong %5 - %14 = OpCopyObject %ulong %15 + %14 = OpLoad %ulong %5 OpStore %7 %14 - %17 = OpLoad %ulong %6 - %16 = OpCopyObject %ulong %17 - OpStore %8 %16 - %19 = OpLoad %ulong %7 - %28 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %19 - %18 = OpLoad %ulong %28 - OpStore %9 %18 + %15 = OpLoad %ulong %6 + OpStore %8 %15 + %17 = OpLoad %ulong %7 + %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %17 + %16 = OpLoad %ulong %22 + OpStore %9 %16 + %18 = OpLoad %ulong %9 + %23 = OpBitcast %_ptr_Function_ulong %10 + %24 = OpCopyObject %ulong %18 + OpStore %23 %24 + %43 = OpFunctionCall %void %1 %11 %10 + %19 = OpLoad %ulong %11 + OpStore %9 %19 + %20 = OpLoad %ulong %8 %21 = OpLoad %ulong %9 - %29 = OpCopyObject %ulong %21 - %30 = OpCopyObject %ulong %29 - %20 = OpCopyObject %ulong %30 - OpStore %10 %20 - %23 = OpLoad %ulong %10 - %22 = OpFunctionCall %ulong %1 %23 - OpStore %11 %22 - %25 = OpLoad %ulong %11 - %24 = OpCopyObject %ulong %25 - OpStore %9 %24 - %26 = OpLoad %ulong %8 - %27 = OpLoad %ulong %9 - %31 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %26 - OpStore %31 %27 + %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %20 + OpStore %25 %21 OpReturn OpFunctionEnd - %1 = OpFunction %ulong None %53 - %36 = OpFunctionParameter %ulong - %45 = OpLabel - %34 = OpVariable %_ptr_Function_ulong Function - %33 = OpVariable %_ptr_Function_ulong Function - %35 = OpVariable %_ptr_Function_ulong Function - OpStore %34 %36 - %38 = OpLoad %ulong %34 - %37 = OpCopyObject %ulong %38 - OpStore %35 %37 - %40 = OpLoad %ulong %35 - %39 = OpIAdd %ulong %40 %ulong_1 - OpStore %35 %39 - %42 = OpLoad %ulong %35 - %41 = OpCopyObject %ulong %42 - OpStore %33 %41 - %43 = OpLoad %ulong %33 - OpReturnValue %43 + %1 = OpFunction %void None %44 + %27 = OpFunctionParameter %_ptr_Function_ulong + %28 = OpFunctionParameter %_ptr_Function_ulong + %35 = OpLabel + %29 = OpVariable %_ptr_Function_ulong Function + %30 = OpLoad %ulong %28 + OpStore %29 %30 + %32 = OpLoad %ulong %29 + %31 = OpIAdd %ulong %32 %ulong_1 + OpStore %29 %31 + %33 = OpLoad %ulong %29 + OpStore %27 %33 + OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/constant_f32.spvtxt b/ptx/src/test/spirv_run/constant_f32.spvtxt index 27c5f4e..46193a2 100644 --- a/ptx/src/test/spirv_run/constant_f32.spvtxt +++ b/ptx/src/test/spirv_run/constant_f32.spvtxt @@ -1,57 +1,48 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 32 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%24 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "constant_f32" -; OpDecorate %1 FunctionDenormModeINTEL 32 Preserve -%25 = OpTypeVoid -%26 = OpTypeInt 64 0 -%27 = OpTypeFunction %25 %26 %26 -%28 = OpTypePointer Function %26 -%29 = OpTypeFloat 32 -%30 = OpTypePointer Function %29 -%31 = OpTypePointer Generic %29 -%19 = OpConstant %29 0.5 -%1 = OpFunction %25 None %27 -%7 = OpFunctionParameter %26 -%8 = OpFunctionParameter %26 -%22 = OpLabel -%2 = OpVariable %28 Function -%3 = OpVariable %28 Function -%4 = OpVariable %28 Function -%5 = OpVariable %28 Function -%6 = OpVariable %30 Function -OpStore %2 %7 -OpStore %3 %8 -%10 = OpLoad %26 %2 -%9 = OpCopyObject %26 %10 -OpStore %4 %9 -%12 = OpLoad %26 %3 -%11 = OpCopyObject %26 %12 -OpStore %5 %11 -%14 = OpLoad %26 %4 -%20 = OpConvertUToPtr %31 %14 -%13 = OpLoad %29 %20 -OpStore %6 %13 -%16 = OpLoad %29 %6 -%15 = OpFMul %29 %16 %19 -OpStore %6 %15 -%17 = OpLoad %26 %5 -%18 = OpLoad %29 %6 -%21 = OpConvertUToPtr %31 %17 -OpStore %21 %18 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %22 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "constant_f32" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %25 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %float_0_5 = OpConstant %float 0.5 + %1 = OpFunction %void None %25 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %20 = 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_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %18 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpFMul %float %14 %float_0_5 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %19 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %19 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/constant_negative.spvtxt b/ptx/src/test/spirv_run/constant_negative.spvtxt index ec2ff72..5532e6e 100644 --- a/ptx/src/test/spirv_run/constant_negative.spvtxt +++ b/ptx/src/test/spirv_run/constant_negative.spvtxt @@ -1,56 +1,48 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 32 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%24 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "constant_negative" -%25 = OpTypeVoid -%26 = OpTypeInt 64 0 -%27 = OpTypeFunction %25 %26 %26 -%28 = OpTypePointer Function %26 -%29 = OpTypeInt 32 0 -%30 = OpTypePointer Function %29 -%31 = OpTypePointer Generic %29 -%19 = OpConstant %29 4294967295 -%1 = OpFunction %25 None %27 -%7 = OpFunctionParameter %26 -%8 = OpFunctionParameter %26 -%22 = OpLabel -%2 = OpVariable %28 Function -%3 = OpVariable %28 Function -%4 = OpVariable %28 Function -%5 = OpVariable %28 Function -%6 = OpVariable %30 Function -OpStore %2 %7 -OpStore %3 %8 -%10 = OpLoad %26 %2 -%9 = OpCopyObject %26 %10 -OpStore %4 %9 -%12 = OpLoad %26 %3 -%11 = OpCopyObject %26 %12 -OpStore %5 %11 -%14 = OpLoad %26 %4 -%20 = OpConvertUToPtr %31 %14 -%13 = OpLoad %29 %20 -OpStore %6 %13 -%16 = OpLoad %29 %6 -%15 = OpIMul %29 %16 %19 -OpStore %6 %15 -%17 = OpLoad %26 %5 -%18 = OpLoad %29 %6 -%21 = OpConvertUToPtr %31 %17 -OpStore %21 %18 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %22 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "constant_negative" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %25 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint +%uint_4294967295 = OpConstant %uint 4294967295 + %1 = OpFunction %void None %25 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %20 = 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_uint Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %18 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpIMul %uint %14 %uint_4294967295 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %uint %6 + %19 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %19 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/cos.ptx b/ptx/src/test/spirv_run/cos.ptx new file mode 100644 index 0000000..7983f20 --- /dev/null +++ b/ptx/src/test/spirv_run/cos.ptx @@ -0,0 +1,21 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry cos(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .f32 temp;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.f32 temp, [in_addr];
+ cos.approx.f32 temp, temp;
+ st.f32 [out_addr], temp;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/cos.spvtxt b/ptx/src/test/spirv_run/cos.spvtxt new file mode 100644 index 0000000..6820142 --- /dev/null +++ b/ptx/src/test/spirv_run/cos.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "cos" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 cos %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/cvt_sat_s_u.spvtxt b/ptx/src/test/spirv_run/cvt_sat_s_u.spvtxt index 208c279..be321eb 100644 --- a/ptx/src/test/spirv_run/cvt_sat_s_u.spvtxt +++ b/ptx/src/test/spirv_run/cvt_sat_s_u.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %27 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %25 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "cvt_sat_s_u" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %30 = OpTypeFunction %void %ulong %ulong + %28 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint - %1 = OpFunction %void None %30 + %1 = OpFunction %void None %28 %9 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong - %25 = OpLabel + %23 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -27,25 +30,23 @@ %8 = OpVariable %_ptr_Function_uint Function OpStore %2 %9 OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %2 OpStore %4 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %5 %13 - %16 = OpLoad %ulong %4 - %23 = OpConvertUToPtr %_ptr_Generic_uint %16 - %15 = OpLoad %uint %23 - OpStore %6 %15 - %18 = OpLoad %uint %6 - %17 = OpSatConvertSToU %uint %18 - OpStore %7 %17 - %20 = OpLoad %uint %7 - %19 = OpBitcast %uint %20 - OpStore %8 %19 - %21 = OpLoad %ulong %5 - %22 = OpLoad %uint %8 - %24 = OpConvertUToPtr %_ptr_Generic_uint %21 - OpStore %24 %22 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %21 = OpConvertUToPtr %_ptr_Generic_uint %14 + %13 = OpLoad %uint %21 + OpStore %6 %13 + %16 = OpLoad %uint %6 + %15 = OpSatConvertSToU %uint %16 + OpStore %7 %15 + %18 = OpLoad %uint %7 + %17 = OpBitcast %uint %18 + OpStore %8 %17 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %8 + %22 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %22 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/cvta.spvtxt b/ptx/src/test/spirv_run/cvta.spvtxt index 84e7eac..cf6ff8b 100644 --- a/ptx/src/test/spirv_run/cvta.spvtxt +++ b/ptx/src/test/spirv_run/cvta.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %29 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %27 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "cvta" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %32 = OpTypeFunction %void %ulong %ulong + %30 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_Function_float = OpTypePointer Function %float %_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float - %1 = OpFunction %void None %32 + %1 = OpFunction %void None %30 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %27 = OpLabel + %25 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -25,29 +28,27 @@ %6 = OpVariable %_ptr_Function_float Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %20 = OpCopyObject %ulong %12 + %19 = OpCopyObject %ulong %20 + %11 = OpCopyObject %ulong %19 + OpStore %4 %11 + %14 = OpLoad %ulong %5 %22 = OpCopyObject %ulong %14 %21 = OpCopyObject %ulong %22 %13 = OpCopyObject %ulong %21 - OpStore %4 %13 - %16 = OpLoad %ulong %5 - %24 = OpCopyObject %ulong %16 - %23 = OpCopyObject %ulong %24 - %15 = OpCopyObject %ulong %23 - OpStore %5 %15 - %18 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %18 - %17 = OpLoad %float %25 - OpStore %6 %17 - %19 = OpLoad %ulong %5 - %20 = OpLoad %float %6 - %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %19 - OpStore %26 %20 + OpStore %5 %13 + %16 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %16 + %15 = OpLoad %float %23 + OpStore %6 %15 + %17 = OpLoad %ulong %5 + %18 = OpLoad %float %6 + %24 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %17 + OpStore %24 %18 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/div_approx.spvtxt b/ptx/src/test/spirv_run/div_approx.spvtxt index 40cc152..c62888c 100644 --- a/ptx/src/test/spirv_run/div_approx.spvtxt +++ b/ptx/src/test/spirv_run/div_approx.spvtxt @@ -1,65 +1,56 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 38 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%30 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "div_approx" -OpDecorate %1 FunctionDenormModeINTEL 32 Preserve -OpDecorate %18 FPFastMathMode AllowRecip -%31 = OpTypeVoid -%32 = OpTypeInt 64 0 -%33 = OpTypeFunction %31 %32 %32 -%34 = OpTypePointer Function %32 -%35 = OpTypeFloat 32 -%36 = OpTypePointer Function %35 -%37 = OpTypePointer Generic %35 -%23 = OpConstant %32 4 -%1 = OpFunction %31 None %33 -%8 = OpFunctionParameter %32 -%9 = OpFunctionParameter %32 -%28 = OpLabel -%2 = OpVariable %34 Function -%3 = OpVariable %34 Function -%4 = OpVariable %34 Function -%5 = OpVariable %34 Function -%6 = OpVariable %36 Function -%7 = OpVariable %36 Function -OpStore %2 %8 -OpStore %3 %9 -%11 = OpLoad %32 %2 -%10 = OpCopyObject %32 %11 -OpStore %4 %10 -%13 = OpLoad %32 %3 -%12 = OpCopyObject %32 %13 -OpStore %5 %12 -%15 = OpLoad %32 %4 -%25 = OpConvertUToPtr %37 %15 -%14 = OpLoad %35 %25 -OpStore %6 %14 -%17 = OpLoad %32 %4 -%24 = OpIAdd %32 %17 %23 -%26 = OpConvertUToPtr %37 %24 -%16 = OpLoad %35 %26 -OpStore %7 %16 -%19 = OpLoad %35 %6 -%20 = OpLoad %35 %7 -%18 = OpFDiv %35 %19 %20 -OpStore %6 %18 -%21 = OpLoad %32 %5 -%22 = OpLoad %35 %6 -%27 = OpConvertUToPtr %37 %21 -OpStore %27 %22 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %28 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "div_approx" + OpDecorate %16 FPFastMathMode AllowRecip + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %31 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %ulong_4 = OpConstant %ulong 4 + %1 = OpFunction %void None %31 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %26 = 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_float Function + %7 = OpVariable %_ptr_Function_float Function + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %4 %10 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_float %13 + %12 = OpLoad %float %23 + OpStore %6 %12 + %15 = OpLoad %ulong %4 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_float %22 + %14 = OpLoad %float %24 + OpStore %7 %14 + %17 = OpLoad %float %6 + %18 = OpLoad %float %7 + %16 = OpFDiv %float %17 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %float %6 + %25 = OpConvertUToPtr %_ptr_Generic_float %19 + OpStore %25 %20 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ex2.ptx b/ptx/src/test/spirv_run/ex2.ptx new file mode 100644 index 0000000..1edbcc6 --- /dev/null +++ b/ptx/src/test/spirv_run/ex2.ptx @@ -0,0 +1,21 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry ex2(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .f32 temp;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.f32 temp, [in_addr];
+ ex2.approx.f32 temp, temp;
+ st.f32 [out_addr], temp;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/ex2.spvtxt b/ptx/src/test/spirv_run/ex2.spvtxt new file mode 100644 index 0000000..e9be3e3 --- /dev/null +++ b/ptx/src/test/spirv_run/ex2.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "ex2" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 exp2 %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/extern_shared.spvtxt b/ptx/src/test/spirv_run/extern_shared.spvtxt index b184980..fca8ee7 100644 --- a/ptx/src/test/spirv_run/extern_shared.spvtxt +++ b/ptx/src/test/spirv_run/extern_shared.spvtxt @@ -7,7 +7,7 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %32 = OpExtInstImport "OpenCL.std" + %30 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "extern_shared" %1 %void = OpTypeVoid @@ -18,51 +18,49 @@ %ulong = OpTypeInt 64 0 %uchar = OpTypeInt 8 0 %_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar - %40 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar + %38 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar %_ptr_Function__ptr_Workgroup_uchar = OpTypePointer Function %_ptr_Workgroup_uchar %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %_ptr_Function__ptr_Workgroup_uint = OpTypePointer Function %_ptr_Workgroup_uint %_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong - %2 = OpFunction %void None %40 + %2 = OpFunction %void None %38 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %28 = OpFunctionParameter %_ptr_Workgroup_uchar - %41 = OpLabel - %29 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function + %26 = OpFunctionParameter %_ptr_Workgroup_uchar + %39 = OpLabel + %27 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function - OpStore %29 %28 - OpBranch %26 - %26 = OpLabel + OpStore %27 %26 + OpBranch %24 + %24 = OpLabel OpStore %3 %8 OpStore %4 %9 - %11 = OpLoad %ulong %3 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %3 OpStore %5 %10 - %13 = OpLoad %ulong %4 - %12 = OpCopyObject %ulong %13 - OpStore %6 %12 - %15 = OpLoad %ulong %5 - %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %15 - %14 = OpLoad %ulong %22 - OpStore %7 %14 - %30 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %29 - %16 = OpLoad %_ptr_Workgroup_uint %30 - %17 = OpLoad %ulong %7 - %23 = OpBitcast %_ptr_Workgroup_ulong %16 - OpStore %23 %17 - %31 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %29 - %19 = OpLoad %_ptr_Workgroup_uint %31 - %24 = OpBitcast %_ptr_Workgroup_ulong %19 - %18 = OpLoad %ulong %24 - OpStore %7 %18 - %20 = OpLoad %ulong %6 - %21 = OpLoad %ulong %7 - %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %20 - OpStore %25 %21 + %11 = OpLoad %ulong %4 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %20 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %13 + %12 = OpLoad %ulong %20 + OpStore %7 %12 + %28 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %27 + %14 = OpLoad %_ptr_Workgroup_uint %28 + %15 = OpLoad %ulong %7 + %21 = OpBitcast %_ptr_Workgroup_ulong %14 + OpStore %21 %15 + %29 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %27 + %17 = OpLoad %_ptr_Workgroup_uint %29 + %22 = OpBitcast %_ptr_Workgroup_ulong %17 + %16 = OpLoad %ulong %22 + OpStore %7 %16 + %18 = OpLoad %ulong %6 + %19 = OpLoad %ulong %7 + %23 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %18 + OpStore %23 %19 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/fma.spvtxt b/ptx/src/test/spirv_run/fma.spvtxt index 4a90d09..9716198 100644 --- a/ptx/src/test/spirv_run/fma.spvtxt +++ b/ptx/src/test/spirv_run/fma.spvtxt @@ -1,72 +1,63 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 45 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%37 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "fma" -; OpDecorate %1 FunctionDenormModeINTEL 32 Preserve -%38 = OpTypeVoid -%39 = OpTypeInt 64 0 -%40 = OpTypeFunction %38 %39 %39 -%41 = OpTypePointer Function %39 -%42 = OpTypeFloat 32 -%43 = OpTypePointer Function %42 -%44 = OpTypePointer Generic %42 -%27 = OpConstant %39 4 -%29 = OpConstant %39 8 -%1 = OpFunction %38 None %40 -%9 = OpFunctionParameter %39 -%10 = OpFunctionParameter %39 -%35 = OpLabel -%2 = OpVariable %41 Function -%3 = OpVariable %41 Function -%4 = OpVariable %41 Function -%5 = OpVariable %41 Function -%6 = OpVariable %43 Function -%7 = OpVariable %43 Function -%8 = OpVariable %43 Function -OpStore %2 %9 -OpStore %3 %10 -%12 = OpLoad %39 %2 -%11 = OpCopyObject %39 %12 -OpStore %4 %11 -%14 = OpLoad %39 %3 -%13 = OpCopyObject %39 %14 -OpStore %5 %13 -%16 = OpLoad %39 %4 -%31 = OpConvertUToPtr %44 %16 -%15 = OpLoad %42 %31 -OpStore %6 %15 -%18 = OpLoad %39 %4 -%28 = OpIAdd %39 %18 %27 -%32 = OpConvertUToPtr %44 %28 -%17 = OpLoad %42 %32 -OpStore %7 %17 -%20 = OpLoad %39 %4 -%30 = OpIAdd %39 %20 %29 -%33 = OpConvertUToPtr %44 %30 -%19 = OpLoad %42 %33 -OpStore %8 %19 -%22 = OpLoad %42 %6 -%23 = OpLoad %42 %7 -%24 = OpLoad %42 %8 -%21 = OpExtInst %42 %37 mad %22 %23 %24 -OpStore %6 %21 -%25 = OpLoad %39 %5 -%26 = OpLoad %42 %6 -%34 = OpConvertUToPtr %44 %25 -OpStore %34 %26 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %35 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "fma" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %38 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %ulong_4 = OpConstant %ulong 4 + %ulong_8 = OpConstant %ulong 8 + %1 = OpFunction %void None %38 + %9 = OpFunctionParameter %ulong + %10 = OpFunctionParameter %ulong + %33 = 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_float Function + %7 = OpVariable %_ptr_Function_float Function + %8 = OpVariable %_ptr_Function_float Function + OpStore %2 %9 + OpStore %3 %10 + %11 = OpLoad %ulong %2 + OpStore %4 %11 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %29 = OpConvertUToPtr %_ptr_Generic_float %14 + %13 = OpLoad %float %29 + OpStore %6 %13 + %16 = OpLoad %ulong %4 + %26 = OpIAdd %ulong %16 %ulong_4 + %30 = OpConvertUToPtr %_ptr_Generic_float %26 + %15 = OpLoad %float %30 + OpStore %7 %15 + %18 = OpLoad %ulong %4 + %28 = OpIAdd %ulong %18 %ulong_8 + %31 = OpConvertUToPtr %_ptr_Generic_float %28 + %17 = OpLoad %float %31 + OpStore %8 %17 + %20 = OpLoad %float %6 + %21 = OpLoad %float %7 + %22 = OpLoad %float %8 + %19 = OpExtInst %float %35 mad %20 %21 %22 + OpStore %6 %19 + %23 = OpLoad %ulong %5 + %24 = OpLoad %float %6 + %32 = OpConvertUToPtr %_ptr_Generic_float %23 + OpStore %32 %24 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/global_array.spvtxt b/ptx/src/test/spirv_run/global_array.spvtxt index a4ed91d..25874ac 100644 --- a/ptx/src/test/spirv_run/global_array.spvtxt +++ b/ptx/src/test/spirv_run/global_array.spvtxt @@ -7,28 +7,28 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %22 = OpExtInstImport "OpenCL.std" + %21 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "global_array" %1 %void = OpTypeVoid %uint = OpTypeInt 32 0 %uint_4 = OpConstant %uint 4 %_arr_uint_uint_4 = OpTypeArray %uint %uint_4 -%_ptr_CrossWorkgroup__arr_uint_uint_4 = OpTypePointer CrossWorkgroup %_arr_uint_uint_4 - %uint_4_0 = OpConstant %uint 4 %uint_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 - %31 = OpConstantComposite %_arr_uint_uint_4 %uint_1 %uint_0 %uint_0 %uint_0 - %1 = OpVariable %_ptr_CrossWorkgroup__arr_uint_uint_4 CrossWorkgroup %31 + %28 = OpConstantComposite %_arr_uint_uint_4 %uint_1 %uint_0 %uint_0 %uint_0 + %uint_4_0 = OpConstant %uint 4 +%_ptr_CrossWorkgroup__arr_uint_uint_4 = OpTypePointer CrossWorkgroup %_arr_uint_uint_4 + %1 = OpVariable %_ptr_CrossWorkgroup__arr_uint_uint_4 CrossWorkgroup %28 %ulong = OpTypeInt 64 0 - %33 = OpTypeFunction %void %ulong %ulong + %32 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint - %2 = OpFunction %void None %33 + %2 = OpFunction %void None %32 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %20 = OpLabel + %19 = OpLabel %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function %5 = OpVariable %_ptr_Function_ulong Function @@ -36,19 +36,18 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %3 %8 OpStore %4 %9 - %17 = OpConvertPtrToU %ulong %1 - %10 = OpCopyObject %ulong %17 + %16 = OpConvertPtrToU %ulong %1 + %10 = OpCopyObject %ulong %16 OpStore %5 %10 - %12 = OpLoad %ulong %4 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %4 OpStore %6 %11 - %14 = OpLoad %ulong %5 + %13 = OpLoad %ulong %5 + %17 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %13 + %12 = OpLoad %uint %17 + OpStore %7 %12 + %14 = OpLoad %ulong %6 + %15 = OpLoad %uint %7 %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %14 - %13 = OpLoad %uint %18 - OpStore %7 %13 - %15 = OpLoad %ulong %6 - %16 = OpLoad %uint %7 - %19 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %15 - OpStore %19 %16 + OpStore %18 %15 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/implicit_param.spvtxt b/ptx/src/test/spirv_run/implicit_param.spvtxt index c30788c..a78e53f 100644 --- a/ptx/src/test/spirv_run/implicit_param.spvtxt +++ b/ptx/src/test/spirv_run/implicit_param.spvtxt @@ -2,25 +2,27 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %28 = OpExtInstImport "OpenCL.std" + %24 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "implicit_param" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %31 = OpTypeFunction %void %ulong %ulong + %27 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_Function_float = OpTypePointer Function %float %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float - %1 = OpFunction %void None %31 + %1 = OpFunction %void None %27 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %26 = OpLabel + %22 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -29,27 +31,23 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %15 - %14 = OpLoad %float %22 - OpStore %6 %14 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %13 + %12 = OpLoad %float %18 + OpStore %6 %12 + %14 = OpLoad %float %6 + %19 = OpBitcast %_ptr_Function_float %7 + OpStore %19 %14 + %20 = OpBitcast %_ptr_Function_float %7 + %15 = OpLoad %float %20 + OpStore %6 %15 + %16 = OpLoad %ulong %5 %17 = OpLoad %float %6 - %23 = OpCopyObject %float %17 - %16 = OpBitcast %uint %23 - OpStore %7 %16 - %19 = OpLoad %uint %7 - %24 = OpBitcast %float %19 - %18 = OpCopyObject %float %24 - OpStore %6 %18 - %20 = OpLoad %ulong %5 - %21 = OpLoad %float %6 - %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %20 - OpStore %25 %21 + %21 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %16 + OpStore %21 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ld_st.spvtxt b/ptx/src/test/spirv_run/ld_st.spvtxt index d36db57..8e3f98d 100644 --- a/ptx/src/test/spirv_run/ld_st.spvtxt +++ b/ptx/src/test/spirv_run/ld_st.spvtxt @@ -2,20 +2,23 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %21 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %19 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "ld_st" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %24 = OpTypeFunction %void %ulong %ulong + %22 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %24 + %1 = OpFunction %void None %22 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %19 = OpLabel + %17 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -23,19 +26,17 @@ %6 = OpVariable %_ptr_Function_ulong Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 - %17 = OpConvertUToPtr %_ptr_Generic_ulong %14 - %13 = OpLoad %ulong %17 - OpStore %6 %13 - %15 = OpLoad %ulong %5 - %16 = OpLoad %ulong %6 - %18 = OpConvertUToPtr %_ptr_Generic_ulong %15 - OpStore %18 %16 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %15 = OpConvertUToPtr %_ptr_Generic_ulong %12 + %11 = OpLoad %ulong %15 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %14 = OpLoad %ulong %6 + %16 = OpConvertUToPtr %_ptr_Generic_ulong %13 + OpStore %16 %14 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ld_st_implicit.spvtxt b/ptx/src/test/spirv_run/ld_st_implicit.spvtxt index d4d9499..35f715b 100644 --- a/ptx/src/test/spirv_run/ld_st_implicit.spvtxt +++ b/ptx/src/test/spirv_run/ld_st_implicit.spvtxt @@ -7,20 +7,20 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %23 = OpExtInstImport "OpenCL.std" + %21 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "ld_st_implicit" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %26 = OpTypeFunction %void %ulong %ulong + %24 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float %uint = OpTypeInt 32 0 - %1 = OpFunction %void None %26 + %1 = OpFunction %void None %24 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %21 = OpLabel + %19 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -28,24 +28,22 @@ %6 = OpVariable %_ptr_Function_ulong Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 - %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %14 - %17 = OpLoad %float %18 - %31 = OpBitcast %uint %17 - %13 = OpUConvert %ulong %31 - OpStore %6 %13 - %15 = OpLoad %ulong %5 - %16 = OpLoad %ulong %6 - %19 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %15 - %32 = OpBitcast %ulong %16 - %33 = OpUConvert %uint %32 - %20 = OpBitcast %float %33 - OpStore %19 %20 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %16 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %12 + %15 = OpLoad %float %16 + %29 = OpBitcast %uint %15 + %11 = OpUConvert %ulong %29 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %14 = OpLoad %ulong %6 + %17 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %13 + %30 = OpBitcast %ulong %14 + %31 = OpUConvert %uint %30 + %18 = OpBitcast %float %31 + OpStore %17 %18 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ld_st_offset.spvtxt b/ptx/src/test/spirv_run/ld_st_offset.spvtxt index 208b53b..963d88a 100644 --- a/ptx/src/test/spirv_run/ld_st_offset.spvtxt +++ b/ptx/src/test/spirv_run/ld_st_offset.spvtxt @@ -2,24 +2,27 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %32 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %30 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "ld_st_offset" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %35 = OpTypeFunction %void %ulong %ulong + %33 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint %ulong_4 = OpConstant %ulong 4 %ulong_4_0 = OpConstant %ulong 4 - %1 = OpFunction %void None %35 + %1 = OpFunction %void None %33 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %30 = OpLabel + %28 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -28,29 +31,27 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %24 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %24 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %26 = OpConvertUToPtr %_ptr_Generic_uint %15 - %14 = OpLoad %uint %26 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %23 = OpIAdd %ulong %17 %ulong_4 - %27 = OpConvertUToPtr %_ptr_Generic_uint %23 - %16 = OpLoad %uint %27 - OpStore %7 %16 + %21 = OpIAdd %ulong %15 %ulong_4 + %25 = OpConvertUToPtr %_ptr_Generic_uint %21 + %14 = OpLoad %uint %25 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %uint %7 + %26 = OpConvertUToPtr %_ptr_Generic_uint %16 + OpStore %26 %17 %18 = OpLoad %ulong %5 - %19 = OpLoad %uint %7 - %28 = OpConvertUToPtr %_ptr_Generic_uint %18 - OpStore %28 %19 - %20 = OpLoad %ulong %5 - %21 = OpLoad %uint %6 - %25 = OpIAdd %ulong %20 %ulong_4_0 - %29 = OpConvertUToPtr %_ptr_Generic_uint %25 - OpStore %29 %21 + %19 = OpLoad %uint %6 + %23 = OpIAdd %ulong %18 %ulong_4_0 + %27 = OpConvertUToPtr %_ptr_Generic_uint %23 + OpStore %27 %19 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/lg2.ptx b/ptx/src/test/spirv_run/lg2.ptx new file mode 100644 index 0000000..c571955 --- /dev/null +++ b/ptx/src/test/spirv_run/lg2.ptx @@ -0,0 +1,21 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry lg2(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .f32 temp;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.f32 temp, [in_addr];
+ lg2.approx.f32 temp, temp;
+ st.f32 [out_addr], temp;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/lg2.spvtxt b/ptx/src/test/spirv_run/lg2.spvtxt new file mode 100644 index 0000000..d30fe8a --- /dev/null +++ b/ptx/src/test/spirv_run/lg2.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "lg2" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 log2 %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/local_align.spvtxt b/ptx/src/test/spirv_run/local_align.spvtxt index 2482a75..915ac6f 100644 --- a/ptx/src/test/spirv_run/local_align.spvtxt +++ b/ptx/src/test/spirv_run/local_align.spvtxt @@ -2,26 +2,29 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %22 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %20 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "local_align" OpDecorate %4 Alignment 8 %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %25 = OpTypeFunction %void %ulong %ulong + %23 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong - %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 %uint_8 = OpConstant %uint 8 %_arr_uchar_uint_8 = OpTypeArray %uchar %uint_8 %_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8 %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %25 + %1 = OpFunction %void None %23 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %20 = OpLabel + %18 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function__arr_uchar_uint_8 Function @@ -30,19 +33,17 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %5 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %6 %12 - %15 = OpLoad %ulong %5 - %18 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %18 - OpStore %7 %14 - %16 = OpLoad %ulong %6 - %17 = OpLoad %ulong %7 - %19 = OpConvertUToPtr %_ptr_Generic_ulong %16 - OpStore %19 %17 + %11 = OpLoad %ulong %3 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %16 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %16 + OpStore %7 %12 + %14 = OpLoad %ulong %6 + %15 = OpLoad %ulong %7 + %17 = OpConvertUToPtr %_ptr_Generic_ulong %14 + OpStore %17 %15 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mad_s32.spvtxt b/ptx/src/test/spirv_run/mad_s32.spvtxt index 3a7153d..9150089 100644 --- a/ptx/src/test/spirv_run/mad_s32.spvtxt +++ b/ptx/src/test/spirv_run/mad_s32.spvtxt @@ -2,15 +2,17 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %48 = OpExtInstImport "OpenCL.std" + %46 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mad_s32" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %51 = OpTypeFunction %void %ulong %ulong + %49 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint @@ -19,10 +21,10 @@ %ulong_8 = OpConstant %ulong 8 %ulong_4_0 = OpConstant %ulong 4 %ulong_8_0 = OpConstant %ulong 8 - %1 = OpFunction %void None %51 + %1 = OpFunction %void None %49 %10 = OpFunctionParameter %ulong %11 = OpFunctionParameter %ulong - %46 = OpLabel + %44 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -33,45 +35,43 @@ %9 = OpVariable %_ptr_Function_uint Function OpStore %2 %10 OpStore %3 %11 - %13 = OpLoad %ulong %2 - %12 = OpCopyObject %ulong %13 + %12 = OpLoad %ulong %2 OpStore %4 %12 - %15 = OpLoad %ulong %3 - %14 = OpCopyObject %ulong %15 - OpStore %5 %14 + %13 = OpLoad %ulong %3 + OpStore %5 %13 + %15 = OpLoad %ulong %4 + %38 = OpConvertUToPtr %_ptr_Generic_uint %15 + %14 = OpLoad %uint %38 + OpStore %7 %14 %17 = OpLoad %ulong %4 - %40 = OpConvertUToPtr %_ptr_Generic_uint %17 - %16 = OpLoad %uint %40 - OpStore %7 %16 + %31 = OpIAdd %ulong %17 %ulong_4 + %39 = OpConvertUToPtr %_ptr_Generic_uint %31 + %16 = OpLoad %uint %39 + OpStore %8 %16 %19 = OpLoad %ulong %4 - %33 = OpIAdd %ulong %19 %ulong_4 - %41 = OpConvertUToPtr %_ptr_Generic_uint %33 - %18 = OpLoad %uint %41 - OpStore %8 %18 - %21 = OpLoad %ulong %4 - %35 = OpIAdd %ulong %21 %ulong_8 - %42 = OpConvertUToPtr %_ptr_Generic_uint %35 - %20 = OpLoad %uint %42 - OpStore %9 %20 - %23 = OpLoad %uint %7 - %24 = OpLoad %uint %8 - %25 = OpLoad %uint %9 - %56 = OpIMul %uint %23 %24 - %22 = OpIAdd %uint %25 %56 - OpStore %6 %22 + %33 = OpIAdd %ulong %19 %ulong_8 + %40 = OpConvertUToPtr %_ptr_Generic_uint %33 + %18 = OpLoad %uint %40 + OpStore %9 %18 + %21 = OpLoad %uint %7 + %22 = OpLoad %uint %8 + %23 = OpLoad %uint %9 + %54 = OpIMul %uint %21 %22 + %20 = OpIAdd %uint %23 %54 + OpStore %6 %20 + %24 = OpLoad %ulong %5 + %25 = OpLoad %uint %6 + %41 = OpConvertUToPtr %_ptr_Generic_uint %24 + OpStore %41 %25 %26 = OpLoad %ulong %5 %27 = OpLoad %uint %6 - %43 = OpConvertUToPtr %_ptr_Generic_uint %26 - OpStore %43 %27 + %35 = OpIAdd %ulong %26 %ulong_4_0 + %42 = OpConvertUToPtr %_ptr_Generic_uint %35 + OpStore %42 %27 %28 = OpLoad %ulong %5 %29 = OpLoad %uint %6 - %37 = OpIAdd %ulong %28 %ulong_4_0 - %44 = OpConvertUToPtr %_ptr_Generic_uint %37 - OpStore %44 %29 - %30 = OpLoad %ulong %5 - %31 = OpLoad %uint %6 - %39 = OpIAdd %ulong %30 %ulong_8_0 - %45 = OpConvertUToPtr %_ptr_Generic_uint %39 - OpStore %45 %31 + %37 = OpIAdd %ulong %28 %ulong_8_0 + %43 = OpConvertUToPtr %_ptr_Generic_uint %37 + OpStore %43 %29 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/max.spvtxt b/ptx/src/test/spirv_run/max.spvtxt index cab9a9a..05eb705 100644 --- a/ptx/src/test/spirv_run/max.spvtxt +++ b/ptx/src/test/spirv_run/max.spvtxt @@ -7,21 +7,21 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %30 = OpExtInstImport "OpenCL.std" + %28 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "max" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %33 = OpTypeFunction %void %ulong %ulong + %31 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint %ulong_4 = OpConstant %ulong 4 - %1 = OpFunction %void None %33 + %1 = OpFunction %void None %31 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %28 = OpLabel + %26 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -30,28 +30,26 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %23 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_uint %15 - %14 = OpLoad %uint %25 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %24 = OpIAdd %ulong %17 %ulong_4 - %26 = OpConvertUToPtr %_ptr_Generic_uint %24 - %16 = OpLoad %uint %26 - OpStore %7 %16 - %19 = OpLoad %uint %6 - %20 = OpLoad %uint %7 - %18 = OpExtInst %uint %30 s_max %19 %20 - OpStore %6 %18 - %21 = OpLoad %ulong %5 - %22 = OpLoad %uint %6 - %27 = OpConvertUToPtr %_ptr_Generic_uint %21 - OpStore %27 %22 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_uint %22 + %14 = OpLoad %uint %24 + OpStore %7 %14 + %17 = OpLoad %uint %6 + %18 = OpLoad %uint %7 + %16 = OpExtInst %uint %28 s_max %17 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %6 + %25 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %25 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/min.spvtxt b/ptx/src/test/spirv_run/min.spvtxt index 119cd15..d0d2b9a 100644 --- a/ptx/src/test/spirv_run/min.spvtxt +++ b/ptx/src/test/spirv_run/min.spvtxt @@ -7,21 +7,21 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %30 = OpExtInstImport "OpenCL.std" + %28 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "min" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %33 = OpTypeFunction %void %ulong %ulong + %31 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint %ulong_4 = OpConstant %ulong 4 - %1 = OpFunction %void None %33 + %1 = OpFunction %void None %31 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %28 = OpLabel + %26 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -30,28 +30,26 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_uint %13 + %12 = OpLoad %uint %23 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_uint %15 - %14 = OpLoad %uint %25 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %24 = OpIAdd %ulong %17 %ulong_4 - %26 = OpConvertUToPtr %_ptr_Generic_uint %24 - %16 = OpLoad %uint %26 - OpStore %7 %16 - %19 = OpLoad %uint %6 - %20 = OpLoad %uint %7 - %18 = OpExtInst %uint %30 s_min %19 %20 - OpStore %6 %18 - %21 = OpLoad %ulong %5 - %22 = OpLoad %uint %6 - %27 = OpConvertUToPtr %_ptr_Generic_uint %21 - OpStore %27 %22 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_uint %22 + %14 = OpLoad %uint %24 + OpStore %7 %14 + %17 = OpLoad %uint %6 + %18 = OpLoad %uint %7 + %16 = OpExtInst %uint %28 s_min %17 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %uint %6 + %25 = OpConvertUToPtr %_ptr_Generic_uint %19 + OpStore %25 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 7ba3c4d..3fa82ba 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -60,8 +60,7 @@ test_ptx!(call, [1u64], [2u64]); test_ptx!(vector, [1u32, 2u32], [3u32, 3u32]);
test_ptx!(ld_st_offset, [1u32, 2u32], [2u32, 1u32]);
test_ptx!(ntid, [3u32], [4u32]);
-// TODO: enable test below
-// test_ptx!(reg_local, [12u64], [13u64]);
+test_ptx!(reg_local, [12u64], [13u64]);
test_ptx!(mov_address, [0xDEADu64], [0u64]);
test_ptx!(b64tof64, [111u64], [111u64]);
test_ptx!(implicit_param, [34u32], [34u32]);
@@ -105,6 +104,10 @@ test_ptx!(div_approx, [1f32, 2f32], [0.5f32]); test_ptx!(sqrt, [0.25f32], [0.5f32]);
test_ptx!(rsqrt, [0.25f64], [2f64]);
test_ptx!(neg, [181i32], [-181i32]);
+test_ptx!(sin, [std::f32::consts::PI/2f32], [1f32]);
+test_ptx!(cos, [std::f32::consts::PI], [-1f32]);
+test_ptx!(lg2, [512f32], [9f32]);
+test_ptx!(ex2, [10f32], [1024f32]);
struct DisplayError<T: Debug> {
err: T,
diff --git a/ptx/src/test/spirv_run/mov.spvtxt b/ptx/src/test/spirv_run/mov.spvtxt index d8a5029..15118aa 100644 --- a/ptx/src/test/spirv_run/mov.spvtxt +++ b/ptx/src/test/spirv_run/mov.spvtxt @@ -2,20 +2,23 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %24 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %22 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mov" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %27 = OpTypeFunction %void %ulong %ulong + %25 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %27 + %1 = OpFunction %void None %25 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %22 = OpLabel + %20 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -24,22 +27,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %20 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %20 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpCopyObject %ulong %17 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %21 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %18 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpCopyObject %ulong %15 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %19 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mul_hi.spvtxt b/ptx/src/test/spirv_run/mul_hi.spvtxt index bea23a9..8449183 100644 --- a/ptx/src/test/spirv_run/mul_hi.spvtxt +++ b/ptx/src/test/spirv_run/mul_hi.spvtxt @@ -2,21 +2,24 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %25 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %23 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mul_hi" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %28 = OpTypeFunction %void %ulong %ulong + %26 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %28 + %1 = OpFunction %void None %26 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %23 = OpLabel + %21 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -25,22 +28,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpExtInst %ulong %25 u_mul_hi %17 %ulong_2 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %22 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpExtInst %ulong %23 u_mul_hi %15 %ulong_2 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %20 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mul_lo.spvtxt b/ptx/src/test/spirv_run/mul_lo.spvtxt index e114374..d4b2566 100644 --- a/ptx/src/test/spirv_run/mul_lo.spvtxt +++ b/ptx/src/test/spirv_run/mul_lo.spvtxt @@ -2,21 +2,24 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %25 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %23 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mul_lo" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %28 = OpTypeFunction %void %ulong %ulong + %26 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %28 + %1 = OpFunction %void None %26 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %23 = OpLabel + %21 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -25,22 +28,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpIMul %ulong %17 %ulong_2 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %22 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpIMul %ulong %15 %ulong_2 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %20 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mul_non_ftz.spvtxt b/ptx/src/test/spirv_run/mul_non_ftz.spvtxt index 78153aa..cb20943 100644 --- a/ptx/src/test/spirv_run/mul_non_ftz.spvtxt +++ b/ptx/src/test/spirv_run/mul_non_ftz.spvtxt @@ -7,25 +7,21 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - OpCapability DenormFlushToZero - OpCapability DenormPreserve - OpExtension "SPV_KHR_float_controls" - %30 = OpExtInstImport "OpenCL.std" + %28 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mul_non_ftz" - OpExecutionMode %1 DenormPreserve 32 %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %33 = OpTypeFunction %void %ulong %ulong + %31 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_Function_float = OpTypePointer Function %float %_ptr_Generic_float = OpTypePointer Generic %float %ulong_4 = OpConstant %ulong 4 - %1 = OpFunction %void None %33 + %1 = OpFunction %void None %31 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %28 = OpLabel + %26 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -34,28 +30,26 @@ %7 = OpVariable %_ptr_Function_float Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_float %13 + %12 = OpLoad %float %23 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_float %15 - %14 = OpLoad %float %25 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %24 = OpIAdd %ulong %17 %ulong_4 - %26 = OpConvertUToPtr %_ptr_Generic_float %24 - %16 = OpLoad %float %26 - OpStore %7 %16 - %19 = OpLoad %float %6 - %20 = OpLoad %float %7 - %18 = OpFMul %float %19 %20 - OpStore %6 %18 - %21 = OpLoad %ulong %5 - %22 = OpLoad %float %6 - %27 = OpConvertUToPtr %_ptr_Generic_float %21 - OpStore %27 %22 + %22 = OpIAdd %ulong %15 %ulong_4 + %24 = OpConvertUToPtr %_ptr_Generic_float %22 + %14 = OpLoad %float %24 + OpStore %7 %14 + %17 = OpLoad %float %6 + %18 = OpLoad %float %7 + %16 = OpFMul %float %17 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %float %6 + %25 = OpConvertUToPtr %_ptr_Generic_float %19 + OpStore %25 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mul_wide.spvtxt b/ptx/src/test/spirv_run/mul_wide.spvtxt index 8ac0459..632fa3b 100644 --- a/ptx/src/test/spirv_run/mul_wide.spvtxt +++ b/ptx/src/test/spirv_run/mul_wide.spvtxt @@ -7,24 +7,24 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %32 = OpExtInstImport "OpenCL.std" + %30 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "mul_wide" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %35 = OpTypeFunction %void %ulong %ulong + %33 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint %ulong_4 = OpConstant %ulong 4 - %_struct_40 = OpTypeStruct %uint %uint + %_struct_38 = OpTypeStruct %uint %uint %v2uint = OpTypeVector %uint 2 %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %35 + %1 = OpFunction %void None %33 %9 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong - %30 = OpLabel + %28 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -34,33 +34,31 @@ %8 = OpVariable %_ptr_Function_ulong Function OpStore %2 %9 OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %2 OpStore %4 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %5 %13 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %24 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %14 + %13 = OpLoad %uint %24 + OpStore %6 %13 %16 = OpLoad %ulong %4 - %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %16 - %15 = OpLoad %uint %26 - OpStore %6 %15 - %18 = OpLoad %ulong %4 - %25 = OpIAdd %ulong %18 %ulong_4 - %27 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %25 - %17 = OpLoad %uint %27 - OpStore %7 %17 - %20 = OpLoad %uint %6 - %21 = OpLoad %uint %7 - %41 = OpSMulExtended %_struct_40 %20 %21 - %42 = OpCompositeExtract %uint %41 0 - %43 = OpCompositeExtract %uint %41 1 - %45 = OpCompositeConstruct %v2uint %42 %43 - %19 = OpBitcast %ulong %45 - OpStore %8 %19 - %22 = OpLoad %ulong %5 - %23 = OpLoad %ulong %8 - %28 = OpConvertUToPtr %_ptr_Generic_ulong %22 - %29 = OpCopyObject %ulong %23 - OpStore %28 %29 + %23 = OpIAdd %ulong %16 %ulong_4 + %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %23 + %15 = OpLoad %uint %25 + OpStore %7 %15 + %18 = OpLoad %uint %6 + %19 = OpLoad %uint %7 + %39 = OpSMulExtended %_struct_38 %18 %19 + %40 = OpCompositeExtract %uint %39 0 + %41 = OpCompositeExtract %uint %39 1 + %43 = OpCompositeConstruct %v2uint %40 %41 + %17 = OpBitcast %ulong %43 + OpStore %8 %17 + %20 = OpLoad %ulong %5 + %21 = OpLoad %ulong %8 + %26 = OpConvertUToPtr %_ptr_Generic_ulong %20 + %27 = OpCopyObject %ulong %21 + OpStore %26 %27 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/neg.spvtxt b/ptx/src/test/spirv_run/neg.spvtxt index b358858..50726aa 100644 --- a/ptx/src/test/spirv_run/neg.spvtxt +++ b/ptx/src/test/spirv_run/neg.spvtxt @@ -2,46 +2,46 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %26 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL - OpEntryPoint Kernel %1 "not" + OpEntryPoint Kernel %1 "neg" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %29 = OpTypeFunction %void %ulong %ulong + %24 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong -%_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %29 + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %9 = OpFunctionParameter %ulong - %24 = OpLabel + %19 = 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_ulong Function - %7 = OpVariable %_ptr_Function_ulong Function - OpStore %2 %8 - OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 - OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %20 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %20 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %22 = OpCopyObject %ulong %17 - %21 = OpNot %ulong %22 - %16 = OpCopyObject %ulong %21 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %23 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %23 %19 + %6 = OpVariable %_ptr_Function_uint Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %17 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpSNegate %uint %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %uint %6 + %18 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %18 %16 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/not.spvtxt b/ptx/src/test/spirv_run/not.spvtxt index b358858..d6bc389 100644 --- a/ptx/src/test/spirv_run/not.spvtxt +++ b/ptx/src/test/spirv_run/not.spvtxt @@ -2,20 +2,23 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %26 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %24 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "not" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %29 = OpTypeFunction %void %ulong %ulong + %27 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong - %1 = OpFunction %void None %29 + %1 = OpFunction %void None %27 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %24 = OpLabel + %22 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -24,24 +27,22 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %20 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %20 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %22 = OpCopyObject %ulong %17 - %21 = OpNot %ulong %22 - %16 = OpCopyObject %ulong %21 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %23 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %23 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %18 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %20 = OpCopyObject %ulong %15 + %19 = OpNot %ulong %20 + %14 = OpCopyObject %ulong %19 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %21 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %21 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/ntid.spvtxt b/ptx/src/test/spirv_run/ntid.spvtxt index be16d2e..d1a3caa 100644 --- a/ptx/src/test/spirv_run/ntid.spvtxt +++ b/ptx/src/test/spirv_run/ntid.spvtxt @@ -2,10 +2,12 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %29 = OpExtInstImport "OpenCL.std" + %27 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "ntid" %gl_WorkGroupSize OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize @@ -15,14 +17,14 @@ %_ptr_UniformConstant_v4uint = OpTypePointer UniformConstant %v4uint %gl_WorkGroupSize = OpVariable %_ptr_UniformConstant_v4uint UniformConstant %ulong = OpTypeInt 64 0 - %35 = OpTypeFunction %void %ulong %ulong + %33 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint - %1 = OpFunction %void None %35 + %1 = OpFunction %void None %33 %9 = OpFunctionParameter %ulong %10 = OpFunctionParameter %ulong - %27 = OpLabel + %25 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -31,27 +33,25 @@ %7 = OpVariable %_ptr_Function_uint Function OpStore %2 %9 OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 + %11 = OpLoad %ulong %2 OpStore %4 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %5 %13 - %16 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_uint %16 - %15 = OpLoad %uint %25 - OpStore %6 %15 - %18 = OpLoad %v4uint %gl_WorkGroupSize - %24 = OpCompositeExtract %uint %18 0 - %17 = OpCopyObject %uint %24 - OpStore %7 %17 - %20 = OpLoad %uint %6 - %21 = OpLoad %uint %7 - %19 = OpIAdd %uint %20 %21 - OpStore %6 %19 - %22 = OpLoad %ulong %5 - %23 = OpLoad %uint %6 - %26 = OpConvertUToPtr %_ptr_Generic_uint %22 - OpStore %26 %23 + %12 = OpLoad %ulong %3 + OpStore %5 %12 + %14 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_uint %14 + %13 = OpLoad %uint %23 + OpStore %6 %13 + %16 = OpLoad %v4uint %gl_WorkGroupSize + %22 = OpCompositeExtract %uint %16 0 + %15 = OpCopyObject %uint %22 + OpStore %7 %15 + %18 = OpLoad %uint %6 + %19 = OpLoad %uint %7 + %17 = OpIAdd %uint %18 %19 + OpStore %6 %17 + %20 = OpLoad %ulong %5 + %21 = OpLoad %uint %6 + %24 = OpConvertUToPtr %_ptr_Generic_uint %20 + OpStore %24 %21 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/or.spvtxt b/ptx/src/test/spirv_run/or.spvtxt index fbf80c5..312b1b3 100644 --- a/ptx/src/test/spirv_run/or.spvtxt +++ b/ptx/src/test/spirv_run/or.spvtxt @@ -7,19 +7,19 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %33 = OpExtInstImport "OpenCL.std" + %31 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "or" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %36 = OpTypeFunction %void %ulong %ulong + %34 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_8 = OpConstant %ulong 8 - %1 = OpFunction %void None %36 + %1 = OpFunction %void None %34 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %31 = OpLabel + %29 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -28,31 +28,29 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %23 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %23 + OpStore %6 %12 %15 = OpLoad %ulong %4 - %25 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %25 - OpStore %6 %14 - %17 = OpLoad %ulong %4 - %24 = OpIAdd %ulong %17 %ulong_8 - %26 = OpConvertUToPtr %_ptr_Generic_ulong %24 - %16 = OpLoad %ulong %26 - OpStore %7 %16 - %19 = OpLoad %ulong %6 - %20 = OpLoad %ulong %7 - %28 = OpCopyObject %ulong %19 - %29 = OpCopyObject %ulong %20 - %27 = OpBitwiseOr %ulong %28 %29 - %18 = OpCopyObject %ulong %27 - OpStore %6 %18 - %21 = OpLoad %ulong %5 - %22 = OpLoad %ulong %6 - %30 = OpConvertUToPtr %_ptr_Generic_ulong %21 - OpStore %30 %22 + %22 = OpIAdd %ulong %15 %ulong_8 + %24 = OpConvertUToPtr %_ptr_Generic_ulong %22 + %14 = OpLoad %ulong %24 + OpStore %7 %14 + %17 = OpLoad %ulong %6 + %18 = OpLoad %ulong %7 + %26 = OpCopyObject %ulong %17 + %27 = OpCopyObject %ulong %18 + %25 = OpBitwiseOr %ulong %26 %27 + %16 = OpCopyObject %ulong %25 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %ulong %6 + %28 = OpConvertUToPtr %_ptr_Generic_ulong %19 + OpStore %28 %20 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/pred_not.spvtxt b/ptx/src/test/spirv_run/pred_not.spvtxt index 410b1e4..178c98f 100644 --- a/ptx/src/test/spirv_run/pred_not.spvtxt +++ b/ptx/src/test/spirv_run/pred_not.spvtxt @@ -2,15 +2,17 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 OpCapability Float64 - %44 = OpExtInstImport "OpenCL.std" + %42 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "pred_not" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %47 = OpTypeFunction %void %ulong %ulong + %45 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %bool = OpTypeBool %_ptr_Function_bool = OpTypePointer Function %bool @@ -20,10 +22,10 @@ %false = OpConstantFalse %bool %ulong_1 = OpConstant %ulong 1 %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %47 + %1 = OpFunction %void None %45 %14 = OpFunctionParameter %ulong %15 = OpFunctionParameter %ulong - %42 = OpLabel + %40 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -34,45 +36,43 @@ %9 = OpVariable %_ptr_Function_bool Function OpStore %2 %14 OpStore %3 %15 - %17 = OpLoad %ulong %2 - %16 = OpCopyObject %ulong %17 + %16 = OpLoad %ulong %2 OpStore %4 %16 - %19 = OpLoad %ulong %3 - %18 = OpCopyObject %ulong %19 - OpStore %5 %18 + %17 = OpLoad %ulong %3 + OpStore %5 %17 + %19 = OpLoad %ulong %4 + %37 = OpConvertUToPtr %_ptr_Generic_ulong %19 + %18 = OpLoad %ulong %37 + OpStore %6 %18 %21 = OpLoad %ulong %4 - %39 = OpConvertUToPtr %_ptr_Generic_ulong %21 - %20 = OpLoad %ulong %39 - OpStore %6 %20 - %23 = OpLoad %ulong %4 - %36 = OpIAdd %ulong %23 %ulong_8 - %40 = OpConvertUToPtr %_ptr_Generic_ulong %36 - %22 = OpLoad %ulong %40 - OpStore %7 %22 - %25 = OpLoad %ulong %6 - %26 = OpLoad %ulong %7 - %24 = OpULessThan %bool %25 %26 - OpStore %9 %24 - %28 = OpLoad %bool %9 - %27 = OpSelect %bool %28 %false %true - OpStore %9 %27 - %29 = OpLoad %bool %9 - OpBranchConditional %29 %10 %11 + %34 = OpIAdd %ulong %21 %ulong_8 + %38 = OpConvertUToPtr %_ptr_Generic_ulong %34 + %20 = OpLoad %ulong %38 + OpStore %7 %20 + %23 = OpLoad %ulong %6 + %24 = OpLoad %ulong %7 + %22 = OpULessThan %bool %23 %24 + OpStore %9 %22 + %26 = OpLoad %bool %9 + %25 = OpSelect %bool %26 %false %true + OpStore %9 %25 + %27 = OpLoad %bool %9 + OpBranchConditional %27 %10 %11 %10 = OpLabel - %30 = OpCopyObject %ulong %ulong_1 - OpStore %8 %30 + %28 = OpCopyObject %ulong %ulong_1 + OpStore %8 %28 OpBranch %11 %11 = OpLabel - %31 = OpLoad %bool %9 - OpBranchConditional %31 %13 %12 + %29 = OpLoad %bool %9 + OpBranchConditional %29 %13 %12 %12 = OpLabel - %32 = OpCopyObject %ulong %ulong_2 - OpStore %8 %32 + %30 = OpCopyObject %ulong %ulong_2 + OpStore %8 %30 OpBranch %13 %13 = OpLabel - %33 = OpLoad %ulong %5 - %34 = OpLoad %ulong %8 - %41 = OpConvertUToPtr %_ptr_Generic_ulong %33 - OpStore %41 %34 + %31 = OpLoad %ulong %5 + %32 = OpLoad %ulong %8 + %39 = OpConvertUToPtr %_ptr_Generic_ulong %31 + OpStore %39 %32 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/rcp.spvtxt b/ptx/src/test/spirv_run/rcp.spvtxt index fd10ff1..0ce2d75 100644 --- a/ptx/src/test/spirv_run/rcp.spvtxt +++ b/ptx/src/test/spirv_run/rcp.spvtxt @@ -7,24 +7,22 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - OpExtension "SPV_KHR_float_controls" - %23 = OpExtInstImport "OpenCL.std" + %21 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "rcp" - OpExecutionMode %1 DenormPreserve 32 - OpDecorate %15 FPFastMathMode AllowRecip + OpDecorate %13 FPFastMathMode AllowRecip %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %26 = OpTypeFunction %void %ulong %ulong + %24 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %float = OpTypeFloat 32 %_ptr_Function_float = OpTypePointer Function %float %_ptr_Generic_float = OpTypePointer Generic %float %float_1 = OpConstant %float 1 - %1 = OpFunction %void None %26 + %1 = OpFunction %void None %24 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %21 = OpLabel + %19 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -32,22 +30,20 @@ %6 = OpVariable %_ptr_Function_float Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 - %19 = OpConvertUToPtr %_ptr_Generic_float %14 - %13 = OpLoad %float %19 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpFDiv %float %float_1 %14 OpStore %6 %13 + %15 = OpLoad %ulong %5 %16 = OpLoad %float %6 - %15 = OpFDiv %float %float_1 %16 - OpStore %6 %15 - %17 = OpLoad %ulong %5 - %18 = OpLoad %float %6 - %20 = OpConvertUToPtr %_ptr_Generic_float %17 - OpStore %20 %18 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/reg_local.ptx b/ptx/src/test/spirv_run/reg_local.ptx index f09b95a..5707cea 100644 --- a/ptx/src/test/spirv_run/reg_local.ptx +++ b/ptx/src/test/spirv_run/reg_local.ptx @@ -11,14 +11,13 @@ .reg .u64 in_addr;
.reg .u64 out_addr;
.reg .b64 temp;
- .reg .s64 unused;
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
ld.global.u64 temp, [in_addr];
st.u64 [local_x], temp + 1;
- ld.u64 temp, [local_x];
- st.global.u64 [out_addr], temp;
+ ld.u64 temp, [local_x+0];
+ st.global.u64 [out_addr+0], temp;
ret;
}
\ No newline at end of file diff --git a/ptx/src/test/spirv_run/reg_local.spvtxt b/ptx/src/test/spirv_run/reg_local.spvtxt index 2d6bd08..596cedc 100644 --- a/ptx/src/test/spirv_run/reg_local.spvtxt +++ b/ptx/src/test/spirv_run/reg_local.spvtxt @@ -2,62 +2,66 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %35 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %34 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "reg_local" OpDecorate %4 Alignment 8 %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %38 = OpTypeFunction %void %ulong %ulong + %37 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong - %uchar = OpTypeInt 8 0 %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 %uint_8 = OpConstant %uint 8 %_arr_uchar_uint_8 = OpTypeArray %uchar %uint_8 %_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8 %_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong %ulong_1 = OpConstant %ulong 1 - %1 = OpFunction %void None %38 + %ulong_0 = OpConstant %ulong 0 + %ulong_0_0 = OpConstant %ulong 0 + %1 = OpFunction %void None %37 + %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %10 = OpFunctionParameter %ulong - %33 = OpLabel + %32 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function__arr_uchar_uint_8 Function %5 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function - %8 = OpVariable %_ptr_Function_ulong Function - OpStore %2 %9 - OpStore %3 %10 - %12 = OpLoad %ulong %2 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %3 - %13 = OpCopyObject %ulong %14 - OpStore %6 %13 - %16 = OpLoad %ulong %5 - %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %16 - %26 = OpLoad %ulong %25 - %15 = OpCopyObject %ulong %26 + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %5 %10 + %11 = OpLoad %ulong %3 + OpStore %6 %11 + %13 = OpLoad %ulong %5 + %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %13 + %24 = OpLoad %ulong %25 + %12 = OpCopyObject %ulong %24 + OpStore %7 %12 + %14 = OpLoad %ulong %7 + %26 = OpCopyObject %ulong %14 + %19 = OpIAdd %ulong %26 %ulong_1 + %27 = OpBitcast %_ptr_Function_ulong %4 + OpStore %27 %19 + %28 = OpBitcast %_ptr_Function_ulong %4 + %45 = OpBitcast %ulong %28 + %46 = OpIAdd %ulong %45 %ulong_0 + %21 = OpBitcast %_ptr_Function_ulong %46 + %29 = OpLoad %ulong %21 + %15 = OpCopyObject %ulong %29 OpStore %7 %15 - %18 = OpLoad %ulong %7 - %27 = OpCopyObject %ulong %18 - %24 = OpIAdd %ulong %27 %ulong_1 - %28 = OpCopyObject %ulong %24 - %17 = OpBitcast %ulong %28 - OpStore %4 %17 - %20 = OpLoad %_arr_uchar_uint_8 %4 - %29 = OpBitcast %ulong %20 - %30 = OpCopyObject %ulong %29 - %19 = OpCopyObject %ulong %30 - OpStore %7 %19 - %21 = OpLoad %ulong %6 - %22 = OpLoad %ulong %7 - %31 = OpCopyObject %ulong %22 - %32 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %21 - OpStore %32 %31 + %16 = OpLoad %ulong %6 + %17 = OpLoad %ulong %7 + %23 = OpIAdd %ulong %16 %ulong_0_0 + %30 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %23 + %31 = OpCopyObject %ulong %17 + OpStore %30 %31 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/rsqrt.spvtxt b/ptx/src/test/spirv_run/rsqrt.spvtxt index 5c3ba97..ed473bc 100644 --- a/ptx/src/test/spirv_run/rsqrt.spvtxt +++ b/ptx/src/test/spirv_run/rsqrt.spvtxt @@ -1,56 +1,47 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 31 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%23 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "rsqrt" -OpDecorate %1 FunctionDenormModeINTEL 64 Preserve -%24 = OpTypeVoid -%25 = OpTypeInt 64 0 -%26 = OpTypeFunction %24 %25 %25 -%27 = OpTypePointer Function %25 -%28 = OpTypeFloat 64 -%29 = OpTypePointer Function %28 -%30 = OpTypePointer Generic %28 -%1 = OpFunction %24 None %26 -%7 = OpFunctionParameter %25 -%8 = OpFunctionParameter %25 -%21 = OpLabel -%2 = OpVariable %27 Function -%3 = OpVariable %27 Function -%4 = OpVariable %27 Function -%5 = OpVariable %27 Function -%6 = OpVariable %29 Function -OpStore %2 %7 -OpStore %3 %8 -%10 = OpLoad %25 %2 -%9 = OpCopyObject %25 %10 -OpStore %4 %9 -%12 = OpLoad %25 %3 -%11 = OpCopyObject %25 %12 -OpStore %5 %11 -%14 = OpLoad %25 %4 -%19 = OpConvertUToPtr %30 %14 -%13 = OpLoad %28 %19 -OpStore %6 %13 -%16 = OpLoad %28 %6 -%15 = OpExtInst %28 %23 native_rsqrt %16 -OpStore %6 %15 -%17 = OpLoad %25 %5 -%18 = OpLoad %28 %6 -%20 = OpConvertUToPtr %30 %17 -OpStore %20 %18 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "rsqrt" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %double = OpTypeFloat 64 +%_ptr_Function_double = OpTypePointer Function %double +%_ptr_Generic_double = OpTypePointer Generic %double + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_double Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_double %12 + %11 = OpLoad %double %17 + OpStore %6 %11 + %14 = OpLoad %double %6 + %13 = OpExtInst %double %21 native_rsqrt %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %double %6 + %18 = OpConvertUToPtr %_ptr_Generic_double %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/selp.spvtxt b/ptx/src/test/spirv_run/selp.spvtxt index 6f73bc2..44e87e4 100644 --- a/ptx/src/test/spirv_run/selp.spvtxt +++ b/ptx/src/test/spirv_run/selp.spvtxt @@ -1,65 +1,57 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 40 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%31 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "selp" -%32 = OpTypeVoid -%33 = OpTypeInt 64 0 -%34 = OpTypeFunction %32 %33 %33 -%35 = OpTypePointer Function %33 -%36 = OpTypeInt 16 0 -%37 = OpTypePointer Function %36 -%38 = OpTypePointer Generic %36 -%23 = OpConstant %33 2 -%39 = OpTypeBool -%25 = OpConstantFalse %39 -%1 = OpFunction %32 None %34 -%8 = OpFunctionParameter %33 -%9 = OpFunctionParameter %33 -%29 = OpLabel -%2 = OpVariable %35 Function -%3 = OpVariable %35 Function -%4 = OpVariable %35 Function -%5 = OpVariable %35 Function -%6 = OpVariable %37 Function -%7 = OpVariable %37 Function -OpStore %2 %8 -OpStore %3 %9 -%11 = OpLoad %33 %2 -%10 = OpCopyObject %33 %11 -OpStore %4 %10 -%13 = OpLoad %33 %3 -%12 = OpCopyObject %33 %13 -OpStore %5 %12 -%15 = OpLoad %33 %4 -%26 = OpConvertUToPtr %38 %15 -%14 = OpLoad %36 %26 -OpStore %6 %14 -%17 = OpLoad %33 %4 -%24 = OpIAdd %33 %17 %23 -%27 = OpConvertUToPtr %38 %24 -%16 = OpLoad %36 %27 -OpStore %7 %16 -%19 = OpLoad %36 %6 -%20 = OpLoad %36 %7 -%18 = OpSelect %36 %25 %20 %20 -OpStore %6 %18 -%21 = OpLoad %33 %5 -%22 = OpLoad %36 %6 -%28 = OpConvertUToPtr %38 %21 -OpStore %28 %22 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %29 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "selp" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %32 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %ushort = OpTypeInt 16 0 +%_ptr_Function_ushort = OpTypePointer Function %ushort +%_ptr_Generic_ushort = OpTypePointer Generic %ushort + %ulong_2 = OpConstant %ulong 2 + %bool = OpTypeBool + %false = OpConstantFalse %bool + %1 = OpFunction %void None %32 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %27 = 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_ushort Function + %7 = OpVariable %_ptr_Function_ushort Function + OpStore %2 %8 + OpStore %3 %9 + %10 = OpLoad %ulong %2 + OpStore %4 %10 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %24 = OpConvertUToPtr %_ptr_Generic_ushort %13 + %12 = OpLoad %ushort %24 + OpStore %6 %12 + %15 = OpLoad %ulong %4 + %22 = OpIAdd %ulong %15 %ulong_2 + %25 = OpConvertUToPtr %_ptr_Generic_ushort %22 + %14 = OpLoad %ushort %25 + OpStore %7 %14 + %17 = OpLoad %ushort %6 + %18 = OpLoad %ushort %7 + %16 = OpSelect %ushort %false %18 %18 + OpStore %6 %16 + %19 = OpLoad %ulong %5 + %20 = OpLoad %ushort %6 + %26 = OpConvertUToPtr %_ptr_Generic_ushort %19 + OpStore %26 %20 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/setp.spvtxt b/ptx/src/test/spirv_run/setp.spvtxt index 5e18377..ec94573 100644 --- a/ptx/src/test/spirv_run/setp.spvtxt +++ b/ptx/src/test/spirv_run/setp.spvtxt @@ -2,14 +2,17 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %42 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %40 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "setp" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %45 = OpTypeFunction %void %ulong %ulong + %43 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %bool = OpTypeBool %_ptr_Function_bool = OpTypePointer Function %bool @@ -17,10 +20,10 @@ %ulong_8 = OpConstant %ulong 8 %ulong_1 = OpConstant %ulong 1 %ulong_2 = OpConstant %ulong 2 - %1 = OpFunction %void None %45 + %1 = OpFunction %void None %43 %14 = OpFunctionParameter %ulong %15 = OpFunctionParameter %ulong - %40 = OpLabel + %38 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -31,42 +34,40 @@ %9 = OpVariable %_ptr_Function_bool Function OpStore %2 %14 OpStore %3 %15 - %17 = OpLoad %ulong %2 - %16 = OpCopyObject %ulong %17 + %16 = OpLoad %ulong %2 OpStore %4 %16 - %19 = OpLoad %ulong %3 - %18 = OpCopyObject %ulong %19 - OpStore %5 %18 + %17 = OpLoad %ulong %3 + OpStore %5 %17 + %19 = OpLoad %ulong %4 + %35 = OpConvertUToPtr %_ptr_Generic_ulong %19 + %18 = OpLoad %ulong %35 + OpStore %6 %18 %21 = OpLoad %ulong %4 - %37 = OpConvertUToPtr %_ptr_Generic_ulong %21 - %20 = OpLoad %ulong %37 - OpStore %6 %20 - %23 = OpLoad %ulong %4 - %34 = OpIAdd %ulong %23 %ulong_8 - %38 = OpConvertUToPtr %_ptr_Generic_ulong %34 - %22 = OpLoad %ulong %38 - OpStore %7 %22 - %25 = OpLoad %ulong %6 - %26 = OpLoad %ulong %7 - %24 = OpULessThan %bool %25 %26 - OpStore %9 %24 - %27 = OpLoad %bool %9 - OpBranchConditional %27 %10 %11 + %32 = OpIAdd %ulong %21 %ulong_8 + %36 = OpConvertUToPtr %_ptr_Generic_ulong %32 + %20 = OpLoad %ulong %36 + OpStore %7 %20 + %23 = OpLoad %ulong %6 + %24 = OpLoad %ulong %7 + %22 = OpULessThan %bool %23 %24 + OpStore %9 %22 + %25 = OpLoad %bool %9 + OpBranchConditional %25 %10 %11 %10 = OpLabel - %28 = OpCopyObject %ulong %ulong_1 - OpStore %8 %28 + %26 = OpCopyObject %ulong %ulong_1 + OpStore %8 %26 OpBranch %11 %11 = OpLabel - %29 = OpLoad %bool %9 - OpBranchConditional %29 %13 %12 + %27 = OpLoad %bool %9 + OpBranchConditional %27 %13 %12 %12 = OpLabel - %30 = OpCopyObject %ulong %ulong_2 - OpStore %8 %30 + %28 = OpCopyObject %ulong %ulong_2 + OpStore %8 %28 OpBranch %13 %13 = OpLabel - %31 = OpLoad %ulong %5 - %32 = OpLoad %ulong %8 - %39 = OpConvertUToPtr %_ptr_Generic_ulong %31 - OpStore %39 %32 + %29 = OpLoad %ulong %5 + %30 = OpLoad %ulong %8 + %37 = OpConvertUToPtr %_ptr_Generic_ulong %29 + OpStore %37 %30 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/shared_ptr_32.spvtxt b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt index 609cc0e..98e2501 100644 --- a/ptx/src/test/spirv_run/shared_ptr_32.spvtxt +++ b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt @@ -1,74 +1,66 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 47 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%34 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "shared_ptr_32" %4 -OpDecorate %4 Alignment 4 -%35 = OpTypeVoid -%36 = OpTypeInt 32 0 -%37 = OpTypeInt 8 0 -%38 = OpConstant %36 128 -%39 = OpTypeArray %37 %38 -%40 = OpTypePointer Workgroup %39 -%4 = OpVariable %40 Workgroup -%41 = OpTypeInt 64 0 -%42 = OpTypeFunction %35 %41 %41 -%43 = OpTypePointer Function %41 -%44 = OpTypePointer Function %36 -%45 = OpTypePointer CrossWorkgroup %41 -%46 = OpTypePointer Workgroup %41 -%25 = OpConstant %36 0 -%1 = OpFunction %35 None %42 -%10 = OpFunctionParameter %41 -%11 = OpFunctionParameter %41 -%32 = OpLabel -%2 = OpVariable %43 Function -%3 = OpVariable %43 Function -%5 = OpVariable %43 Function -%6 = OpVariable %43 Function -%7 = OpVariable %44 Function -%8 = OpVariable %43 Function -%9 = OpVariable %43 Function -OpStore %2 %10 -OpStore %3 %11 -%13 = OpLoad %41 %2 -%12 = OpCopyObject %41 %13 -OpStore %5 %12 -%15 = OpLoad %41 %3 -%14 = OpCopyObject %41 %15 -OpStore %6 %14 -%27 = OpConvertPtrToU %36 %4 -%16 = OpCopyObject %36 %27 -OpStore %7 %16 -%18 = OpLoad %41 %5 -%28 = OpConvertUToPtr %45 %18 -%17 = OpLoad %41 %28 -OpStore %8 %17 -%19 = OpLoad %36 %7 -%20 = OpLoad %41 %8 -%29 = OpConvertUToPtr %46 %19 -OpStore %29 %20 -%22 = OpLoad %36 %7 -%26 = OpIAdd %36 %22 %25 -%30 = OpConvertUToPtr %46 %26 -%21 = OpLoad %41 %30 -OpStore %9 %21 -%23 = OpLoad %41 %6 -%24 = OpLoad %41 %9 -%31 = OpConvertUToPtr %45 %23 -OpStore %31 %24 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %32 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "shared_ptr_32" %4 + OpDecorate %4 Alignment 4 + %void = OpTypeVoid + %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 + %uint_128 = OpConstant %uint 128 +%_arr_uchar_uint_128 = OpTypeArray %uchar %uint_128 +%_ptr_Workgroup__arr_uchar_uint_128 = OpTypePointer Workgroup %_arr_uchar_uint_128 + %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup + %ulong = OpTypeInt 64 0 + %40 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong +%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong + %uint_0 = OpConstant %uint 0 + %1 = OpFunction %void None %40 + %10 = OpFunctionParameter %ulong + %11 = OpFunctionParameter %ulong + %30 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_ulong Function + %7 = OpVariable %_ptr_Function_uint Function + %8 = OpVariable %_ptr_Function_ulong Function + %9 = OpVariable %_ptr_Function_ulong Function + OpStore %2 %10 + OpStore %3 %11 + %12 = OpLoad %ulong %2 + OpStore %5 %12 + %13 = OpLoad %ulong %3 + OpStore %6 %13 + %25 = OpConvertPtrToU %uint %4 + %14 = OpCopyObject %uint %25 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %16 + %15 = OpLoad %ulong %26 + OpStore %8 %15 + %17 = OpLoad %uint %7 + %18 = OpLoad %ulong %8 + %27 = OpConvertUToPtr %_ptr_Workgroup_ulong %17 + OpStore %27 %18 + %20 = OpLoad %uint %7 + %24 = OpIAdd %uint %20 %uint_0 + %28 = OpConvertUToPtr %_ptr_Workgroup_ulong %24 + %19 = OpLoad %ulong %28 + OpStore %9 %19 + %21 = OpLoad %ulong %6 + %22 = OpLoad %ulong %9 + %29 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %21 + OpStore %29 %22 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/shared_variable.spvtxt b/ptx/src/test/spirv_run/shared_variable.spvtxt index 1af2bd1..ffd6bd6 100644 --- a/ptx/src/test/spirv_run/shared_variable.spvtxt +++ b/ptx/src/test/spirv_run/shared_variable.spvtxt @@ -1,65 +1,57 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 39 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%27 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "shared_variable" %4 -OpDecorate %4 Alignment 4 -%28 = OpTypeVoid -%29 = OpTypeInt 32 0 -%30 = OpTypeInt 8 0 -%31 = OpConstant %29 128 -%32 = OpTypeArray %30 %31 -%33 = OpTypePointer Workgroup %32 -%4 = OpVariable %33 Workgroup -%34 = OpTypeInt 64 0 -%35 = OpTypeFunction %28 %34 %34 -%36 = OpTypePointer Function %34 -%37 = OpTypePointer CrossWorkgroup %34 -%38 = OpTypePointer Workgroup %34 -%1 = OpFunction %28 None %35 -%9 = OpFunctionParameter %34 -%10 = OpFunctionParameter %34 -%25 = OpLabel -%2 = OpVariable %36 Function -%3 = OpVariable %36 Function -%5 = OpVariable %36 Function -%6 = OpVariable %36 Function -%7 = OpVariable %36 Function -%8 = OpVariable %36 Function -OpStore %2 %9 -OpStore %3 %10 -%12 = OpLoad %34 %2 -%11 = OpCopyObject %34 %12 -OpStore %5 %11 -%14 = OpLoad %34 %3 -%13 = OpCopyObject %34 %14 -OpStore %6 %13 -%16 = OpLoad %34 %5 -%21 = OpConvertUToPtr %37 %16 -%15 = OpLoad %34 %21 -OpStore %7 %15 -%17 = OpLoad %34 %7 -%22 = OpBitcast %38 %4 -OpStore %22 %17 -%23 = OpBitcast %38 %4 -%18 = OpLoad %34 %23 -OpStore %8 %18 -%19 = OpLoad %34 %6 -%20 = OpLoad %34 %8 -%24 = OpConvertUToPtr %37 %19 -OpStore %24 %20 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %25 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "shared_variable" %4 + OpDecorate %4 Alignment 4 + %void = OpTypeVoid + %uint = OpTypeInt 32 0 + %uchar = OpTypeInt 8 0 + %uint_128 = OpConstant %uint 128 +%_arr_uchar_uint_128 = OpTypeArray %uchar %uint_128 +%_ptr_Workgroup__arr_uchar_uint_128 = OpTypePointer Workgroup %_arr_uchar_uint_128 + %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_128 Workgroup + %ulong = OpTypeInt 64 0 + %33 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong +%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong +%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong + %1 = OpFunction %void None %33 + %9 = OpFunctionParameter %ulong + %10 = OpFunctionParameter %ulong + %23 = OpLabel + %2 = OpVariable %_ptr_Function_ulong Function + %3 = OpVariable %_ptr_Function_ulong Function + %5 = OpVariable %_ptr_Function_ulong Function + %6 = OpVariable %_ptr_Function_ulong Function + %7 = OpVariable %_ptr_Function_ulong Function + %8 = OpVariable %_ptr_Function_ulong Function + OpStore %2 %9 + OpStore %3 %10 + %11 = OpLoad %ulong %2 + OpStore %5 %11 + %12 = OpLoad %ulong %3 + OpStore %6 %12 + %14 = OpLoad %ulong %5 + %19 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %14 + %13 = OpLoad %ulong %19 + OpStore %7 %13 + %15 = OpLoad %ulong %7 + %20 = OpBitcast %_ptr_Workgroup_ulong %4 + OpStore %20 %15 + %21 = OpBitcast %_ptr_Workgroup_ulong %4 + %16 = OpLoad %ulong %21 + OpStore %8 %16 + %17 = OpLoad %ulong %6 + %18 = OpLoad %ulong %8 + %22 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %17 + OpStore %22 %18 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/shl.spvtxt b/ptx/src/test/spirv_run/shl.spvtxt index 4843a65..ce19fa5 100644 --- a/ptx/src/test/spirv_run/shl.spvtxt +++ b/ptx/src/test/spirv_run/shl.spvtxt @@ -2,22 +2,25 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %27 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %25 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "shl" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %30 = OpTypeFunction %void %ulong %ulong + %28 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %uint = OpTypeInt 32 0 %uint_2 = OpConstant %uint 2 - %1 = OpFunction %void None %30 + %1 = OpFunction %void None %28 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %25 = OpLabel + %23 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -26,24 +29,22 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %23 = OpCopyObject %ulong %17 - %22 = OpShiftLeftLogical %ulong %23 %uint_2 - %16 = OpCopyObject %ulong %22 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %24 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %24 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %21 = OpCopyObject %ulong %15 + %20 = OpShiftLeftLogical %ulong %21 %uint_2 + %14 = OpCopyObject %ulong %20 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %22 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %22 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/shr.spvtxt b/ptx/src/test/spirv_run/shr.spvtxt index 417839d..893dbf3 100644 --- a/ptx/src/test/spirv_run/shr.spvtxt +++ b/ptx/src/test/spirv_run/shr.spvtxt @@ -7,21 +7,21 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %24 = OpExtInstImport "OpenCL.std" + %22 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "shr" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %27 = OpTypeFunction %void %ulong %ulong + %25 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uint = OpTypeInt 32 0 %_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Generic_uint = OpTypePointer Generic %uint %uint_1 = OpConstant %uint 1 - %1 = OpFunction %void None %27 + %1 = OpFunction %void None %25 %7 = OpFunctionParameter %ulong %8 = OpFunctionParameter %ulong - %22 = OpLabel + %20 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -29,22 +29,20 @@ %6 = OpVariable %_ptr_Function_uint Function OpStore %2 %7 OpStore %3 %8 - %10 = OpLoad %ulong %2 - %9 = OpCopyObject %ulong %10 + %9 = OpLoad %ulong %2 OpStore %4 %9 - %12 = OpLoad %ulong %3 - %11 = OpCopyObject %ulong %12 - OpStore %5 %11 - %14 = OpLoad %ulong %4 - %20 = OpConvertUToPtr %_ptr_Generic_uint %14 - %13 = OpLoad %uint %20 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %18 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %18 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpShiftRightArithmetic %uint %14 %uint_1 OpStore %6 %13 + %15 = OpLoad %ulong %5 %16 = OpLoad %uint %6 - %15 = OpShiftRightArithmetic %uint %16 %uint_1 - OpStore %6 %15 - %17 = OpLoad %ulong %5 - %18 = OpLoad %uint %6 - %21 = OpConvertUToPtr %_ptr_Generic_uint %17 - OpStore %21 %18 + %19 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %19 %16 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/sin.ptx b/ptx/src/test/spirv_run/sin.ptx new file mode 100644 index 0000000..fe94cac --- /dev/null +++ b/ptx/src/test/spirv_run/sin.ptx @@ -0,0 +1,21 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry sin(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .f32 temp;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.f32 temp, [in_addr];
+ sin.approx.f32 temp, temp;
+ st.f32 [out_addr], temp;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/sin.spvtxt b/ptx/src/test/spirv_run/sin.spvtxt new file mode 100644 index 0000000..6656a43 --- /dev/null +++ b/ptx/src/test/spirv_run/sin.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "sin" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 sin %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/sqrt.spvtxt b/ptx/src/test/spirv_run/sqrt.spvtxt index d2c5b20..6d1cfd2 100644 --- a/ptx/src/test/spirv_run/sqrt.spvtxt +++ b/ptx/src/test/spirv_run/sqrt.spvtxt @@ -1,56 +1,47 @@ -; SPIR-V -; Version: 1.3 -; Generator: rspirv -; Bound: 31 -OpCapability GenericPointer -OpCapability Linkage -OpCapability Addresses -OpCapability Kernel -OpCapability Int8 -OpCapability Int16 -OpCapability Int64 -OpCapability Float16 -OpCapability Float64 -; OpCapability FunctionFloatControlINTEL -; OpExtension "SPV_INTEL_float_controls2" -%23 = OpExtInstImport "OpenCL.std" -OpMemoryModel Physical64 OpenCL -OpEntryPoint Kernel %1 "sqrt" -OpDecorate %1 FunctionDenormModeINTEL 32 Preserve -%24 = OpTypeVoid -%25 = OpTypeInt 64 0 -%26 = OpTypeFunction %24 %25 %25 -%27 = OpTypePointer Function %25 -%28 = OpTypeFloat 32 -%29 = OpTypePointer Function %28 -%30 = OpTypePointer Generic %28 -%1 = OpFunction %24 None %26 -%7 = OpFunctionParameter %25 -%8 = OpFunctionParameter %25 -%21 = OpLabel -%2 = OpVariable %27 Function -%3 = OpVariable %27 Function -%4 = OpVariable %27 Function -%5 = OpVariable %27 Function -%6 = OpVariable %29 Function -OpStore %2 %7 -OpStore %3 %8 -%10 = OpLoad %25 %2 -%9 = OpCopyObject %25 %10 -OpStore %4 %9 -%12 = OpLoad %25 %3 -%11 = OpCopyObject %25 %12 -OpStore %5 %11 -%14 = OpLoad %25 %4 -%19 = OpConvertUToPtr %30 %14 -%13 = OpLoad %28 %19 -OpStore %6 %13 -%16 = OpLoad %28 %6 -%15 = OpExtInst %28 %23 native_sqrt %16 -OpStore %6 %15 -%17 = OpLoad %25 %5 -%18 = OpLoad %28 %6 -%20 = OpConvertUToPtr %30 %17 -OpStore %20 %18 -OpReturn -OpFunctionEnd
\ No newline at end of file + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "sqrt" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %float = OpTypeFloat 32 +%_ptr_Function_float = OpTypePointer Function %float +%_ptr_Generic_float = OpTypePointer Generic %float + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_float Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_float %12 + %11 = OpLoad %float %17 + OpStore %6 %11 + %14 = OpLoad %float %6 + %13 = OpExtInst %float %21 native_sqrt %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %float %6 + %18 = OpConvertUToPtr %_ptr_Generic_float %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/sub.spvtxt b/ptx/src/test/spirv_run/sub.spvtxt index 8520168..88017ae 100644 --- a/ptx/src/test/spirv_run/sub.spvtxt +++ b/ptx/src/test/spirv_run/sub.spvtxt @@ -7,19 +7,19 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %25 = OpExtInstImport "OpenCL.std" + %23 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "sub" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %28 = OpTypeFunction %void %ulong %ulong + %26 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong %ulong_1 = OpConstant %ulong 1 - %1 = OpFunction %void None %28 + %1 = OpFunction %void None %26 %8 = OpFunctionParameter %ulong %9 = OpFunctionParameter %ulong - %23 = OpLabel + %21 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -28,22 +28,20 @@ %7 = OpVariable %_ptr_Function_ulong Function OpStore %2 %8 OpStore %3 %9 - %11 = OpLoad %ulong %2 - %10 = OpCopyObject %ulong %11 + %10 = OpLoad %ulong %2 OpStore %4 %10 - %13 = OpLoad %ulong %3 - %12 = OpCopyObject %ulong %13 - OpStore %5 %12 - %15 = OpLoad %ulong %4 - %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 - %14 = OpLoad %ulong %21 - OpStore %6 %14 - %17 = OpLoad %ulong %6 - %16 = OpISub %ulong %17 %ulong_1 - OpStore %7 %16 - %18 = OpLoad %ulong %5 - %19 = OpLoad %ulong %7 - %22 = OpConvertUToPtr %_ptr_Generic_ulong %18 - OpStore %22 %19 + %11 = OpLoad %ulong %3 + OpStore %5 %11 + %13 = OpLoad %ulong %4 + %19 = OpConvertUToPtr %_ptr_Generic_ulong %13 + %12 = OpLoad %ulong %19 + OpStore %6 %12 + %15 = OpLoad %ulong %6 + %14 = OpISub %ulong %15 %ulong_1 + OpStore %7 %14 + %16 = OpLoad %ulong %5 + %17 = OpLoad %ulong %7 + %20 = OpConvertUToPtr %_ptr_Generic_ulong %16 + OpStore %20 %17 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/vector.spvtxt b/ptx/src/test/spirv_run/vector.spvtxt index ff0ee97..535e480 100644 --- a/ptx/src/test/spirv_run/vector.spvtxt +++ b/ptx/src/test/spirv_run/vector.spvtxt @@ -2,26 +2,29 @@ OpCapability Linkage OpCapability Addresses OpCapability Kernel - OpCapability Int64 OpCapability Int8 - %60 = OpExtInstImport "OpenCL.std" + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %57 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %31 "vector" %void = OpTypeVoid %uint = OpTypeInt 32 0 %v2uint = OpTypeVector %uint 2 - %64 = OpTypeFunction %v2uint %v2uint + %61 = OpTypeFunction %v2uint %v2uint %_ptr_Function_v2uint = OpTypePointer Function %v2uint %_ptr_Function_uint = OpTypePointer Function %uint %ulong = OpTypeInt 64 0 - %68 = OpTypeFunction %void %ulong %ulong + %65 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Generic_v2uint = OpTypePointer Generic %v2uint - %1 = OpFunction %v2uint None %64 + %1 = OpFunction %v2uint None %61 %7 = OpFunctionParameter %v2uint %30 = OpLabel - %3 = OpVariable %_ptr_Function_v2uint Function %2 = OpVariable %_ptr_Function_v2uint Function + %3 = OpVariable %_ptr_Function_v2uint Function %4 = OpVariable %_ptr_Function_v2uint Function %5 = OpVariable %_ptr_Function_uint Function %6 = OpVariable %_ptr_Function_uint Function @@ -57,10 +60,10 @@ %26 = OpLoad %v2uint %2 OpReturnValue %26 OpFunctionEnd - %31 = OpFunction %void None %68 + %31 = OpFunction %void None %65 %40 = OpFunctionParameter %ulong %41 = OpFunctionParameter %ulong - %58 = OpLabel + %55 = OpLabel %32 = OpVariable %_ptr_Function_ulong Function %33 = OpVariable %_ptr_Function_ulong Function %34 = OpVariable %_ptr_Function_ulong Function @@ -71,27 +74,24 @@ %39 = OpVariable %_ptr_Function_ulong Function OpStore %32 %40 OpStore %33 %41 - %43 = OpLoad %ulong %32 - %42 = OpCopyObject %ulong %43 + %42 = OpLoad %ulong %32 OpStore %34 %42 - %45 = OpLoad %ulong %33 - %44 = OpCopyObject %ulong %45 - OpStore %35 %44 - %47 = OpLoad %ulong %34 - %54 = OpConvertUToPtr %_ptr_Generic_v2uint %47 - %46 = OpLoad %v2uint %54 + %43 = OpLoad %ulong %33 + OpStore %35 %43 + %45 = OpLoad %ulong %34 + %52 = OpConvertUToPtr %_ptr_Generic_v2uint %45 + %44 = OpLoad %v2uint %52 + OpStore %36 %44 + %47 = OpLoad %v2uint %36 + %46 = OpFunctionCall %v2uint %1 %47 OpStore %36 %46 %49 = OpLoad %v2uint %36 - %48 = OpFunctionCall %v2uint %1 %49 - OpStore %36 %48 + %53 = OpBitcast %ulong %49 + %48 = OpCopyObject %ulong %53 + OpStore %39 %48 + %50 = OpLoad %ulong %35 %51 = OpLoad %v2uint %36 - %55 = OpBitcast %ulong %51 - %56 = OpCopyObject %ulong %55 - %50 = OpCopyObject %ulong %56 - OpStore %39 %50 - %52 = OpLoad %ulong %35 - %53 = OpLoad %v2uint %36 - %57 = OpConvertUToPtr %_ptr_Generic_v2uint %52 - OpStore %57 %53 + %54 = OpConvertUToPtr %_ptr_Generic_v2uint %50 + OpStore %54 %51 OpReturn OpFunctionEnd diff --git a/ptx/src/test/spirv_run/vector_extract.spvtxt b/ptx/src/test/spirv_run/vector_extract.spvtxt index 45df3a8..4943189 100644 --- a/ptx/src/test/spirv_run/vector_extract.spvtxt +++ b/ptx/src/test/spirv_run/vector_extract.spvtxt @@ -7,12 +7,12 @@ OpCapability Int64 OpCapability Float16 OpCapability Float64 - %75 = OpExtInstImport "OpenCL.std" + %73 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %1 "vector_extract" %void = OpTypeVoid %ulong = OpTypeInt 64 0 - %78 = OpTypeFunction %void %ulong %ulong + %76 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %ushort = OpTypeInt 16 0 %_ptr_Function_ushort = OpTypePointer Function %ushort @@ -21,10 +21,10 @@ %uchar = OpTypeInt 8 0 %v4uchar = OpTypeVector %uchar 4 %_ptr_CrossWorkgroup_v4uchar = OpTypePointer CrossWorkgroup %v4uchar - %1 = OpFunction %void None %78 + %1 = OpFunction %void None %76 %11 = OpFunctionParameter %ulong %12 = OpFunctionParameter %ulong - %73 = OpLabel + %71 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function %4 = OpVariable %_ptr_Function_ulong Function @@ -36,89 +36,87 @@ %10 = OpVariable %_ptr_Function_v4ushort Function OpStore %2 %11 OpStore %3 %12 - %14 = OpLoad %ulong %2 - %13 = OpCopyObject %ulong %14 + %13 = OpLoad %ulong %2 OpStore %4 %13 - %16 = OpLoad %ulong %3 - %15 = OpCopyObject %ulong %16 - OpStore %5 %15 - %21 = OpLoad %ulong %4 - %63 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %21 - %45 = OpLoad %v4uchar %63 - %64 = OpCompositeExtract %uchar %45 0 + %14 = OpLoad %ulong %3 + OpStore %5 %14 + %19 = OpLoad %ulong %4 + %61 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %19 + %43 = OpLoad %v4uchar %61 + %62 = OpCompositeExtract %uchar %43 0 + %85 = OpBitcast %uchar %62 + %15 = OpUConvert %ushort %85 + %63 = OpCompositeExtract %uchar %43 1 + %86 = OpBitcast %uchar %63 + %16 = OpUConvert %ushort %86 + %64 = OpCompositeExtract %uchar %43 2 %87 = OpBitcast %uchar %64 %17 = OpUConvert %ushort %87 - %65 = OpCompositeExtract %uchar %45 1 + %65 = OpCompositeExtract %uchar %43 3 %88 = OpBitcast %uchar %65 %18 = OpUConvert %ushort %88 - %66 = OpCompositeExtract %uchar %45 2 - %89 = OpBitcast %uchar %66 - %19 = OpUConvert %ushort %89 - %67 = OpCompositeExtract %uchar %45 3 - %90 = OpBitcast %uchar %67 - %20 = OpUConvert %ushort %90 - OpStore %6 %17 - OpStore %7 %18 - OpStore %8 %19 - OpStore %9 %20 - %23 = OpLoad %ushort %7 - %24 = OpLoad %ushort %8 - %25 = OpLoad %ushort %9 - %26 = OpLoad %ushort %6 - %46 = OpUndef %v4ushort - %47 = OpCompositeInsert %v4ushort %23 %46 0 - %48 = OpCompositeInsert %v4ushort %24 %47 1 - %49 = OpCompositeInsert %v4ushort %25 %48 2 - %50 = OpCompositeInsert %v4ushort %26 %49 3 - %22 = OpCopyObject %v4ushort %50 - OpStore %10 %22 - %31 = OpLoad %v4ushort %10 - %51 = OpCopyObject %v4ushort %31 - %27 = OpCompositeExtract %ushort %51 0 - %28 = OpCompositeExtract %ushort %51 1 - %29 = OpCompositeExtract %ushort %51 2 - %30 = OpCompositeExtract %ushort %51 3 - OpStore %8 %27 - OpStore %9 %28 - OpStore %6 %29 - OpStore %7 %30 - %36 = OpLoad %ushort %8 - %37 = OpLoad %ushort %9 - %38 = OpLoad %ushort %6 - %39 = OpLoad %ushort %7 - %53 = OpUndef %v4ushort - %54 = OpCompositeInsert %v4ushort %36 %53 0 - %55 = OpCompositeInsert %v4ushort %37 %54 1 - %56 = OpCompositeInsert %v4ushort %38 %55 2 - %57 = OpCompositeInsert %v4ushort %39 %56 3 - %52 = OpCopyObject %v4ushort %57 - %32 = OpCompositeExtract %ushort %52 0 - %33 = OpCompositeExtract %ushort %52 1 - %34 = OpCompositeExtract %ushort %52 2 - %35 = OpCompositeExtract %ushort %52 3 - OpStore %9 %32 - OpStore %6 %33 - OpStore %7 %34 - OpStore %8 %35 - %40 = OpLoad %ulong %5 - %41 = OpLoad %ushort %6 - %42 = OpLoad %ushort %7 - %43 = OpLoad %ushort %8 - %44 = OpLoad %ushort %9 - %58 = OpUndef %v4uchar + OpStore %6 %15 + OpStore %7 %16 + OpStore %8 %17 + OpStore %9 %18 + %21 = OpLoad %ushort %7 + %22 = OpLoad %ushort %8 + %23 = OpLoad %ushort %9 + %24 = OpLoad %ushort %6 + %44 = OpUndef %v4ushort + %45 = OpCompositeInsert %v4ushort %21 %44 0 + %46 = OpCompositeInsert %v4ushort %22 %45 1 + %47 = OpCompositeInsert %v4ushort %23 %46 2 + %48 = OpCompositeInsert %v4ushort %24 %47 3 + %20 = OpCopyObject %v4ushort %48 + OpStore %10 %20 + %29 = OpLoad %v4ushort %10 + %49 = OpCopyObject %v4ushort %29 + %25 = OpCompositeExtract %ushort %49 0 + %26 = OpCompositeExtract %ushort %49 1 + %27 = OpCompositeExtract %ushort %49 2 + %28 = OpCompositeExtract %ushort %49 3 + OpStore %8 %25 + OpStore %9 %26 + OpStore %6 %27 + OpStore %7 %28 + %34 = OpLoad %ushort %8 + %35 = OpLoad %ushort %9 + %36 = OpLoad %ushort %6 + %37 = OpLoad %ushort %7 + %51 = OpUndef %v4ushort + %52 = OpCompositeInsert %v4ushort %34 %51 0 + %53 = OpCompositeInsert %v4ushort %35 %52 1 + %54 = OpCompositeInsert %v4ushort %36 %53 2 + %55 = OpCompositeInsert %v4ushort %37 %54 3 + %50 = OpCopyObject %v4ushort %55 + %30 = OpCompositeExtract %ushort %50 0 + %31 = OpCompositeExtract %ushort %50 1 + %32 = OpCompositeExtract %ushort %50 2 + %33 = OpCompositeExtract %ushort %50 3 + OpStore %9 %30 + OpStore %6 %31 + OpStore %7 %32 + OpStore %8 %33 + %38 = OpLoad %ulong %5 + %39 = OpLoad %ushort %6 + %40 = OpLoad %ushort %7 + %41 = OpLoad %ushort %8 + %42 = OpLoad %ushort %9 + %56 = OpUndef %v4uchar + %89 = OpBitcast %ushort %39 + %66 = OpUConvert %uchar %89 + %57 = OpCompositeInsert %v4uchar %66 %56 0 + %90 = OpBitcast %ushort %40 + %67 = OpUConvert %uchar %90 + %58 = OpCompositeInsert %v4uchar %67 %57 1 %91 = OpBitcast %ushort %41 %68 = OpUConvert %uchar %91 - %59 = OpCompositeInsert %v4uchar %68 %58 0 + %59 = OpCompositeInsert %v4uchar %68 %58 2 %92 = OpBitcast %ushort %42 %69 = OpUConvert %uchar %92 - %60 = OpCompositeInsert %v4uchar %69 %59 1 - %93 = OpBitcast %ushort %43 - %70 = OpUConvert %uchar %93 - %61 = OpCompositeInsert %v4uchar %70 %60 2 - %94 = OpBitcast %ushort %44 - %71 = OpUConvert %uchar %94 - %62 = OpCompositeInsert %v4uchar %71 %61 3 - %72 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %40 - OpStore %72 %62 + %60 = OpCompositeInsert %v4uchar %69 %59 3 + %70 = OpConvertUToPtr %_ptr_CrossWorkgroup_v4uchar %38 + OpStore %70 %60 OpReturn OpFunctionEnd |