Deal with mismatched signs in S/U/F conversion opcodes.
This commit is contained in:
Hans-Kristian Arntzen
parent
7bb74c99aa
commit
9ae91c2d1e
@ -0,0 +1,42 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
int s32;
|
||||
uint u32;
|
||||
short s16;
|
||||
ushort u16;
|
||||
float f32;
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO& _4 [[buffer(0)]])
|
||||
{
|
||||
int _29 = _4.s32;
|
||||
uint _30 = _4.u32;
|
||||
short _31 = _4.s16;
|
||||
ushort _32 = _4.u16;
|
||||
float _33 = _4.f32;
|
||||
_4.s32 = int(_31);
|
||||
_4.u32 = uint(_31);
|
||||
_4.s32 = int(short(_32));
|
||||
_4.u32 = uint(short(_32));
|
||||
_4.u32 = uint(ushort(_31));
|
||||
_4.u32 = uint(_32);
|
||||
_4.s16 = short(_29);
|
||||
_4.u16 = ushort(_29);
|
||||
_4.s16 = short(_30);
|
||||
_4.u16 = ushort(_30);
|
||||
_4.u16 = ushort(_29);
|
||||
_4.u16 = ushort(_30);
|
||||
_4.f32 = float(_31);
|
||||
_4.f32 = float(short(_32));
|
||||
_4.f32 = float(ushort(_31));
|
||||
_4.f32 = float(_32);
|
||||
_4.s16 = short(_33);
|
||||
_4.u16 = ushort(short(_33));
|
||||
_4.u16 = ushort(_33);
|
||||
}
|
||||
|
||||
@ -0,0 +1,42 @@
|
||||
#version 450
|
||||
#extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
|
||||
#extension GL_EXT_shader_16bit_storage : require
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
layout(set = 0, binding = 0, std430) buffer SSBO
|
||||
{
|
||||
int s32;
|
||||
uint u32;
|
||||
int16_t s16;
|
||||
uint16_t u16;
|
||||
float f32;
|
||||
} _4;
|
||||
|
||||
void main()
|
||||
{
|
||||
int _29 = _4.s32;
|
||||
uint _30 = _4.u32;
|
||||
int16_t _31 = _4.s16;
|
||||
uint16_t _32 = _4.u16;
|
||||
float _33 = _4.f32;
|
||||
_4.s32 = int(_31);
|
||||
_4.u32 = uint(_31);
|
||||
_4.s32 = int(int16_t(_32));
|
||||
_4.u32 = uint(int16_t(_32));
|
||||
_4.u32 = uint(uint16_t(_31));
|
||||
_4.u32 = uint(_32);
|
||||
_4.s16 = int16_t(_29);
|
||||
_4.u16 = uint16_t(_29);
|
||||
_4.s16 = int16_t(_30);
|
||||
_4.u16 = uint16_t(_30);
|
||||
_4.u16 = uint16_t(_29);
|
||||
_4.u16 = uint16_t(_30);
|
||||
_4.f32 = float(_31);
|
||||
_4.f32 = float(int16_t(_32));
|
||||
_4.f32 = float(uint16_t(_31));
|
||||
_4.f32 = float(_32);
|
||||
_4.s16 = int16_t(_33);
|
||||
_4.u16 = uint16_t(int16_t(_33));
|
||||
_4.u16 = uint16_t(_33);
|
||||
}
|
||||
|
||||
@ -0,0 +1,33 @@
|
||||
#version 450
|
||||
#extension GL_ARB_gpu_shader_int64 : require
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
#ifndef SPIRV_CROSS_CONSTANT_ID_0
|
||||
#define SPIRV_CROSS_CONSTANT_ID_0 1
|
||||
#endif
|
||||
const int ConstantInt = SPIRV_CROSS_CONSTANT_ID_0;
|
||||
#ifndef SPIRV_CROSS_CONSTANT_ID_1
|
||||
#define SPIRV_CROSS_CONSTANT_ID_1 2u
|
||||
#endif
|
||||
const uint ConstantUint = SPIRV_CROSS_CONSTANT_ID_1;
|
||||
const int64_t ConstantInt64_1 = int64_t(ConstantInt);
|
||||
const int64_t ConstantInt64_2 = int64_t(int(ConstantUint));
|
||||
const uint64_t ConstantUint64_1 = uint64_t(ConstantInt);
|
||||
const uint64_t ConstantUint64_2 = uint64_t(int(ConstantUint));
|
||||
const int64_t _20 = (ConstantInt64_1 + ConstantInt64_2);
|
||||
const uint64_t _21 = (ConstantUint64_1 + ConstantUint64_2);
|
||||
const int _22 = int(_20);
|
||||
const uint _23 = uint(_21);
|
||||
|
||||
layout(binding = 0, std430) buffer SSBO
|
||||
{
|
||||
int s64;
|
||||
uint u64;
|
||||
} _4;
|
||||
|
||||
void main()
|
||||
{
|
||||
_4.s64 = _22;
|
||||
_4.u64 = _23;
|
||||
}
|
||||
|
||||
131
shaders-msl-no-opt/asm/comp/arithmetic-conversion-signs.asm.comp
Normal file
131
shaders-msl-no-opt/asm/comp/arithmetic-conversion-signs.asm.comp
Normal file
@ -0,0 +1,131 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 76
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability Int16
|
||||
OpCapability StorageBuffer16BitAccess
|
||||
OpExtension "SPV_KHR_16bit_storage"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpSource GLSL 450
|
||||
OpSourceExtension "GL_EXT_shader_explicit_arithmetic_types_int16"
|
||||
OpName %main "main"
|
||||
OpName %SSBO "SSBO"
|
||||
OpMemberName %SSBO 0 "s32"
|
||||
OpMemberName %SSBO 1 "u32"
|
||||
OpMemberName %SSBO 2 "s16"
|
||||
OpMemberName %SSBO 3 "u16"
|
||||
OpMemberName %SSBO 4 "f32"
|
||||
OpName %_ ""
|
||||
OpMemberDecorate %SSBO 0 Offset 0
|
||||
OpMemberDecorate %SSBO 1 Offset 4
|
||||
OpMemberDecorate %SSBO 2 Offset 8
|
||||
OpMemberDecorate %SSBO 3 Offset 10
|
||||
OpMemberDecorate %SSBO 4 Offset 12
|
||||
OpDecorate %SSBO BufferBlock
|
||||
OpDecorate %_ DescriptorSet 0
|
||||
OpDecorate %_ Binding 0
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%uint = OpTypeInt 32 0
|
||||
%short = OpTypeInt 16 1
|
||||
%ushort = OpTypeInt 16 0
|
||||
%float = OpTypeFloat 32
|
||||
%SSBO = OpTypeStruct %int %uint %short %ushort %float
|
||||
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
|
||||
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
|
||||
%int_2 = OpConstant %int 2
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||
%_ptr_Uniform_short = OpTypePointer Uniform %short
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%int_3 = OpConstant %int 3
|
||||
%_ptr_Uniform_ushort = OpTypePointer Uniform %ushort
|
||||
%int_4 = OpConstant %int 4
|
||||
%_ptr_Uniform_float = OpTypePointer Uniform %float
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%ptr_s32 = OpAccessChain %_ptr_Uniform_int %_ %int_0
|
||||
%ptr_u32 = OpAccessChain %_ptr_Uniform_uint %_ %int_1
|
||||
%ptr_s16 = OpAccessChain %_ptr_Uniform_short %_ %int_2
|
||||
%ptr_u16 = OpAccessChain %_ptr_Uniform_ushort %_ %int_3
|
||||
%ptr_f32 = OpAccessChain %_ptr_Uniform_float %_ %int_4
|
||||
%s32 = OpLoad %int %ptr_s32
|
||||
%u32 = OpLoad %uint %ptr_u32
|
||||
%s16 = OpLoad %short %ptr_s16
|
||||
%u16 = OpLoad %ushort %ptr_u16
|
||||
%f32 = OpLoad %float %ptr_f32
|
||||
|
||||
; Sign-extend
|
||||
%s16_to_s32_signed = OpSConvert %int %s16
|
||||
OpStore %ptr_s32 %s16_to_s32_signed
|
||||
%s16_to_u32_signed = OpSConvert %uint %s16
|
||||
OpStore %ptr_u32 %s16_to_u32_signed
|
||||
|
||||
%u16_to_s32_signed = OpSConvert %int %u16
|
||||
OpStore %ptr_s32 %u16_to_s32_signed
|
||||
%u16_to_u32_signed = OpSConvert %uint %u16
|
||||
OpStore %ptr_u32 %u16_to_u32_signed
|
||||
|
||||
; Zero-extend
|
||||
; Result must be unsigned for OpUConvert.
|
||||
;%s16_to_s32_unsigned = OpUConvert %int %s16
|
||||
;OpStore %ptr_s32 %s16_to_s32_unsigned
|
||||
%s16_to_u32_unsigned = OpUConvert %uint %s16
|
||||
OpStore %ptr_u32 %s16_to_u32_unsigned
|
||||
|
||||
;%u16_to_s32_unsigned = OpUConvert %int %u16
|
||||
;OpStore %ptr_s32 %u16_to_s32_unsigned
|
||||
%u16_to_u32_unsigned = OpUConvert %uint %u16
|
||||
OpStore %ptr_u32 %u16_to_u32_unsigned
|
||||
|
||||
; Truncate (SConvert == UConvert)
|
||||
%s32_to_s16_signed = OpSConvert %short %s32
|
||||
OpStore %ptr_s16 %s32_to_s16_signed
|
||||
%s32_to_u16_signed = OpSConvert %ushort %s32
|
||||
OpStore %ptr_u16 %s32_to_u16_signed
|
||||
|
||||
%u32_to_s16_signed = OpSConvert %short %u32
|
||||
OpStore %ptr_s16 %u32_to_s16_signed
|
||||
%u32_to_u16_signed = OpSConvert %ushort %u32
|
||||
OpStore %ptr_u16 %u32_to_u16_signed
|
||||
|
||||
;%s32_to_s16_unsigned = OpUConvert %short %s32
|
||||
;OpStore %ptr_s16 %s32_to_s16_unsigned
|
||||
%s32_to_u16_unsigned = OpUConvert %ushort %s32
|
||||
OpStore %ptr_u16 %s32_to_u16_unsigned
|
||||
|
||||
;%u32_to_s16_unsigned = OpUConvert %short %u32
|
||||
;OpStore %ptr_s16 %u32_to_s16_unsigned
|
||||
%u32_to_u16_unsigned = OpUConvert %ushort %u32
|
||||
OpStore %ptr_u16 %u32_to_u16_unsigned
|
||||
|
||||
; SToF
|
||||
%s16_to_f32_signed = OpConvertSToF %float %s16
|
||||
OpStore %ptr_f32 %s16_to_f32_signed
|
||||
%u16_to_f32_signed = OpConvertSToF %float %u16
|
||||
OpStore %ptr_f32 %u16_to_f32_signed
|
||||
%s16_to_f32_unsigned = OpConvertUToF %float %s16
|
||||
OpStore %ptr_f32 %s16_to_f32_unsigned
|
||||
%u16_to_f32_unsigned = OpConvertUToF %float %u16
|
||||
OpStore %ptr_f32 %u16_to_f32_unsigned
|
||||
|
||||
; FToS
|
||||
%f32_to_s16_signed = OpConvertFToS %short %f32
|
||||
OpStore %ptr_s16 %f32_to_s16_signed
|
||||
%f32_to_u16_signed = OpConvertFToS %ushort %f32
|
||||
OpStore %ptr_u16 %f32_to_u16_signed
|
||||
|
||||
; FToU
|
||||
%f32_to_u16_unsigned = OpConvertFToU %ushort %f32
|
||||
OpStore %ptr_u16 %f32_to_u16_unsigned
|
||||
; Result must be unsigned for FToU, so don't bother testing that.
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -0,0 +1,131 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 76
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability Int16
|
||||
OpCapability StorageBuffer16BitAccess
|
||||
OpExtension "SPV_KHR_16bit_storage"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpSource GLSL 450
|
||||
OpSourceExtension "GL_EXT_shader_explicit_arithmetic_types_int16"
|
||||
OpName %main "main"
|
||||
OpName %SSBO "SSBO"
|
||||
OpMemberName %SSBO 0 "s32"
|
||||
OpMemberName %SSBO 1 "u32"
|
||||
OpMemberName %SSBO 2 "s16"
|
||||
OpMemberName %SSBO 3 "u16"
|
||||
OpMemberName %SSBO 4 "f32"
|
||||
OpName %_ ""
|
||||
OpMemberDecorate %SSBO 0 Offset 0
|
||||
OpMemberDecorate %SSBO 1 Offset 4
|
||||
OpMemberDecorate %SSBO 2 Offset 8
|
||||
OpMemberDecorate %SSBO 3 Offset 10
|
||||
OpMemberDecorate %SSBO 4 Offset 12
|
||||
OpDecorate %SSBO BufferBlock
|
||||
OpDecorate %_ DescriptorSet 0
|
||||
OpDecorate %_ Binding 0
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%uint = OpTypeInt 32 0
|
||||
%short = OpTypeInt 16 1
|
||||
%ushort = OpTypeInt 16 0
|
||||
%float = OpTypeFloat 32
|
||||
%SSBO = OpTypeStruct %int %uint %short %ushort %float
|
||||
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
|
||||
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
|
||||
%int_2 = OpConstant %int 2
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||
%_ptr_Uniform_short = OpTypePointer Uniform %short
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%int_3 = OpConstant %int 3
|
||||
%_ptr_Uniform_ushort = OpTypePointer Uniform %ushort
|
||||
%int_4 = OpConstant %int 4
|
||||
%_ptr_Uniform_float = OpTypePointer Uniform %float
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%ptr_s32 = OpAccessChain %_ptr_Uniform_int %_ %int_0
|
||||
%ptr_u32 = OpAccessChain %_ptr_Uniform_uint %_ %int_1
|
||||
%ptr_s16 = OpAccessChain %_ptr_Uniform_short %_ %int_2
|
||||
%ptr_u16 = OpAccessChain %_ptr_Uniform_ushort %_ %int_3
|
||||
%ptr_f32 = OpAccessChain %_ptr_Uniform_float %_ %int_4
|
||||
%s32 = OpLoad %int %ptr_s32
|
||||
%u32 = OpLoad %uint %ptr_u32
|
||||
%s16 = OpLoad %short %ptr_s16
|
||||
%u16 = OpLoad %ushort %ptr_u16
|
||||
%f32 = OpLoad %float %ptr_f32
|
||||
|
||||
; Sign-extend
|
||||
%s16_to_s32_signed = OpSConvert %int %s16
|
||||
OpStore %ptr_s32 %s16_to_s32_signed
|
||||
%s16_to_u32_signed = OpSConvert %uint %s16
|
||||
OpStore %ptr_u32 %s16_to_u32_signed
|
||||
|
||||
%u16_to_s32_signed = OpSConvert %int %u16
|
||||
OpStore %ptr_s32 %u16_to_s32_signed
|
||||
%u16_to_u32_signed = OpSConvert %uint %u16
|
||||
OpStore %ptr_u32 %u16_to_u32_signed
|
||||
|
||||
; Zero-extend
|
||||
; Result must be unsigned for OpUConvert.
|
||||
;%s16_to_s32_unsigned = OpUConvert %int %s16
|
||||
;OpStore %ptr_s32 %s16_to_s32_unsigned
|
||||
%s16_to_u32_unsigned = OpUConvert %uint %s16
|
||||
OpStore %ptr_u32 %s16_to_u32_unsigned
|
||||
|
||||
;%u16_to_s32_unsigned = OpUConvert %int %u16
|
||||
;OpStore %ptr_s32 %u16_to_s32_unsigned
|
||||
%u16_to_u32_unsigned = OpUConvert %uint %u16
|
||||
OpStore %ptr_u32 %u16_to_u32_unsigned
|
||||
|
||||
; Truncate (SConvert == UConvert)
|
||||
%s32_to_s16_signed = OpSConvert %short %s32
|
||||
OpStore %ptr_s16 %s32_to_s16_signed
|
||||
%s32_to_u16_signed = OpSConvert %ushort %s32
|
||||
OpStore %ptr_u16 %s32_to_u16_signed
|
||||
|
||||
%u32_to_s16_signed = OpSConvert %short %u32
|
||||
OpStore %ptr_s16 %u32_to_s16_signed
|
||||
%u32_to_u16_signed = OpSConvert %ushort %u32
|
||||
OpStore %ptr_u16 %u32_to_u16_signed
|
||||
|
||||
;%s32_to_s16_unsigned = OpUConvert %short %s32
|
||||
;OpStore %ptr_s16 %s32_to_s16_unsigned
|
||||
%s32_to_u16_unsigned = OpUConvert %ushort %s32
|
||||
OpStore %ptr_u16 %s32_to_u16_unsigned
|
||||
|
||||
;%u32_to_s16_unsigned = OpUConvert %short %u32
|
||||
;OpStore %ptr_s16 %u32_to_s16_unsigned
|
||||
%u32_to_u16_unsigned = OpUConvert %ushort %u32
|
||||
OpStore %ptr_u16 %u32_to_u16_unsigned
|
||||
|
||||
; SToF
|
||||
%s16_to_f32_signed = OpConvertSToF %float %s16
|
||||
OpStore %ptr_f32 %s16_to_f32_signed
|
||||
%u16_to_f32_signed = OpConvertSToF %float %u16
|
||||
OpStore %ptr_f32 %u16_to_f32_signed
|
||||
%s16_to_f32_unsigned = OpConvertUToF %float %s16
|
||||
OpStore %ptr_f32 %s16_to_f32_unsigned
|
||||
%u16_to_f32_unsigned = OpConvertUToF %float %u16
|
||||
OpStore %ptr_f32 %u16_to_f32_unsigned
|
||||
|
||||
; FToS
|
||||
%f32_to_s16_signed = OpConvertFToS %short %f32
|
||||
OpStore %ptr_s16 %f32_to_s16_signed
|
||||
%f32_to_u16_signed = OpConvertFToS %ushort %f32
|
||||
OpStore %ptr_u16 %f32_to_u16_signed
|
||||
|
||||
; FToU
|
||||
%f32_to_u16_unsigned = OpConvertFToU %ushort %f32
|
||||
OpStore %ptr_u16 %f32_to_u16_unsigned
|
||||
; Result must be unsigned for FToU, so don't bother testing that.
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -0,0 +1,63 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 30
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability Int64
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpSource GLSL 450
|
||||
OpSourceExtension "GL_ARB_gpu_shader_int64"
|
||||
OpName %main "main"
|
||||
OpName %SSBO "SSBO"
|
||||
OpMemberName %SSBO 0 "s64"
|
||||
OpMemberName %SSBO 1 "u64"
|
||||
OpName %_ ""
|
||||
OpName %ConstantInt "ConstantInt"
|
||||
OpName %ConstantInt64_1 "ConstantInt64_1"
|
||||
OpName %ConstantUint "ConstantUint"
|
||||
OpName %ConstantInt64_2 "ConstantInt64_2"
|
||||
OpName %ConstantUint64_1 "ConstantUint64_1"
|
||||
OpName %ConstantUint64_2 "ConstantUint64_2"
|
||||
OpMemberDecorate %SSBO 0 Offset 0
|
||||
OpMemberDecorate %SSBO 1 Offset 4
|
||||
OpDecorate %SSBO BufferBlock
|
||||
OpDecorate %_ DescriptorSet 0
|
||||
OpDecorate %_ Binding 0
|
||||
OpDecorate %ConstantInt SpecId 0
|
||||
OpDecorate %ConstantUint SpecId 1
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%uint = OpTypeInt 32 0
|
||||
%long = OpTypeInt 64 1
|
||||
%ulong = OpTypeInt 64 0
|
||||
%SSBO = OpTypeStruct %int %uint
|
||||
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
|
||||
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
|
||||
%int_0 = OpConstant %int 0
|
||||
%ulong_0 = OpConstant %ulong 0
|
||||
%ConstantInt = OpSpecConstant %int 1
|
||||
%ConstantUint = OpSpecConstant %uint 2
|
||||
%ConstantInt64_1 = OpSpecConstantOp %long SConvert %ConstantInt
|
||||
%ConstantInt64_2 = OpSpecConstantOp %long SConvert %ConstantUint
|
||||
%ConstantUint64_1 = OpSpecConstantOp %ulong SConvert %ConstantInt
|
||||
%ConstantUint64_2 = OpSpecConstantOp %ulong SConvert %ConstantUint
|
||||
%added_long = OpSpecConstantOp %long IAdd %ConstantInt64_1 %ConstantInt64_2
|
||||
%added_ulong = OpSpecConstantOp %ulong IAdd %ConstantUint64_1 %ConstantUint64_2
|
||||
%trunc_long = OpSpecConstantOp %int SConvert %added_long
|
||||
%trunc_ulong = OpSpecConstantOp %uint SConvert %added_ulong
|
||||
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%22 = OpAccessChain %_ptr_Uniform_int %_ %int_0
|
||||
OpStore %22 %trunc_long
|
||||
%29 = OpAccessChain %_ptr_Uniform_uint %_ %int_1
|
||||
OpStore %29 %trunc_ulong
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -2868,7 +2868,7 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
|
||||
}
|
||||
|
||||
uint32_t bit_width = 0;
|
||||
if (unary || binary)
|
||||
if (unary || binary || cop.opcode == OpSConvert || cop.opcode == OpUConvert)
|
||||
bit_width = expression_type(cop.arguments[0]).width;
|
||||
|
||||
SPIRType::BaseType input_type;
|
||||
@ -2888,6 +2888,8 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
|
||||
case OpSMod:
|
||||
case OpSDiv:
|
||||
case OpShiftRightArithmetic:
|
||||
case OpSConvert:
|
||||
case OpSNegate:
|
||||
input_type = to_signed_basetype(bit_width);
|
||||
break;
|
||||
|
||||
@ -2898,6 +2900,7 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
|
||||
case OpUMod:
|
||||
case OpUDiv:
|
||||
case OpShiftRightLogical:
|
||||
case OpUConvert:
|
||||
input_type = to_unsigned_basetype(bit_width);
|
||||
break;
|
||||
|
||||
@ -2939,6 +2942,21 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
|
||||
// Works around various casting scenarios in glslang as there is no OpBitcast for specialization constants.
|
||||
return join("(", op, bitcast_glsl(type, cop.arguments[0]), ")");
|
||||
}
|
||||
else if (cop.opcode == OpSConvert || cop.opcode == OpUConvert)
|
||||
{
|
||||
if (cop.arguments.size() < 1)
|
||||
SPIRV_CROSS_THROW("Not enough arguments to OpSpecConstantOp.");
|
||||
|
||||
auto &arg_type = expression_type(cop.arguments[0]);
|
||||
if (arg_type.width < type.width && input_type != arg_type.basetype)
|
||||
{
|
||||
auto expected = arg_type;
|
||||
expected.basetype = input_type;
|
||||
return join(op, "(", bitcast_glsl(expected, cop.arguments[0]), ")");
|
||||
}
|
||||
else
|
||||
return join(op, "(", to_expression(cop.arguments[0]), ")");
|
||||
}
|
||||
else
|
||||
{
|
||||
if (cop.arguments.size() < 1)
|
||||
@ -3822,15 +3840,19 @@ void CompilerGLSL::emit_unary_func_op_cast(uint32_t result_type, uint32_t result
|
||||
SPIRType::BaseType input_type, SPIRType::BaseType expected_result_type)
|
||||
{
|
||||
auto &out_type = get<SPIRType>(result_type);
|
||||
auto &expr_type = expression_type(op0);
|
||||
auto expected_type = out_type;
|
||||
|
||||
// Bit-widths might be different in unary cases because we use it for SConvert/UConvert and friends.
|
||||
expected_type.basetype = input_type;
|
||||
string cast_op =
|
||||
expression_type(op0).basetype != input_type ? bitcast_glsl(expected_type, op0) : to_expression(op0);
|
||||
expected_type.width = expr_type.width;
|
||||
string cast_op = expr_type.basetype != input_type ? bitcast_glsl(expected_type, op0) : to_unpacked_expression(op0);
|
||||
|
||||
string expr;
|
||||
if (out_type.basetype != expected_result_type)
|
||||
{
|
||||
expected_type.basetype = expected_result_type;
|
||||
expected_type.width = out_type.width;
|
||||
expr = bitcast_glsl_op(out_type, expected_type);
|
||||
expr += '(';
|
||||
expr += join(op, "(", cast_op, ")");
|
||||
@ -3852,11 +3874,11 @@ void CompilerGLSL::emit_trinary_func_op_cast(uint32_t result_type, uint32_t resu
|
||||
auto expected_type = out_type;
|
||||
expected_type.basetype = input_type;
|
||||
string cast_op0 =
|
||||
expression_type(op0).basetype != input_type ? bitcast_glsl(expected_type, op0) : to_expression(op0);
|
||||
expression_type(op0).basetype != input_type ? bitcast_glsl(expected_type, op0) : to_unpacked_expression(op0);
|
||||
string cast_op1 =
|
||||
expression_type(op1).basetype != input_type ? bitcast_glsl(expected_type, op1) : to_expression(op1);
|
||||
expression_type(op1).basetype != input_type ? bitcast_glsl(expected_type, op1) : to_unpacked_expression(op1);
|
||||
string cast_op2 =
|
||||
expression_type(op2).basetype != input_type ? bitcast_glsl(expected_type, op2) : to_expression(op2);
|
||||
expression_type(op2).basetype != input_type ? bitcast_glsl(expected_type, op2) : to_unpacked_expression(op2);
|
||||
|
||||
string expr;
|
||||
if (out_type.basetype != input_type)
|
||||
@ -5622,9 +5644,9 @@ string CompilerGLSL::bitcast_glsl(const SPIRType &result_type, uint32_t argument
|
||||
{
|
||||
auto op = bitcast_glsl_op(result_type, expression_type(argument));
|
||||
if (op.empty())
|
||||
return to_enclosed_expression(argument);
|
||||
return to_enclosed_unpacked_expression(argument);
|
||||
else
|
||||
return join(op, "(", to_expression(argument), ")");
|
||||
return join(op, "(", to_unpacked_expression(argument), ")");
|
||||
}
|
||||
|
||||
std::string CompilerGLSL::bitcast_expression(SPIRType::BaseType target_type, uint32_t arg)
|
||||
@ -7077,6 +7099,10 @@ uint32_t CompilerGLSL::get_integer_width_for_instruction(const Instruction &inst
|
||||
|
||||
switch (instr.op)
|
||||
{
|
||||
case OpSConvert:
|
||||
case OpConvertSToF:
|
||||
case OpUConvert:
|
||||
case OpConvertUToF:
|
||||
case OpIEqual:
|
||||
case OpINotEqual:
|
||||
case OpSLessThan:
|
||||
@ -8150,12 +8176,45 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
}
|
||||
|
||||
// Conversion
|
||||
case OpSConvert:
|
||||
case OpConvertSToF:
|
||||
case OpUConvert:
|
||||
case OpConvertUToF:
|
||||
{
|
||||
auto input_type = opcode == OpSConvert || opcode == OpConvertSToF ? int_type : uint_type;
|
||||
uint32_t result_type = ops[0];
|
||||
uint32_t id = ops[1];
|
||||
|
||||
auto &type = get<SPIRType>(result_type);
|
||||
auto &arg_type = expression_type(ops[2]);
|
||||
auto func = type_to_glsl_constructor(type);
|
||||
|
||||
// If we're sign-extending or zero-extending, we need to make sure we cast from the correct type.
|
||||
// For truncation, it does not matter, so don't emit useless casts.
|
||||
if (arg_type.width < type.width)
|
||||
emit_unary_func_op_cast(result_type, id, ops[2], func.c_str(), input_type, type.basetype);
|
||||
else
|
||||
emit_unary_func_op(result_type, id, ops[2], func.c_str());
|
||||
break;
|
||||
}
|
||||
|
||||
case OpConvertFToU:
|
||||
case OpConvertFToS:
|
||||
case OpConvertSToF:
|
||||
case OpConvertUToF:
|
||||
case OpUConvert:
|
||||
case OpSConvert:
|
||||
{
|
||||
// Cast to expected arithmetic type, then potentially bitcast away to desired signedness.
|
||||
uint32_t result_type = ops[0];
|
||||
uint32_t id = ops[1];
|
||||
auto &type = get<SPIRType>(result_type);
|
||||
auto expected_type = type;
|
||||
auto &float_type = expression_type(ops[2]);
|
||||
expected_type.basetype =
|
||||
opcode == OpConvertFToS ? to_signed_basetype(type.width) : to_unsigned_basetype(type.width);
|
||||
|
||||
auto func = type_to_glsl_constructor(expected_type);
|
||||
emit_unary_func_op_cast(result_type, id, ops[2], func.c_str(), float_type.basetype, expected_type.basetype);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpFConvert:
|
||||
{
|
||||
uint32_t result_type = ops[0];
|
||||
|
||||
@ -404,7 +404,7 @@ def regression_check_reflect(shader, json_file, args):
|
||||
with open(reference) as f:
|
||||
expected = json.load(f)
|
||||
if (json_compare(actual, expected) != True):
|
||||
if update:
|
||||
if args.update:
|
||||
print('Generated reflection json has changed for {}!'.format(reference))
|
||||
# If we expect changes, update the reference file.
|
||||
if os.path.exists(reference):
|
||||
@ -440,7 +440,7 @@ def regression_check(shader, glsl, args):
|
||||
|
||||
if os.path.exists(reference):
|
||||
if md5_for_file(glsl) != md5_for_file(reference):
|
||||
if update:
|
||||
if args.update:
|
||||
print('Generated source code has changed for {}!'.format(reference))
|
||||
# If we expect changes, update the reference file.
|
||||
if os.path.exists(reference):
|
||||
|
||||
Loading…
Reference in New Issue
Block a user