src/tint/lang/spirv/reader/parser/convert_test.cc

// Copyright 2025 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
//    list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
//    this list of conditions and the following disclaimer in the documentation
//    and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
//    contributors may be used to endorse or promote products derived from
//    this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

#include "src/tint/lang/spirv/reader/parser/helper_test.h"

namespace tint::spirv::reader {
namespace {

TEST_F(SpirvParserTest, ConvertFToS_ScalarSigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
        %int = OpTypeInt 32 1
      %float = OpTypeFloat 32
        %two = OpConstant %float 2
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertFToS %int %two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:i32 = spirv.convert_f_to_s<i32> 2.0f
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertFToS_ScalarUnsigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
      %float = OpTypeFloat 32
        %two = OpConstant %float 2
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertFToS %uint %two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:u32 = spirv.convert_f_to_s<u32> 2.0f
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertFToS_VectorSigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
        %int = OpTypeInt 32 1
      %v2int = OpTypeVector %int 2
      %float = OpTypeFloat 32
    %v2float = OpTypeVector %float 2
        %two = OpConstant %float 2
     %v2_two = OpConstantComposite %v2float %two %two
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertFToS %v2int %v2_two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:vec2<i32> = spirv.convert_f_to_s<i32> vec2<f32>(2.0f)
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertFToS_VectorUnsigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
     %v2uint = OpTypeVector %uint 2
      %float = OpTypeFloat 32
    %v2float = OpTypeVector %float 2
        %two = OpConstant %float 2
     %v2_two = OpConstantComposite %v2float %two %two
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertFToS %v2uint %v2_two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:vec2<u32> = spirv.convert_f_to_s<u32> vec2<f32>(2.0f)
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertFToU_ScalarUnsigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
      %float = OpTypeFloat 32
        %two = OpConstant %float 2
    %void_fn = OpTypeFunction %void
        %ptr = OpTypePointer Function %float

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %v = OpVariable %ptr Function %two
              %ld = OpLoad %float %v
               %1 = OpConvertFToU %uint %ld
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:ptr<function, f32, read_write> = var 2.0f
    %3:f32 = load %2
    %4:u32 = convert %3
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertFToU_VectorUnsigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
     %v2uint = OpTypeVector %uint 2
      %float = OpTypeFloat 32
    %v2float = OpTypeVector %float 2
        %two = OpConstant %float 2
     %v2_two = OpConstantComposite %v2float %two %two
    %void_fn = OpTypeFunction %void
        %ptr = OpTypePointer Function %v2float

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %v = OpVariable %ptr Function %v2_two
              %ld = OpLoad %v2float %v
               %1 = OpConvertFToU %v2uint %ld
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:ptr<function, vec2<f32>, read_write> = var vec2<f32>(2.0f)
    %3:vec2<f32> = load %2
    %4:vec2<u32> = convert %3
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertSToF_ScalarSigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
        %int = OpTypeInt 32 1
      %float = OpTypeFloat 32
        %two = OpConstant %int 2
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertSToF %float %two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:f32 = spirv.convert_s_to_f<f32> 2i
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertSToF_ScalarUnsigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
      %float = OpTypeFloat 32
        %two = OpConstant %uint 2
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertSToF %float %two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:f32 = spirv.convert_s_to_f<f32> 2u
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertSToF_VectorSigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
        %int = OpTypeInt 32 1
      %v2int = OpTypeVector %int 2
      %float = OpTypeFloat 32
    %v2float = OpTypeVector %float 2
        %two = OpConstant %int 2
     %v2_two = OpConstantComposite %v2int %two %two
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertSToF %v2float %v2_two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:vec2<f32> = spirv.convert_s_to_f<f32> vec2<i32>(2i)
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertSToF_VectorUnsigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
     %v2uint = OpTypeVector %uint 2
      %float = OpTypeFloat 32
    %v2float = OpTypeVector %float 2
        %two = OpConstant %uint 2
     %v2_two = OpConstantComposite %v2uint %two %two
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertSToF %v2float %v2_two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:vec2<f32> = spirv.convert_s_to_f<f32> vec2<u32>(2u)
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertUToF_ScalarSigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
        %int = OpTypeInt 32 1
      %float = OpTypeFloat 32
        %two = OpConstant %int 2
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertUToF %float %two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:f32 = spirv.convert_u_to_f<f32> 2i
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertUToF_ScalarUnsigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
      %float = OpTypeFloat 32
        %two = OpConstant %uint 2
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertUToF %float %two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:f32 = spirv.convert_u_to_f<f32> 2u
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertUToF_VectorSigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
        %int = OpTypeInt 32 1
      %v2int = OpTypeVector %int 2
      %float = OpTypeFloat 32
    %v2float = OpTypeVector %float 2
        %two = OpConstant %int 2
     %v2_two = OpConstantComposite %v2int %two %two
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertUToF %v2float %v2_two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:vec2<f32> = spirv.convert_u_to_f<f32> vec2<i32>(2i)
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, ConvertUToF_VectorUnsigned) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
     %v2uint = OpTypeVector %uint 2
      %float = OpTypeFloat 32
    %v2float = OpTypeVector %float 2
        %two = OpConstant %uint 2
     %v2_two = OpConstantComposite %v2uint %two %two
    %void_fn = OpTypeFunction %void

