Search Unity

Question 'disable_fastmath' : unknown pragma ignored in compute shader

Discussion in 'Editor & General Support' started by jg482, Aug 7, 2022.

  1. jg482

    jg482

    Joined:
    Apr 5, 2018
    Posts:
    8
    Hey there,

    I'm trying to dsiable fastmath on metal in my compute shader. I added the pragma by adding:

    Code (CSharp):
    1. #pragma disable_fastmath
    in my compute shader. But when compiling I get a shader warning:

    Shader warning in 'Test': 'disable_fastmath' : unknown pragma ignored at kernel CSMain at Assets/Mesh/Test.compute(1) (on metal)

    Any idea why it is not working? I tried to add it in different lines e.g. before the kernel pragma and after but it didn't change anything.
     
  2. jg482

    jg482

    Joined:
    Apr 5, 2018
    Posts:
    8
    I worked into the HLSLcc code on github. The HLSLCC_FLAG_DISABLE_FASTMATH flag in psContext->flags gets checked in HLSLcc/src/toMetal.cpp in line 167. And there it should add #define UNITY_DISABLE_FASTMATH for alle shader types (including compute shaders) except hull and domain shaders.

    So my concern that the flag is not used in compute shaders is unfounded. Obviously it is not set correctly in the previous step in the psContext->flags or is not even recognized correctly by the precompiler. Maybe it is a bug? Any thoughts on this?
     
  3. jg482

    jg482

    Joined:
    Apr 5, 2018
    Posts:
    8
    Okay, I found out that despite the warning it gets set in the compiled shader

    Code (CSharp):
    1. **** Platform Metal:
    2. Compiled code for kernel CSMain
    3. keywords: <none>
    4. #define UNITY_DISABLE_FASTMATH
    5. #include <metal_stdlib>
    6. #include <metal_texture>
    7. using namespace metal;
    8. struct _Buffer_Type
    9. {
    10.     uint value[2];
    11. };
    12.  
    13. kernel void computeMain(
    14.     device _Buffer_Type *_Buffer [[ buffer(0) ]],
    15.     uint3 mtl_ThreadID [[ thread_position_in_grid ]])
    16. {
    17.     float2 u_xlat0;
    18.     u_xlat0.xy = float2(as_type<float>(_Buffer[mtl_ThreadID.x].value[(0x0 >> 2) + 0]), as_type<float>(_Buffer[mtl_ThreadID.x].value[(0x0 >> 2) + 1]));
    19.     u_xlat0.x = u_xlat0.y + u_xlat0.x;
    20.     u_xlat0.y = 0.0;
    21.     _Buffer[mtl_ThreadID.x].value[(0x0 >> 2)] = as_type<uint>(u_xlat0.x);
    22.     _Buffer[mtl_ThreadID.x].value[(0x0 >> 2) + 1] = as_type<uint>(u_xlat0.y);
    23.     return;
    24. }
    So the problem probably lies somewhere else