From a61feea65677c2bbaf479a413d5f28a1e3464920 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Tue, 11 Feb 2025 08:55:39 -0800 Subject: [PATCH] adds SPIR-V tests for scalar printf operands (#2211) Adds targeted SPIR-V tests for printf with scalar operands. See: * https://github.com/KhronosGroup/OpenCL-Docs/issues/1211 * https://github.com/KhronosGroup/OpenCL-Docs/pull/1236 The fp32 test is likely to be the most interesting, especially on devices that support fp64, because printf with scalar fp32 operands is not generated by default in this case with Clang and the SPIR-V LLVM Translator. --- test_common/harness/os_helpers.cpp | 2 +- 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 | 259 ++++++++++++++++++ 11 files changed, 1051 insertions(+), 1 deletion(-) 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_common/harness/os_helpers.cpp b/test_common/harness/os_helpers.cpp index c64c5901..b7087511 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/CMakeLists.txt b/test_conformance/spirv_new/CMakeLists.txt index 805e851b..c635e924 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 00000000..58631498 --- /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 00000000..ba415d09 --- /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 00000000..1b31cf49 --- /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 00000000..a947e5ec --- /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 00000000..61fb8cd0 --- /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 00000000..91ad8e1e --- /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 00000000..ec19e9f8 --- /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 00000000..8401d1fe --- /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 00000000..54bb8326 --- /dev/null +++ b/test_conformance/spirv_new/test_extinst_printf.cpp @@ -0,0 +1,259 @@ +// +// 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() + { + char* tmp = get_temp_filename(); + tempFileName = tmp; + free(tmp); + } + ~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; +}; + +// 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_device_id device, + const char* spirvFileName, + const char* kernelName, + const char* expectedResults, T value) +{ + 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"); + + 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(device, "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(device, "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(device, "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(device, "printf_operands_scalar_fp64", + "printf_operands_scalar_fp64", expected, 8.0); +}