mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
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.
This commit is contained in:
Ben Ashbaugh
@@ -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
|
||||
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
259
test_conformance/spirv_new/test_extinst_printf.cpp
Normal file
259
test_conformance/spirv_new/test_extinst_printf.cpp
Normal file
@@ -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 <io.h>
|
||||
#define streamDup(fd1) _dup(fd1)
|
||||
#define streamDup2(fd1, fd2) _dup2(fd1, fd2)
|
||||
#else
|
||||
#if defined(__APPLE__)
|
||||
#include <sys/sysctl.h>
|
||||
#endif
|
||||
#include <unistd.h>
|
||||
#define streamDup(fd1) dup(fd1)
|
||||
#define streamDup2(fd1, fd2) dup2(fd1, fd2)
|
||||
#endif
|
||||
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
|
||||
// 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 <typename T>
|
||||
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);
|
||||
}
|
||||
Reference in New Issue
Block a user