From 78c736371ca0cd5582e30872602cafbff1ec3d33 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 6 Jan 2025 18:23:08 -0800 Subject: [PATCH 1/4] add tests for scalar printf operands --- test_conformance/spirv_new/CMakeLists.txt | 1 + .../printf_operands_scalar_fp32.spvasm32 | 85 +++++++ .../printf_operands_scalar_fp32.spvasm64 | 85 +++++++ .../printf_operands_scalar_fp64.spvasm32 | 93 +++++++ .../printf_operands_scalar_fp64.spvasm64 | 93 +++++++ .../printf_operands_scalar_int32.spvasm32 | 140 +++++++++++ .../printf_operands_scalar_int32.spvasm64 | 140 +++++++++++ .../printf_operands_scalar_int64.spvasm32 | 77 ++++++ .../printf_operands_scalar_int64.spvasm64 | 77 ++++++ .../spirv_new/test_extinst_printf.cpp | 230 ++++++++++++++++++ 10 files changed, 1021 insertions(+) create mode 100644 test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp32.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp32.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp64.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp64.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int32.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int32.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int64.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int64.spvasm64 create mode 100644 test_conformance/spirv_new/test_extinst_printf.cpp diff --git a/test_conformance/spirv_new/CMakeLists.txt b/test_conformance/spirv_new/CMakeLists.txt index 67faecf895..5f3fe84255 100644 --- a/test_conformance/spirv_new/CMakeLists.txt +++ b/test_conformance/spirv_new/CMakeLists.txt @@ -5,6 +5,7 @@ set(${MODULE_NAME}_SOURCES test_basic_versions.cpp test_cl_khr_expect_assume.cpp test_decorate.cpp + test_extinst_printf.cpp test_get_program_il.cpp test_linkage.cpp test_no_integer_wrap_decoration.cpp diff --git a/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp32.spvasm32 b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp32.spvasm32 new file mode 100644 index 0000000000..58631498ad --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp32.spvasm32 @@ -0,0 +1,85 @@ +; kernel void printf_operands_scalar_fp32(float f) +; { +; printf("a = %.1a\n", f); +; printf("A = %.1A\n", f); +; printf("e = %.1e\n", f); +; printf("E = %.1E\n", f); +; printf("f = %.1f\n", f); +; printf("F = %.1F\n", f); +; printf("g = %.1g\n", f); +; printf("G = %.1G\n", f); +; } + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int8 + %clext = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %test "printf_operands_scalar_fp32" + %uchar = OpTypeInt 8 0 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 + %uint_10 = OpConstant %uint 10 + %uint_12 = OpConstant %uint 12 + %uchar_nul = OpConstant %uchar 0 + %uchar_nl = OpConstant %uchar 10 + %uchar_sp = OpConstant %uchar 32 + %uchar_pct = OpConstant %uchar 37 + %uchar_dot = OpConstant %uchar 46 + %uchar_1 = OpConstant %uchar 49 + %uchar_eq = OpConstant %uchar 61 + %uchar_A = OpConstant %uchar 65 + %uchar_E = OpConstant %uchar 69 + %uchar_F = OpConstant %uchar 70 + %uchar_G = OpConstant %uchar 71 + %uchar_a = OpConstant %uchar 97 + %uchar_e = OpConstant %uchar 101 + %uchar_f = OpConstant %uchar 102 + %uchar_g = OpConstant %uchar 103 + %string_10 = OpTypeArray %uchar %uint_10 +%cptr_string_10 = OpTypePointer UniformConstant %string_10 + %void = OpTypeVoid + %float = OpTypeFloat 32 + %kernel_sig = OpTypeFunction %void %float + %cptr_char = OpTypePointer UniformConstant %uchar + + %array_a = OpConstantComposite %string_10 %uchar_a %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_a %uchar_nl %uchar_nul ; "a = %.1a\n" + %string_a = OpVariable %cptr_string_10 UniformConstant %array_a + %array_A = OpConstantComposite %string_10 %uchar_A %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_A %uchar_nl %uchar_nul ; "A = %.1A\n" + %string_A = OpVariable %cptr_string_10 UniformConstant %array_A + %array_e = OpConstantComposite %string_10 %uchar_e %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_e %uchar_nl %uchar_nul ; "e = %.1e\n" + %string_e = OpVariable %cptr_string_10 UniformConstant %array_e + %array_E = OpConstantComposite %string_10 %uchar_E %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_E %uchar_nl %uchar_nul ; "E = %.1E\n" + %string_E = OpVariable %cptr_string_10 UniformConstant %array_E + %array_f = OpConstantComposite %string_10 %uchar_f %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_f %uchar_nl %uchar_nul ; "f = %.1f\n" + %string_f = OpVariable %cptr_string_10 UniformConstant %array_f + %array_F = OpConstantComposite %string_10 %uchar_F %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_F %uchar_nl %uchar_nul ; "F = %.1F\n" + %string_F = OpVariable %cptr_string_10 UniformConstant %array_F + %array_g = OpConstantComposite %string_10 %uchar_g %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_g %uchar_nl %uchar_nul ; "g = %.1g\n" + %string_g = OpVariable %cptr_string_10 UniformConstant %array_g + %array_G = OpConstantComposite %string_10 %uchar_G %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_G %uchar_nl %uchar_nul ; "G = %.1G\n" + %string_G = OpVariable %cptr_string_10 UniformConstant %array_G + + %test = OpFunction %void None %kernel_sig + %f = OpFunctionParameter %float + %entry = OpLabel + + %fmt_a = OpBitcast %cptr_char %string_a + %printf_a = OpExtInst %uint %clext printf %fmt_a %f + %fmt_A = OpBitcast %cptr_char %string_A + %printf_A = OpExtInst %uint %clext printf %fmt_A %f + %fmt_e = OpBitcast %cptr_char %string_e + %printf_e = OpExtInst %uint %clext printf %fmt_e %f + %fmt_E = OpBitcast %cptr_char %string_E + %printf_E = OpExtInst %uint %clext printf %fmt_E %f + %fmt_f = OpBitcast %cptr_char %string_f + %printf_f = OpExtInst %uint %clext printf %fmt_f %f + %fmt_F = OpBitcast %cptr_char %string_F + %printf_F = OpExtInst %uint %clext printf %fmt_F %f + %fmt_g = OpBitcast %cptr_char %string_g + %printf_g = OpExtInst %uint %clext printf %fmt_g %f + %fmt_G = OpBitcast %cptr_char %string_G + %printf_G = OpExtInst %uint %clext printf %fmt_G %f + + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp32.spvasm64 b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp32.spvasm64 new file mode 100644 index 0000000000..ba415d0986 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp32.spvasm64 @@ -0,0 +1,85 @@ +; kernel void printf_operands_scalar_fp32(float f) +; { +; printf("a = %.1a\n", f); +; printf("A = %.1A\n", f); +; printf("e = %.1e\n", f); +; printf("E = %.1E\n", f); +; printf("f = %.1f\n", f); +; printf("F = %.1F\n", f); +; printf("g = %.1g\n", f); +; printf("G = %.1G\n", f); +; } + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int8 + %clext = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %test "printf_operands_scalar_fp32" + %uchar = OpTypeInt 8 0 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 + %uint_10 = OpConstant %uint 10 + %uint_12 = OpConstant %uint 12 + %uchar_nul = OpConstant %uchar 0 + %uchar_nl = OpConstant %uchar 10 + %uchar_sp = OpConstant %uchar 32 + %uchar_pct = OpConstant %uchar 37 + %uchar_dot = OpConstant %uchar 46 + %uchar_1 = OpConstant %uchar 49 + %uchar_eq = OpConstant %uchar 61 + %uchar_A = OpConstant %uchar 65 + %uchar_E = OpConstant %uchar 69 + %uchar_F = OpConstant %uchar 70 + %uchar_G = OpConstant %uchar 71 + %uchar_a = OpConstant %uchar 97 + %uchar_e = OpConstant %uchar 101 + %uchar_f = OpConstant %uchar 102 + %uchar_g = OpConstant %uchar 103 + %string_10 = OpTypeArray %uchar %uint_10 +%cptr_string_10 = OpTypePointer UniformConstant %string_10 + %void = OpTypeVoid + %float = OpTypeFloat 32 + %kernel_sig = OpTypeFunction %void %float + %cptr_char = OpTypePointer UniformConstant %uchar + + %array_a = OpConstantComposite %string_10 %uchar_a %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_a %uchar_nl %uchar_nul ; "a = %.1a\n" + %string_a = OpVariable %cptr_string_10 UniformConstant %array_a + %array_A = OpConstantComposite %string_10 %uchar_A %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_A %uchar_nl %uchar_nul ; "A = %.1A\n" + %string_A = OpVariable %cptr_string_10 UniformConstant %array_A + %array_e = OpConstantComposite %string_10 %uchar_e %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_e %uchar_nl %uchar_nul ; "e = %.1e\n" + %string_e = OpVariable %cptr_string_10 UniformConstant %array_e + %array_E = OpConstantComposite %string_10 %uchar_E %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_E %uchar_nl %uchar_nul ; "E = %.1E\n" + %string_E = OpVariable %cptr_string_10 UniformConstant %array_E + %array_f = OpConstantComposite %string_10 %uchar_f %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_f %uchar_nl %uchar_nul ; "f = %.1f\n" + %string_f = OpVariable %cptr_string_10 UniformConstant %array_f + %array_F = OpConstantComposite %string_10 %uchar_F %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_F %uchar_nl %uchar_nul ; "F = %.1F\n" + %string_F = OpVariable %cptr_string_10 UniformConstant %array_F + %array_g = OpConstantComposite %string_10 %uchar_g %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_g %uchar_nl %uchar_nul ; "g = %.1g\n" + %string_g = OpVariable %cptr_string_10 UniformConstant %array_g + %array_G = OpConstantComposite %string_10 %uchar_G %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_G %uchar_nl %uchar_nul ; "G = %.1G\n" + %string_G = OpVariable %cptr_string_10 UniformConstant %array_G + + %test = OpFunction %void None %kernel_sig + %f = OpFunctionParameter %float + %entry = OpLabel + + %fmt_a = OpBitcast %cptr_char %string_a + %printf_a = OpExtInst %uint %clext printf %fmt_a %f + %fmt_A = OpBitcast %cptr_char %string_A + %printf_A = OpExtInst %uint %clext printf %fmt_A %f + %fmt_e = OpBitcast %cptr_char %string_e + %printf_e = OpExtInst %uint %clext printf %fmt_e %f + %fmt_E = OpBitcast %cptr_char %string_E + %printf_E = OpExtInst %uint %clext printf %fmt_E %f + %fmt_f = OpBitcast %cptr_char %string_f + %printf_f = OpExtInst %uint %clext printf %fmt_f %f + %fmt_F = OpBitcast %cptr_char %string_F + %printf_F = OpExtInst %uint %clext printf %fmt_F %f + %fmt_g = OpBitcast %cptr_char %string_g + %printf_g = OpExtInst %uint %clext printf %fmt_g %f + %fmt_G = OpBitcast %cptr_char %string_G + %printf_G = OpExtInst %uint %clext printf %fmt_G %f + + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp64.spvasm32 b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp64.spvasm32 new file mode 100644 index 0000000000..1b31cf49dd --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp64.spvasm32 @@ -0,0 +1,93 @@ +; kernel void printf_operands_scalar_fp64(double d) +; { +; printf("a = %.1a\n", d); +; printf("A = %.1A\n", d); +; printf("e = %.1e\n", d); +; printf("E = %.1E\n", d); +; printf("f = %.1f\n", d); +; printf("F = %.1F\n", d); +; printf("g = %.1g\n", d); +; printf("G = %.1G\n", d); +; } + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Float64 + OpCapability Int8 + %clext = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %test "printf_operands_scalar_fp64" + %uchar = OpTypeInt 8 0 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 + %uint_10 = OpConstant %uint 10 + %uint_12 = OpConstant %uint 12 + %uchar_nul = OpConstant %uchar 0 + %uchar_nl = OpConstant %uchar 10 + %uchar_sp = OpConstant %uchar 32 + %uchar_pct = OpConstant %uchar 37 + %uchar_dot = OpConstant %uchar 46 + %uchar_1 = OpConstant %uchar 49 + %uchar_eq = OpConstant %uchar 61 + %uchar_A = OpConstant %uchar 65 + %uchar_E = OpConstant %uchar 69 + %uchar_F = OpConstant %uchar 70 + %uchar_G = OpConstant %uchar 71 + %uchar_X = OpConstant %uchar 88 + %uchar_a = OpConstant %uchar 97 + %uchar_d = OpConstant %uchar 100 + %uchar_e = OpConstant %uchar 101 + %uchar_f = OpConstant %uchar 102 + %uchar_g = OpConstant %uchar 103 + %uchar_h = OpConstant %uchar 104 + %uchar_i = OpConstant %uchar 105 + %uchar_o = OpConstant %uchar 111 + %uchar_u = OpConstant %uchar 117 + %uchar_x = OpConstant %uchar 120 + %string_10 = OpTypeArray %uchar %uint_10 +%cptr_string_10 = OpTypePointer UniformConstant %string_10 + %void = OpTypeVoid + %double = OpTypeFloat 64 + %kernel_sig = OpTypeFunction %void %double + %cptr_char = OpTypePointer UniformConstant %uchar + + %array_a = OpConstantComposite %string_10 %uchar_a %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_a %uchar_nl %uchar_nul ; "a = %.1a\n" + %string_a = OpVariable %cptr_string_10 UniformConstant %array_a + %array_A = OpConstantComposite %string_10 %uchar_A %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_A %uchar_nl %uchar_nul ; "A = %.1A\n" + %string_A = OpVariable %cptr_string_10 UniformConstant %array_A + %array_e = OpConstantComposite %string_10 %uchar_e %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_e %uchar_nl %uchar_nul ; "e = %.1e\n" + %string_e = OpVariable %cptr_string_10 UniformConstant %array_e + %array_E = OpConstantComposite %string_10 %uchar_E %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_E %uchar_nl %uchar_nul ; "E = %.1E\n" + %string_E = OpVariable %cptr_string_10 UniformConstant %array_E + %array_f = OpConstantComposite %string_10 %uchar_f %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_f %uchar_nl %uchar_nul ; "f = %.1f\n" + %string_f = OpVariable %cptr_string_10 UniformConstant %array_f + %array_F = OpConstantComposite %string_10 %uchar_F %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_F %uchar_nl %uchar_nul ; "F = %.1F\n" + %string_F = OpVariable %cptr_string_10 UniformConstant %array_F + %array_g = OpConstantComposite %string_10 %uchar_g %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_g %uchar_nl %uchar_nul ; "g = %.1g\n" + %string_g = OpVariable %cptr_string_10 UniformConstant %array_g + %array_G = OpConstantComposite %string_10 %uchar_G %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_G %uchar_nl %uchar_nul ; "G = %.1G\n" + %string_G = OpVariable %cptr_string_10 UniformConstant %array_G + + %test = OpFunction %void None %kernel_sig + %d = OpFunctionParameter %double + %entry = OpLabel + + %fmt_a = OpBitcast %cptr_char %string_a + %printf_a = OpExtInst %uint %clext printf %fmt_a %d + %fmt_A = OpBitcast %cptr_char %string_A + %printf_A = OpExtInst %uint %clext printf %fmt_A %d + %fmt_e = OpBitcast %cptr_char %string_e + %printf_e = OpExtInst %uint %clext printf %fmt_e %d + %fmt_E = OpBitcast %cptr_char %string_E + %printf_E = OpExtInst %uint %clext printf %fmt_E %d + %fmt_f = OpBitcast %cptr_char %string_f + %printf_f = OpExtInst %uint %clext printf %fmt_f %d + %fmt_F = OpBitcast %cptr_char %string_F + %printf_F = OpExtInst %uint %clext printf %fmt_F %d + %fmt_g = OpBitcast %cptr_char %string_g + %printf_g = OpExtInst %uint %clext printf %fmt_g %d + %fmt_G = OpBitcast %cptr_char %string_G + %printf_G = OpExtInst %uint %clext printf %fmt_G %d + + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp64.spvasm64 b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp64.spvasm64 new file mode 100644 index 0000000000..a947e5ec88 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_fp64.spvasm64 @@ -0,0 +1,93 @@ +; kernel void printf_operands_scalar_fp64(double d) +; { +; printf("a = %.1a\n", d); +; printf("A = %.1A\n", d); +; printf("e = %.1e\n", d); +; printf("E = %.1E\n", d); +; printf("f = %.1f\n", d); +; printf("F = %.1F\n", d); +; printf("g = %.1g\n", d); +; printf("G = %.1G\n", d); +; } + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Float64 + OpCapability Int8 + %clext = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %test "printf_operands_scalar_fp64" + %uchar = OpTypeInt 8 0 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 + %uint_10 = OpConstant %uint 10 + %uint_12 = OpConstant %uint 12 + %uchar_nul = OpConstant %uchar 0 + %uchar_nl = OpConstant %uchar 10 + %uchar_sp = OpConstant %uchar 32 + %uchar_pct = OpConstant %uchar 37 + %uchar_dot = OpConstant %uchar 46 + %uchar_1 = OpConstant %uchar 49 + %uchar_eq = OpConstant %uchar 61 + %uchar_A = OpConstant %uchar 65 + %uchar_E = OpConstant %uchar 69 + %uchar_F = OpConstant %uchar 70 + %uchar_G = OpConstant %uchar 71 + %uchar_X = OpConstant %uchar 88 + %uchar_a = OpConstant %uchar 97 + %uchar_d = OpConstant %uchar 100 + %uchar_e = OpConstant %uchar 101 + %uchar_f = OpConstant %uchar 102 + %uchar_g = OpConstant %uchar 103 + %uchar_h = OpConstant %uchar 104 + %uchar_i = OpConstant %uchar 105 + %uchar_o = OpConstant %uchar 111 + %uchar_u = OpConstant %uchar 117 + %uchar_x = OpConstant %uchar 120 + %string_10 = OpTypeArray %uchar %uint_10 +%cptr_string_10 = OpTypePointer UniformConstant %string_10 + %void = OpTypeVoid + %double = OpTypeFloat 64 + %kernel_sig = OpTypeFunction %void %double + %cptr_char = OpTypePointer UniformConstant %uchar + + %array_a = OpConstantComposite %string_10 %uchar_a %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_a %uchar_nl %uchar_nul ; "a = %.1a\n" + %string_a = OpVariable %cptr_string_10 UniformConstant %array_a + %array_A = OpConstantComposite %string_10 %uchar_A %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_A %uchar_nl %uchar_nul ; "A = %.1A\n" + %string_A = OpVariable %cptr_string_10 UniformConstant %array_A + %array_e = OpConstantComposite %string_10 %uchar_e %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_e %uchar_nl %uchar_nul ; "e = %.1e\n" + %string_e = OpVariable %cptr_string_10 UniformConstant %array_e + %array_E = OpConstantComposite %string_10 %uchar_E %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_E %uchar_nl %uchar_nul ; "E = %.1E\n" + %string_E = OpVariable %cptr_string_10 UniformConstant %array_E + %array_f = OpConstantComposite %string_10 %uchar_f %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_f %uchar_nl %uchar_nul ; "f = %.1f\n" + %string_f = OpVariable %cptr_string_10 UniformConstant %array_f + %array_F = OpConstantComposite %string_10 %uchar_F %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_F %uchar_nl %uchar_nul ; "F = %.1F\n" + %string_F = OpVariable %cptr_string_10 UniformConstant %array_F + %array_g = OpConstantComposite %string_10 %uchar_g %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_g %uchar_nl %uchar_nul ; "g = %.1g\n" + %string_g = OpVariable %cptr_string_10 UniformConstant %array_g + %array_G = OpConstantComposite %string_10 %uchar_G %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_dot %uchar_1 %uchar_G %uchar_nl %uchar_nul ; "G = %.1G\n" + %string_G = OpVariable %cptr_string_10 UniformConstant %array_G + + %test = OpFunction %void None %kernel_sig + %d = OpFunctionParameter %double + %entry = OpLabel + + %fmt_a = OpBitcast %cptr_char %string_a + %printf_a = OpExtInst %uint %clext printf %fmt_a %d + %fmt_A = OpBitcast %cptr_char %string_A + %printf_A = OpExtInst %uint %clext printf %fmt_A %d + %fmt_e = OpBitcast %cptr_char %string_e + %printf_e = OpExtInst %uint %clext printf %fmt_e %d + %fmt_E = OpBitcast %cptr_char %string_E + %printf_E = OpExtInst %uint %clext printf %fmt_E %d + %fmt_f = OpBitcast %cptr_char %string_f + %printf_f = OpExtInst %uint %clext printf %fmt_f %d + %fmt_F = OpBitcast %cptr_char %string_F + %printf_F = OpExtInst %uint %clext printf %fmt_F %d + %fmt_g = OpBitcast %cptr_char %string_g + %printf_g = OpExtInst %uint %clext printf %fmt_g %d + %fmt_G = OpBitcast %cptr_char %string_G + %printf_G = OpExtInst %uint %clext printf %fmt_G %d + + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int32.spvasm32 b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int32.spvasm32 new file mode 100644 index 0000000000..61fb8cd02f --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int32.spvasm32 @@ -0,0 +1,140 @@ +; kernel void printf_operands_scalar_int32(int i) +; { +; printf("d = %d\n", i); +; printf("i = %i\n", i); +; printf("o = %o\n", i); +; printf("u = %u\n", i); +; printf("x = %x\n", i); +; printf("X = %X\n", i); +; +; printf("hd = %hd\n", i); +; printf("hi = %hi\n", i); +; printf("ho = %ho\n", i); +; printf("hu = %hu\n", i); +; printf("hx = %hx\n", i); +; printf("hX = %hX\n", i); +; +; printf("hhd = %hhd\n", i); +; printf("hhi = %hhi\n", i); +; printf("hho = %hho\n", i); +; printf("hhu = %hhu\n", i); +; printf("hhx = %hhx\n", i); +; printf("hhX = %hhX\n", i); +; } + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int8 + %clext = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %test "printf_operands_scalar_int32" + %uchar = OpTypeInt 8 0 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 + %uint_10 = OpConstant %uint 10 + %uint_12 = OpConstant %uint 12 + %uchar_nul = OpConstant %uchar 0 + %uchar_nl = OpConstant %uchar 10 + %uchar_sp = OpConstant %uchar 32 + %uchar_pct = OpConstant %uchar 37 + %uchar_eq = OpConstant %uchar 61 + %uchar_X = OpConstant %uchar 88 + %uchar_d = OpConstant %uchar 100 + %uchar_h = OpConstant %uchar 104 + %uchar_i = OpConstant %uchar 105 + %uchar_o = OpConstant %uchar 111 + %uchar_u = OpConstant %uchar 117 + %uchar_x = OpConstant %uchar 120 + %string_8 = OpTypeArray %uchar %uint_8 +%cptr_string_8 = OpTypePointer UniformConstant %string_8 + %string_10 = OpTypeArray %uchar %uint_10 +%cptr_string_10 = OpTypePointer UniformConstant %string_10 + %string_12 = OpTypeArray %uchar %uint_12 +%cptr_string_12 = OpTypePointer UniformConstant %string_12 + %void = OpTypeVoid + %kernel_sig = OpTypeFunction %void %uint + %cptr_char = OpTypePointer UniformConstant %uchar + + %array_d = OpConstantComposite %string_8 %uchar_d %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_d %uchar_nl %uchar_nul ; "d = %d\n" + %string_d = OpVariable %cptr_string_8 UniformConstant %array_d + %array_i = OpConstantComposite %string_8 %uchar_i %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_i %uchar_nl %uchar_nul ; "i = %i\n" + %string_i = OpVariable %cptr_string_8 UniformConstant %array_i + %array_o = OpConstantComposite %string_8 %uchar_o %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_o %uchar_nl %uchar_nul ; "o = %o\n" + %string_o = OpVariable %cptr_string_8 UniformConstant %array_o + %array_u = OpConstantComposite %string_8 %uchar_u %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_u %uchar_nl %uchar_nul ; "u = %u\n" + %string_u = OpVariable %cptr_string_8 UniformConstant %array_u + %array_x = OpConstantComposite %string_8 %uchar_x %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_x %uchar_nl %uchar_nul ; "x = %x\n" + %string_x = OpVariable %cptr_string_8 UniformConstant %array_x + %array_X = OpConstantComposite %string_8 %uchar_X %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_X %uchar_nl %uchar_nul ; "X = %X\n" + %string_X = OpVariable %cptr_string_8 UniformConstant %array_X + + %array_hd = OpConstantComposite %string_10 %uchar_h %uchar_d %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_d %uchar_nl %uchar_nul ; "hd = %hd\n" + %string_hd = OpVariable %cptr_string_10 UniformConstant %array_hd + %array_hi = OpConstantComposite %string_10 %uchar_h %uchar_i %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_i %uchar_nl %uchar_nul ; "hi = %hi\n" + %string_hi = OpVariable %cptr_string_10 UniformConstant %array_hi + %array_ho = OpConstantComposite %string_10 %uchar_h %uchar_o %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_o %uchar_nl %uchar_nul ; "ho = %ho\n" + %string_ho = OpVariable %cptr_string_10 UniformConstant %array_ho + %array_hu = OpConstantComposite %string_10 %uchar_h %uchar_u %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_u %uchar_nl %uchar_nul ; "hu = %hu\n" + %string_hu = OpVariable %cptr_string_10 UniformConstant %array_hu + %array_hx = OpConstantComposite %string_10 %uchar_h %uchar_x %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_x %uchar_nl %uchar_nul ; "hx = %hx\n" + %string_hx = OpVariable %cptr_string_10 UniformConstant %array_hx + %array_hX = OpConstantComposite %string_10 %uchar_h %uchar_X %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_X %uchar_nl %uchar_nul ; "hX = %hX\n" + %string_hX = OpVariable %cptr_string_10 UniformConstant %array_hX + + %array_hhd = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_d %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_d %uchar_nl %uchar_nul ; "hhd = %hhd\n" + %string_hhd = OpVariable %cptr_string_12 UniformConstant %array_hhd + %array_hhi = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_i %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_i %uchar_nl %uchar_nul ; "hhi = %hhi\n" + %string_hhi = OpVariable %cptr_string_12 UniformConstant %array_hhi + %array_hho = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_o %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_o %uchar_nl %uchar_nul ; "hho = %hho\n" + %string_hho = OpVariable %cptr_string_12 UniformConstant %array_hho + %array_hhu = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_u %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_u %uchar_nl %uchar_nul ; "hhu = %hhu\n" + %string_hhu = OpVariable %cptr_string_12 UniformConstant %array_hhu + %array_hhx = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_x %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_x %uchar_nl %uchar_nul ; "hhx = %hhx\n" + %string_hhx = OpVariable %cptr_string_12 UniformConstant %array_hhx + %array_hhX = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_X %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_X %uchar_nl %uchar_nul ; "hhX = %hhX\n" + %string_hhX = OpVariable %cptr_string_12 UniformConstant %array_hhX + + %test = OpFunction %void None %kernel_sig + %i = OpFunctionParameter %uint + %entry = OpLabel + %fmt_d = OpBitcast %cptr_char %string_d + %printf_d = OpExtInst %uint %clext printf %fmt_d %i + %fmt_i = OpBitcast %cptr_char %string_i + %printf_i = OpExtInst %uint %clext printf %fmt_i %i + %fmt_o = OpBitcast %cptr_char %string_o + %printf_o = OpExtInst %uint %clext printf %fmt_o %i + %fmt_u = OpBitcast %cptr_char %string_u + %printf_u = OpExtInst %uint %clext printf %fmt_u %i + %fmt_x = OpBitcast %cptr_char %string_x + %printf_x = OpExtInst %uint %clext printf %fmt_x %i + %fmt_X = OpBitcast %cptr_char %string_X + %printf_X = OpExtInst %uint %clext printf %fmt_X %i + + %fmt_hd = OpBitcast %cptr_char %string_hd + %printf_hd = OpExtInst %uint %clext printf %fmt_hd %i + %fmt_hi = OpBitcast %cptr_char %string_hi + %printf_hi = OpExtInst %uint %clext printf %fmt_hi %i + %fmt_ho = OpBitcast %cptr_char %string_ho + %printf_ho = OpExtInst %uint %clext printf %fmt_ho %i + %fmt_hu = OpBitcast %cptr_char %string_hu + %printf_hu = OpExtInst %uint %clext printf %fmt_hu %i + %fmt_hx = OpBitcast %cptr_char %string_hx + %printf_hx = OpExtInst %uint %clext printf %fmt_hx %i + %fmt_hX = OpBitcast %cptr_char %string_hX + %printf_hX = OpExtInst %uint %clext printf %fmt_hX %i + + %fmt_hhd = OpBitcast %cptr_char %string_hhd + %printf_hhd = OpExtInst %uint %clext printf %fmt_hhd %i + %fmt_hhi = OpBitcast %cptr_char %string_hhi + %printf_hhi = OpExtInst %uint %clext printf %fmt_hhi %i + %fmt_hho = OpBitcast %cptr_char %string_hho + %printf_hho = OpExtInst %uint %clext printf %fmt_hho %i + %fmt_hhu = OpBitcast %cptr_char %string_hhu + %printf_hhu = OpExtInst %uint %clext printf %fmt_hhu %i + %fmt_hhx = OpBitcast %cptr_char %string_hhx + %printf_hhx = OpExtInst %uint %clext printf %fmt_hhx %i + %fmt_hhX = OpBitcast %cptr_char %string_hhX + %printf_hhX = OpExtInst %uint %clext printf %fmt_hhX %i + + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int32.spvasm64 b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int32.spvasm64 new file mode 100644 index 0000000000..91ad8e1e96 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int32.spvasm64 @@ -0,0 +1,140 @@ +; kernel void printf_operands_scalar_int32(int i) +; { +; printf("d = %d\n", i); +; printf("i = %i\n", i); +; printf("o = %o\n", i); +; printf("u = %u\n", i); +; printf("x = %x\n", i); +; printf("X = %X\n", i); +; +; printf("hd = %hd\n", i); +; printf("hi = %hi\n", i); +; printf("ho = %ho\n", i); +; printf("hu = %hu\n", i); +; printf("hx = %hx\n", i); +; printf("hX = %hX\n", i); +; +; printf("hhd = %hhd\n", i); +; printf("hhi = %hhi\n", i); +; printf("hho = %hho\n", i); +; printf("hhu = %hhu\n", i); +; printf("hhx = %hhx\n", i); +; printf("hhX = %hhX\n", i); +; } + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int8 + %clext = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %test "printf_operands_scalar_int32" + %uchar = OpTypeInt 8 0 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 + %uint_10 = OpConstant %uint 10 + %uint_12 = OpConstant %uint 12 + %uchar_nul = OpConstant %uchar 0 + %uchar_nl = OpConstant %uchar 10 + %uchar_sp = OpConstant %uchar 32 + %uchar_pct = OpConstant %uchar 37 + %uchar_eq = OpConstant %uchar 61 + %uchar_X = OpConstant %uchar 88 + %uchar_d = OpConstant %uchar 100 + %uchar_h = OpConstant %uchar 104 + %uchar_i = OpConstant %uchar 105 + %uchar_o = OpConstant %uchar 111 + %uchar_u = OpConstant %uchar 117 + %uchar_x = OpConstant %uchar 120 + %string_8 = OpTypeArray %uchar %uint_8 +%cptr_string_8 = OpTypePointer UniformConstant %string_8 + %string_10 = OpTypeArray %uchar %uint_10 +%cptr_string_10 = OpTypePointer UniformConstant %string_10 + %string_12 = OpTypeArray %uchar %uint_12 +%cptr_string_12 = OpTypePointer UniformConstant %string_12 + %void = OpTypeVoid + %kernel_sig = OpTypeFunction %void %uint + %cptr_char = OpTypePointer UniformConstant %uchar + + %array_d = OpConstantComposite %string_8 %uchar_d %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_d %uchar_nl %uchar_nul ; "d = %d\n" + %string_d = OpVariable %cptr_string_8 UniformConstant %array_d + %array_i = OpConstantComposite %string_8 %uchar_i %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_i %uchar_nl %uchar_nul ; "i = %i\n" + %string_i = OpVariable %cptr_string_8 UniformConstant %array_i + %array_o = OpConstantComposite %string_8 %uchar_o %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_o %uchar_nl %uchar_nul ; "o = %o\n" + %string_o = OpVariable %cptr_string_8 UniformConstant %array_o + %array_u = OpConstantComposite %string_8 %uchar_u %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_u %uchar_nl %uchar_nul ; "u = %u\n" + %string_u = OpVariable %cptr_string_8 UniformConstant %array_u + %array_x = OpConstantComposite %string_8 %uchar_x %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_x %uchar_nl %uchar_nul ; "x = %x\n" + %string_x = OpVariable %cptr_string_8 UniformConstant %array_x + %array_X = OpConstantComposite %string_8 %uchar_X %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_X %uchar_nl %uchar_nul ; "X = %X\n" + %string_X = OpVariable %cptr_string_8 UniformConstant %array_X + + %array_hd = OpConstantComposite %string_10 %uchar_h %uchar_d %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_d %uchar_nl %uchar_nul ; "hd = %hd\n" + %string_hd = OpVariable %cptr_string_10 UniformConstant %array_hd + %array_hi = OpConstantComposite %string_10 %uchar_h %uchar_i %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_i %uchar_nl %uchar_nul ; "hi = %hi\n" + %string_hi = OpVariable %cptr_string_10 UniformConstant %array_hi + %array_ho = OpConstantComposite %string_10 %uchar_h %uchar_o %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_o %uchar_nl %uchar_nul ; "ho = %ho\n" + %string_ho = OpVariable %cptr_string_10 UniformConstant %array_ho + %array_hu = OpConstantComposite %string_10 %uchar_h %uchar_u %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_u %uchar_nl %uchar_nul ; "hu = %hu\n" + %string_hu = OpVariable %cptr_string_10 UniformConstant %array_hu + %array_hx = OpConstantComposite %string_10 %uchar_h %uchar_x %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_x %uchar_nl %uchar_nul ; "hx = %hx\n" + %string_hx = OpVariable %cptr_string_10 UniformConstant %array_hx + %array_hX = OpConstantComposite %string_10 %uchar_h %uchar_X %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_X %uchar_nl %uchar_nul ; "hX = %hX\n" + %string_hX = OpVariable %cptr_string_10 UniformConstant %array_hX + + %array_hhd = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_d %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_d %uchar_nl %uchar_nul ; "hhd = %hhd\n" + %string_hhd = OpVariable %cptr_string_12 UniformConstant %array_hhd + %array_hhi = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_i %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_i %uchar_nl %uchar_nul ; "hhi = %hhi\n" + %string_hhi = OpVariable %cptr_string_12 UniformConstant %array_hhi + %array_hho = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_o %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_o %uchar_nl %uchar_nul ; "hho = %hho\n" + %string_hho = OpVariable %cptr_string_12 UniformConstant %array_hho + %array_hhu = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_u %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_u %uchar_nl %uchar_nul ; "hhu = %hhu\n" + %string_hhu = OpVariable %cptr_string_12 UniformConstant %array_hhu + %array_hhx = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_x %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_x %uchar_nl %uchar_nul ; "hhx = %hhx\n" + %string_hhx = OpVariable %cptr_string_12 UniformConstant %array_hhx + %array_hhX = OpConstantComposite %string_12 %uchar_h %uchar_h %uchar_X %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_h %uchar_h %uchar_X %uchar_nl %uchar_nul ; "hhX = %hhX\n" + %string_hhX = OpVariable %cptr_string_12 UniformConstant %array_hhX + + %test = OpFunction %void None %kernel_sig + %i = OpFunctionParameter %uint + %entry = OpLabel + %fmt_d = OpBitcast %cptr_char %string_d + %printf_d = OpExtInst %uint %clext printf %fmt_d %i + %fmt_i = OpBitcast %cptr_char %string_i + %printf_i = OpExtInst %uint %clext printf %fmt_i %i + %fmt_o = OpBitcast %cptr_char %string_o + %printf_o = OpExtInst %uint %clext printf %fmt_o %i + %fmt_u = OpBitcast %cptr_char %string_u + %printf_u = OpExtInst %uint %clext printf %fmt_u %i + %fmt_x = OpBitcast %cptr_char %string_x + %printf_x = OpExtInst %uint %clext printf %fmt_x %i + %fmt_X = OpBitcast %cptr_char %string_X + %printf_X = OpExtInst %uint %clext printf %fmt_X %i + + %fmt_hd = OpBitcast %cptr_char %string_hd + %printf_hd = OpExtInst %uint %clext printf %fmt_hd %i + %fmt_hi = OpBitcast %cptr_char %string_hi + %printf_hi = OpExtInst %uint %clext printf %fmt_hi %i + %fmt_ho = OpBitcast %cptr_char %string_ho + %printf_ho = OpExtInst %uint %clext printf %fmt_ho %i + %fmt_hu = OpBitcast %cptr_char %string_hu + %printf_hu = OpExtInst %uint %clext printf %fmt_hu %i + %fmt_hx = OpBitcast %cptr_char %string_hx + %printf_hx = OpExtInst %uint %clext printf %fmt_hx %i + %fmt_hX = OpBitcast %cptr_char %string_hX + %printf_hX = OpExtInst %uint %clext printf %fmt_hX %i + + %fmt_hhd = OpBitcast %cptr_char %string_hhd + %printf_hhd = OpExtInst %uint %clext printf %fmt_hhd %i + %fmt_hhi = OpBitcast %cptr_char %string_hhi + %printf_hhi = OpExtInst %uint %clext printf %fmt_hhi %i + %fmt_hho = OpBitcast %cptr_char %string_hho + %printf_hho = OpExtInst %uint %clext printf %fmt_hho %i + %fmt_hhu = OpBitcast %cptr_char %string_hhu + %printf_hhu = OpExtInst %uint %clext printf %fmt_hhu %i + %fmt_hhx = OpBitcast %cptr_char %string_hhx + %printf_hhx = OpExtInst %uint %clext printf %fmt_hhx %i + %fmt_hhX = OpBitcast %cptr_char %string_hhX + %printf_hhX = OpExtInst %uint %clext printf %fmt_hhX %i + + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int64.spvasm32 b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int64.spvasm32 new file mode 100644 index 0000000000..ec19e9f8a6 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int64.spvasm32 @@ -0,0 +1,77 @@ +; kernel void printf_operands_scalar_int64(long l) +; { +; printf("ld = %ld\n", l); +; printf("li = %li\n", l); +; printf("lo = %lo\n", l); +; printf("lu = %lu\n", l); +; printf("lx = %lx\n", l); +; printf("lX = %lX\n", l); +; } + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpCapability Int8 + %clext = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %test "printf_operands_scalar_int64" + %uchar = OpTypeInt 8 0 + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 + %uint_10 = OpConstant %uint 10 + %uint_12 = OpConstant %uint 12 + %uchar_nul = OpConstant %uchar 0 + %uchar_nl = OpConstant %uchar 10 + %uchar_sp = OpConstant %uchar 32 + %uchar_pct = OpConstant %uchar 37 + %uchar_eq = OpConstant %uchar 61 + %uchar_X = OpConstant %uchar 88 + %uchar_d = OpConstant %uchar 100 + %uchar_i = OpConstant %uchar 105 + %uchar_l = OpConstant %uchar 108 + %uchar_o = OpConstant %uchar 111 + %uchar_u = OpConstant %uchar 117 + %uchar_x = OpConstant %uchar 120 + %string_8 = OpTypeArray %uchar %uint_8 +%cptr_string_8 = OpTypePointer UniformConstant %string_8 + %string_10 = OpTypeArray %uchar %uint_10 +%cptr_string_10 = OpTypePointer UniformConstant %string_10 + %string_12 = OpTypeArray %uchar %uint_12 +%cptr_string_12 = OpTypePointer UniformConstant %string_12 + %void = OpTypeVoid + %float = OpTypeFloat 32 + %kernel_sig = OpTypeFunction %void %ulong + %cptr_char = OpTypePointer UniformConstant %uchar + + %array_ld = OpConstantComposite %string_10 %uchar_l %uchar_d %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_d %uchar_nl %uchar_nul ; "ld = %ld\n" + %string_ld = OpVariable %cptr_string_10 UniformConstant %array_ld + %array_li = OpConstantComposite %string_10 %uchar_l %uchar_i %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_i %uchar_nl %uchar_nul ; "li = %li\n" + %string_li = OpVariable %cptr_string_10 UniformConstant %array_li + %array_lo = OpConstantComposite %string_10 %uchar_l %uchar_o %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_o %uchar_nl %uchar_nul ; "lo = %lo\n" + %string_lo = OpVariable %cptr_string_10 UniformConstant %array_lo + %array_lu = OpConstantComposite %string_10 %uchar_l %uchar_u %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_u %uchar_nl %uchar_nul ; "lu = %lu\n" + %string_lu = OpVariable %cptr_string_10 UniformConstant %array_lu + %array_lx = OpConstantComposite %string_10 %uchar_l %uchar_x %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_x %uchar_nl %uchar_nul ; "lx = %lx\n" + %string_lx = OpVariable %cptr_string_10 UniformConstant %array_lx + %array_lX = OpConstantComposite %string_10 %uchar_l %uchar_X %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_X %uchar_nl %uchar_nul ; "lX = %lX\n" + %string_lX = OpVariable %cptr_string_10 UniformConstant %array_lX + + %test = OpFunction %void None %kernel_sig + %l = OpFunctionParameter %ulong + %entry = OpLabel + + %fmt_ld = OpBitcast %cptr_char %string_ld + %printf_ld = OpExtInst %uint %clext printf %fmt_ld %l + %fmt_li = OpBitcast %cptr_char %string_li + %printf_li = OpExtInst %uint %clext printf %fmt_li %l + %fmt_lo = OpBitcast %cptr_char %string_lo + %printf_lo = OpExtInst %uint %clext printf %fmt_lo %l + %fmt_lu = OpBitcast %cptr_char %string_lu + %printf_lu = OpExtInst %uint %clext printf %fmt_lu %l + %fmt_lx = OpBitcast %cptr_char %string_lx + %printf_lx = OpExtInst %uint %clext printf %fmt_lx %l + %fmt_lX = OpBitcast %cptr_char %string_lX + %printf_lX = OpExtInst %uint %clext printf %fmt_lX %l + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int64.spvasm64 b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int64.spvasm64 new file mode 100644 index 0000000000..8401d1fe9d --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/printf_operands_scalar_int64.spvasm64 @@ -0,0 +1,77 @@ +; kernel void printf_operands_scalar_int64(long l) +; { +; printf("ld = %ld\n", l); +; printf("li = %li\n", l); +; printf("lo = %lo\n", l); +; printf("lu = %lu\n", l); +; printf("lx = %lx\n", l); +; printf("lX = %lX\n", l); +; } + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpCapability Int8 + %clext = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %test "printf_operands_scalar_int64" + %uchar = OpTypeInt 8 0 + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 + %uint_10 = OpConstant %uint 10 + %uint_12 = OpConstant %uint 12 + %uchar_nul = OpConstant %uchar 0 + %uchar_nl = OpConstant %uchar 10 + %uchar_sp = OpConstant %uchar 32 + %uchar_pct = OpConstant %uchar 37 + %uchar_eq = OpConstant %uchar 61 + %uchar_X = OpConstant %uchar 88 + %uchar_d = OpConstant %uchar 100 + %uchar_i = OpConstant %uchar 105 + %uchar_l = OpConstant %uchar 108 + %uchar_o = OpConstant %uchar 111 + %uchar_u = OpConstant %uchar 117 + %uchar_x = OpConstant %uchar 120 + %string_8 = OpTypeArray %uchar %uint_8 +%cptr_string_8 = OpTypePointer UniformConstant %string_8 + %string_10 = OpTypeArray %uchar %uint_10 +%cptr_string_10 = OpTypePointer UniformConstant %string_10 + %string_12 = OpTypeArray %uchar %uint_12 +%cptr_string_12 = OpTypePointer UniformConstant %string_12 + %void = OpTypeVoid + %float = OpTypeFloat 32 + %kernel_sig = OpTypeFunction %void %ulong + %cptr_char = OpTypePointer UniformConstant %uchar + + %array_ld = OpConstantComposite %string_10 %uchar_l %uchar_d %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_d %uchar_nl %uchar_nul ; "ld = %ld\n" + %string_ld = OpVariable %cptr_string_10 UniformConstant %array_ld + %array_li = OpConstantComposite %string_10 %uchar_l %uchar_i %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_i %uchar_nl %uchar_nul ; "li = %li\n" + %string_li = OpVariable %cptr_string_10 UniformConstant %array_li + %array_lo = OpConstantComposite %string_10 %uchar_l %uchar_o %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_o %uchar_nl %uchar_nul ; "lo = %lo\n" + %string_lo = OpVariable %cptr_string_10 UniformConstant %array_lo + %array_lu = OpConstantComposite %string_10 %uchar_l %uchar_u %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_u %uchar_nl %uchar_nul ; "lu = %lu\n" + %string_lu = OpVariable %cptr_string_10 UniformConstant %array_lu + %array_lx = OpConstantComposite %string_10 %uchar_l %uchar_x %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_x %uchar_nl %uchar_nul ; "lx = %lx\n" + %string_lx = OpVariable %cptr_string_10 UniformConstant %array_lx + %array_lX = OpConstantComposite %string_10 %uchar_l %uchar_X %uchar_sp %uchar_eq %uchar_sp %uchar_pct %uchar_l %uchar_X %uchar_nl %uchar_nul ; "lX = %lX\n" + %string_lX = OpVariable %cptr_string_10 UniformConstant %array_lX + + %test = OpFunction %void None %kernel_sig + %l = OpFunctionParameter %ulong + %entry = OpLabel + + %fmt_ld = OpBitcast %cptr_char %string_ld + %printf_ld = OpExtInst %uint %clext printf %fmt_ld %l + %fmt_li = OpBitcast %cptr_char %string_li + %printf_li = OpExtInst %uint %clext printf %fmt_li %l + %fmt_lo = OpBitcast %cptr_char %string_lo + %printf_lo = OpExtInst %uint %clext printf %fmt_lo %l + %fmt_lu = OpBitcast %cptr_char %string_lu + %printf_lu = OpExtInst %uint %clext printf %fmt_lu %l + %fmt_lx = OpBitcast %cptr_char %string_lx + %printf_lx = OpExtInst %uint %clext printf %fmt_lx %l + %fmt_lX = OpBitcast %cptr_char %string_lX + %printf_lX = OpExtInst %uint %clext printf %fmt_lX %l + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/test_extinst_printf.cpp b/test_conformance/spirv_new/test_extinst_printf.cpp new file mode 100644 index 0000000000..dc2bbf97ea --- /dev/null +++ b/test_conformance/spirv_new/test_extinst_printf.cpp @@ -0,0 +1,230 @@ +// +// Copyright (c) 2025 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + +#include "harness/os_helpers.h" +#include "testBase.h" + +#if defined(_WIN32) +#include +#define streamDup(fd1) _dup(fd1) +#define streamDup2(fd1, fd2) _dup2(fd1, fd2) +#else +#if defined(__APPLE__) +#include +#endif +#include +#define streamDup(fd1) dup(fd1) +#define streamDup2(fd1, fd2) dup2(fd1, fd2) +#endif + +#include +#include + +// TODO: Unify with test_printf. +struct StreamGrabber +{ + StreamGrabber(): tempFileName(get_temp_filename()) {} + ~StreamGrabber() + { + if (acquired) + { + release(); + } + } + + int acquire(void) + { + if (acquired == false) + { + old_fd = streamDup(fileno(stdout)); + if (!freopen(tempFileName.c_str(), "w", stdout)) + { + release(); + return -1; + } + acquired = true; + } + return 0; + } + + int release(void) + { + if (acquired == true) + { + fflush(stdout); + streamDup2(old_fd, fileno(stdout)); + close(old_fd); + acquired = false; + } + return 0; + } + + int get_results(std::string& results) + { + if (acquired == false) + { + std::ifstream is(tempFileName, std::ios::binary); + if (is.good()) + { + size_t filesize = 0; + is.seekg(0, std::ios::end); + filesize = (size_t)is.tellg(); + is.seekg(0, std::ios::beg); + + results.clear(); + results.resize(filesize); + is.read(&results[0], filesize); + + return 0; + } + } + return -1; + } + + std::string tempFileName; + int old_fd = 0; + bool acquired = false; +}; + +template +static int printf_operands_helper(cl_context context, cl_device_id device, + cl_command_queue queue, + const char* spirvFileName, + const char* kernelName, + const char* expectedResults, T value) +{ + StreamGrabber grabber; + cl_int error; + + clProgramWrapper program; + error = get_program_with_il(program, device, context, spirvFileName); + test_error(error, "Unable to build SPIR-V program"); + + clKernelWrapper kernel = clCreateKernel(program, kernelName, &error); + test_error(error, "Unable to create SPIR-V kernel"); + + error = clSetKernelArg(kernel, 0, sizeof(value), &value); + test_error(error, "Unable to set kernel arguments"); + + size_t global = 1; + grabber.acquire(); + error |= clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, + NULL, NULL); + error |= clFinish(queue); + grabber.release(); + test_error(error, "unable to enqueue kernel"); + + std::string results; + grabber.get_results(results); + + if (results != std::string(expectedResults)) + { + log_error("Results do not match.\n"); + log_error("Expected: \n---\n%s---\n", expectedResults); + log_error("Got: \n---\n%s---\n", results.c_str()); + return TEST_FAIL; + } + + return TEST_PASS; +} + +REGISTER_TEST(extinst_printf_operands_scalar_int32) +{ + static const char* expected = R"(d = 1 +i = 1 +o = 1 +u = 1 +x = 1 +X = 1 +hd = 1 +hi = 1 +ho = 1 +hu = 1 +hx = 1 +hX = 1 +hhd = 1 +hhi = 1 +hho = 1 +hhu = 1 +hhx = 1 +hhX = 1 +)"; + + return printf_operands_helper(context, device, queue, + "printf_operands_scalar_int32", + "printf_operands_scalar_int32", expected, 1); +} + +REGISTER_TEST(extinst_printf_operands_scalar_fp32) +{ + static const char* expected = R"(a = 0x1.0p+1 +A = 0X1.0P+1 +e = 2.0e+00 +E = 2.0E+00 +f = 2.0 +F = 2.0 +g = 2 +G = 2 +)"; + + return printf_operands_helper(context, device, queue, + "printf_operands_scalar_fp32", + "printf_operands_scalar_fp32", expected, 2.0f); +} + +REGISTER_TEST(extinst_printf_operands_scalar_int64) +{ + static const char* expected = R"(ld = 4 +li = 4 +lo = 4 +lu = 4 +lx = 4 +lX = 4 +)"; + + if (!gHasLong) + { + log_info("Device does not support 64-bit integers. Skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + return printf_operands_helper(context, device, queue, + "printf_operands_scalar_int64", + "printf_operands_scalar_int64", expected, 4L); +} + +REGISTER_TEST(extinst_printf_operands_scalar_fp64) +{ +static const char* expected = R"(a = 0x1.0p+3 +A = 0X1.0P+3 +e = 8.0e+00 +E = 8.0E+00 +f = 8.0 +F = 8.0 +g = 8 +G = 8 +)"; + + if (!is_extension_available(device, "cl_khr_fp64")) + { + log_info("Device does not support fp64. Skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + return printf_operands_helper(context, device, queue, + "printf_operands_scalar_fp64", + "printf_operands_scalar_fp64", expected, 8.0); +} From 0b06b490d0a3e37a5bbea418b2e1c555f3b0664a Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 6 Jan 2025 21:22:29 -0800 Subject: [PATCH 2/4] fix formatting --- test_conformance/spirv_new/test_extinst_printf.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/test_conformance/spirv_new/test_extinst_printf.cpp b/test_conformance/spirv_new/test_extinst_printf.cpp index dc2bbf97ea..e3a3201660 100644 --- a/test_conformance/spirv_new/test_extinst_printf.cpp +++ b/test_conformance/spirv_new/test_extinst_printf.cpp @@ -180,9 +180,9 @@ g = 2 G = 2 )"; - return printf_operands_helper(context, device, queue, - "printf_operands_scalar_fp32", - "printf_operands_scalar_fp32", expected, 2.0f); + return printf_operands_helper( + context, device, queue, "printf_operands_scalar_fp32", + "printf_operands_scalar_fp32", expected, 2.0f); } REGISTER_TEST(extinst_printf_operands_scalar_int64) @@ -208,7 +208,7 @@ lX = 4 REGISTER_TEST(extinst_printf_operands_scalar_fp64) { -static const char* expected = R"(a = 0x1.0p+3 + static const char* expected = R"(a = 0x1.0p+3 A = 0X1.0P+3 e = 8.0e+00 E = 8.0E+00 From 6947e333ffbf5b037c41e424e9c7450e40e32c98 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Fri, 10 Jan 2025 07:58:34 -0800 Subject: [PATCH 3/4] fix memory leak --- test_common/harness/os_helpers.cpp | 2 +- test_conformance/spirv_new/test_extinst_printf.cpp | 7 ++++++- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/test_common/harness/os_helpers.cpp b/test_common/harness/os_helpers.cpp index c64c59011e..b708751167 100644 --- a/test_common/harness/os_helpers.cpp +++ b/test_common/harness/os_helpers.cpp @@ -577,7 +577,7 @@ char* get_temp_filename() close(fd); #elif defined(_WIN32) UINT ret = GetTempFileName(".", "tmp", 0, gFileName); - if (ret == 0) return gFileName; + if (ret == 0) return strdup(gFileName); #else MTdata d = init_genrand((cl_uint)time(NULL)); sprintf(gFileName, "tmpfile.%u", genrand_int32(d)); diff --git a/test_conformance/spirv_new/test_extinst_printf.cpp b/test_conformance/spirv_new/test_extinst_printf.cpp index e3a3201660..53c794aa5f 100644 --- a/test_conformance/spirv_new/test_extinst_printf.cpp +++ b/test_conformance/spirv_new/test_extinst_printf.cpp @@ -36,7 +36,12 @@ // TODO: Unify with test_printf. struct StreamGrabber { - StreamGrabber(): tempFileName(get_temp_filename()) {} + StreamGrabber() + { + char* tmp = get_temp_filename(); + tempFileName = tmp; + free(tmp); + } ~StreamGrabber() { if (acquired) From 4ec8cdb74c7d711cdba766a87423742d79e343d4 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Tue, 14 Jan 2025 17:57:27 -0800 Subject: [PATCH 4/4] add support for a printf callback for cl_arm_printf --- .../spirv_new/test_extinst_printf.cpp | 46 ++++++++++++++----- 1 file changed, 35 insertions(+), 11 deletions(-) diff --git a/test_conformance/spirv_new/test_extinst_printf.cpp b/test_conformance/spirv_new/test_extinst_printf.cpp index 53c794aa5f..54bb832671 100644 --- a/test_conformance/spirv_new/test_extinst_printf.cpp +++ b/test_conformance/spirv_new/test_extinst_printf.cpp @@ -104,9 +104,15 @@ struct StreamGrabber bool acquired = false; }; +// printf callback, for cl_arm_printf +void CL_CALLBACK printfCallBack(const char* printf_data, size_t len, + size_t final, void* user_data) +{ + fwrite(printf_data, 1, len, stdout); +} + template -static int printf_operands_helper(cl_context context, cl_device_id device, - cl_command_queue queue, +static int printf_operands_helper(cl_device_id device, const char* spirvFileName, const char* kernelName, const char* expectedResults, T value) @@ -114,6 +120,27 @@ static int printf_operands_helper(cl_context context, cl_device_id device, StreamGrabber grabber; cl_int error; + // Create a context and a queue to test with. + // We cannot use the context and queue from the harness because some + // implementations require a printf callback to be set at context creation. + + cl_context_properties printf_properties[] = { + CL_PRINTF_CALLBACK_ARM, (cl_context_properties)printfCallBack, + CL_PRINTF_BUFFERSIZE_ARM, 256, 0 + }; + + cl_context_properties* props = + is_extension_available(device, "cl_arm_printf") ? printf_properties + : nullptr; + + clContextWrapper context = + clCreateContext(props, 1, &device, notify_callback, nullptr, &error); + test_error(error, "Unable to create printf context"); + + clCommandQueueWrapper queue = + clCreateCommandQueue(context, device, 0, &error); + test_error(error, "Unable to create printf queue"); + clProgramWrapper program; error = get_program_with_il(program, device, context, spirvFileName); test_error(error, "Unable to build SPIR-V program"); @@ -168,8 +195,7 @@ hhx = 1 hhX = 1 )"; - return printf_operands_helper(context, device, queue, - "printf_operands_scalar_int32", + return printf_operands_helper(device, "printf_operands_scalar_int32", "printf_operands_scalar_int32", expected, 1); } @@ -185,9 +211,9 @@ g = 2 G = 2 )"; - return printf_operands_helper( - context, device, queue, "printf_operands_scalar_fp32", - "printf_operands_scalar_fp32", expected, 2.0f); + return printf_operands_helper(device, "printf_operands_scalar_fp32", + "printf_operands_scalar_fp32", expected, + 2.0f); } REGISTER_TEST(extinst_printf_operands_scalar_int64) @@ -206,8 +232,7 @@ lX = 4 return TEST_SKIPPED_ITSELF; } - return printf_operands_helper(context, device, queue, - "printf_operands_scalar_int64", + return printf_operands_helper(device, "printf_operands_scalar_int64", "printf_operands_scalar_int64", expected, 4L); } @@ -229,7 +254,6 @@ G = 8 return TEST_SKIPPED_ITSELF; } - return printf_operands_helper(context, device, queue, - "printf_operands_scalar_fp64", + return printf_operands_helper(device, "printf_operands_scalar_fp64", "printf_operands_scalar_fp64", expected, 8.0); }