       %main = OpFunction %void None %void_fn
 %main_start = OpLabel
               %1 = OpConvertUToF %v2float %v2_two
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:vec2<f32> = spirv.convert_u_to_f<f32> vec2<u32>(2u)
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, FConvert_ScalarF32ToF16) {
    EXPECT_IR(R"(
             OpCapability Shader
             OpCapability Float16
             OpMemoryModel Logical GLSL450
             OpEntryPoint GLCompute %main "main"
             OpExecutionMode %main LocalSize 1 1 1
     %void    = OpTypeVoid
     %float16 = OpTypeFloat 16
     %float32 = OpTypeFloat 32
     %two     = OpConstant %float32 2
     %void_fn = OpTypeFunction %void
         %ptr = OpTypePointer Function %float32

     %main    = OpFunction %void None %void_fn
     %main_start = OpLabel
             %v = OpVariable %ptr Function %two
            %ld = OpLoad %float32 %v
             %1 = OpFConvert %float16 %ld
             OpReturn
     OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:ptr<function, f32, read_write> = var 2.0f
    %3:f32 = load %2
    %4:f16 = convert %3
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, FConvert_ScalarF16ToF32) {
    EXPECT_IR(R"(
             OpCapability Shader
             OpCapability Float16
             OpMemoryModel Logical GLSL450
             OpEntryPoint GLCompute %main "main"
             OpExecutionMode %main LocalSize 1 1 1
     %void    = OpTypeVoid
     %float16 = OpTypeFloat 16
     %float32 = OpTypeFloat 32
     %two     = OpConstant %float16 2
     %void_fn = OpTypeFunction %void
         %ptr = OpTypePointer Function %float16

     %main    = OpFunction %void None %void_fn
     %main_start = OpLabel
             %v = OpVariable %ptr Function %two
            %ld = OpLoad %float16 %v
             %1 = OpFConvert %float32 %ld
             OpReturn
     OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:ptr<function, f16, read_write> = var 2.0h
    %3:f16 = load %2
    %4:f32 = convert %3
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, FConvert_VectorF32ToF16) {
    EXPECT_IR(R"(
             OpCapability Shader
             OpCapability Float16
             OpMemoryModel Logical GLSL450
             OpEntryPoint GLCompute %main "main"
             OpExecutionMode %main LocalSize 1 1 1
     %void       = OpTypeVoid
     %float16    = OpTypeFloat 16
     %float32    = OpTypeFloat 32
     %v2float16  = OpTypeVector %float16 2
     %v2float32  = OpTypeVector %float32 2
     %two        = OpConstant %float32 2
     %v2_two     = OpConstantComposite %v2float32 %two %two
     %void_fn    = OpTypeFunction %void
            %ptr = OpTypePointer Function %v2float32

     %main       = OpFunction %void None %void_fn
     %main_start = OpLabel
             %v = OpVariable %ptr Function %v2_two
            %ld = OpLoad %v2float32 %v
             %1 = OpFConvert %v2float16 %ld
             OpReturn
     OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:ptr<function, vec2<f32>, read_write> = var vec2<f32>(2.0f)
    %3:vec2<f32> = load %2
    %4:vec2<f16> = convert %3
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, FConvert_VectorF16ToF32) {
    EXPECT_IR(R"(
             OpCapability Shader
             OpCapability Float16
             OpMemoryModel Logical GLSL450
             OpEntryPoint GLCompute %main "main"
             OpExecutionMode %main LocalSize 1 1 1
     %void       = OpTypeVoid
     %float16    = OpTypeFloat 16
     %float32    = OpTypeFloat 32
     %v2float16  = OpTypeVector %float16 2
     %v2float32  = OpTypeVector %float32 2
     %two       = OpConstant %float16 2
     %v2_two    = OpConstantComposite %v2float16 %two %two
     %void_fn    = OpTypeFunction %void
            %ptr = OpTypePointer Function %v2float16

     %main       = OpFunction %void None %void_fn
     %main_start = OpLabel
             %v = OpVariable %ptr Function %v2_two
            %ld = OpLoad %v2float16 %v
             %1 = OpFConvert %v2float32 %ld
             OpReturn
     OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:ptr<function, vec2<f16>, read_write> = var vec2<f16>(2.0h)
    %3:vec2<f16> = load %2
    %4:vec2<f32> = convert %3
    ret
  }
}
)");
}

}  // namespace
}  // namespace tint::spirv::reader