search

KhronosGroup / SPIR

Other
178 stars 49 forks source link

How to force to generate "TypeFloat 32" instead of "TypeFloat 64" when calling "printf("%f") #57

Open ncybmh99 opened 7 years ago

ncybmh99 commented 7 years ago

Given a simple 'printf' test kernel as:

__kernel void test(const float x) { printf("%f\n", x); }

My question is:

Suppose the platform doesn't support 'double'. Is there any way for the tool to generate "TypeFloat 32" instead of "TypeFloat 64" and not to introduce "FConvert" when calling "printf"? Is there any compilation options to control it?

The tool generates the following SPV code:

119734787 65536 393230 25 0 2 Capability Addresses 2 Capability Kernel 2 Capability Float64 5 ExtInstImport 1 "OpenCL.std" 3 MemoryModel 1 2 5 EntryPoint 6 16 "test" 3 Source 3 102000 4 Name 12 ".str" 3 Name 17 "x" 4 Name 18 "entry" 4 Name 20 "conv" 4 Name 24 "call" 3 Decorate 12 Constant 4 Decorate 12 Alignment 1 4 TypeInt 2 8 0 4 TypeInt 7 32 0 4 Constant 2 3 37 4 Constant 2 4 102 4 Constant 2 5 10 4 Constant 2 6 0 4 Constant 7 8 4 4 Constant 7 21 0 4 TypeArray 9 2 8 4 TypePointer 11 0 9 2 TypeVoid 13 3 TypeFloat 14 32 4 TypeFunction 15 13 14 3 TypeFloat 19 64 4 TypePointer 22 0 2 7 ConstantComposite 9 10 3 4 5 6 5 Variable 11 12 0 10

5 Function 13 16 0 15 3 FunctionParameter 14 17

2 Label 18 4 FConvert 19 20 17 6 InBoundsPtrAccessChain 22 23 12 21 21 7 ExtInst 7 24 1 printf 23 20 1 Return

1 FunctionEnd

yxsamliu commented 7 years ago

This seems to be a clang bug. spec v1.2 s6.12.13:

The conversion specifiers f, F, e, E, g, G, a, A convert a float argument to a double only if the double data type is supported. Refer to the description of CL_DEVICE_DOUBLE_FP_CONFIG in table 4.3. If the double data type is not supported, the argument will be a float instead of a double.

yxsamliu commented 7 years ago

Or maybe it is not a bug but due to fp64 is on by default. You may try disable fp64 by using

pragma OPENCL EXTENSIONS cl_khr_fp64 : disable

and see if the default promotion of float to double still happens.