-
Notifications
You must be signed in to change notification settings - Fork 193
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
* [hlsl-out] Write `mad` intrinsic for `fma` function - This should be enough because we only support f32 for now. - Adds a new test for WGSL functions, in the spirit of operators.wgsl. - Closes #1579 * Add FMA feature to glsl backend - I think this is right. Just iterate all known expressions in all functions and entry points to locate any `fma` function call. Should not need to walk the statement DAG. * Transform GLSL fma function into an airthmetic expression when necessary * Add tests for GLSL fma function tranformation * Remove the hazard comment from the webgl test input * Add helper method for fma function support checks * Address review comment
- Loading branch information
Showing
15 changed files
with
222 additions
and
4 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,7 @@ | ||
( | ||
glsl: ( | ||
version: Embedded(300), | ||
writer_flags: (bits: 0), | ||
binding_map: {}, | ||
), | ||
) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,13 @@ | ||
fn test_fma() -> vec2<f32> { | ||
let a = vec2<f32>(2.0, 2.0); | ||
let b = vec2<f32>(0.5, 0.5); | ||
let c = vec2<f32>(0.5, 0.5); | ||
|
||
return fma(a, b, c); | ||
} | ||
|
||
|
||
[[stage(vertex)]] | ||
fn main() { | ||
let a = test_fma(); | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,2 @@ | ||
( | ||
) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,15 @@ | ||
fn test_fma() -> vec2<f32> { | ||
let a = vec2<f32>(2.0, 2.0); | ||
let b = vec2<f32>(0.5, 0.5); | ||
let c = vec2<f32>(0.5, 0.5); | ||
|
||
// Hazard: HLSL needs a different intrinsic function for f32 and f64 | ||
// See: https://github.com/gfx-rs/naga/issues/1579 | ||
return fma(a, b, c); | ||
} | ||
|
||
|
||
[[stage(compute), workgroup_size(1)]] | ||
fn main() { | ||
let a = test_fma(); | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,18 @@ | ||
#version 300 es | ||
|
||
precision highp float; | ||
precision highp int; | ||
|
||
|
||
vec2 test_fma() { | ||
vec2 a = vec2(2.0, 2.0); | ||
vec2 b = vec2(0.5, 0.5); | ||
vec2 c = vec2(0.5, 0.5); | ||
return (a * b + c); | ||
} | ||
|
||
void main() { | ||
vec2 _e0 = test_fma(); | ||
return; | ||
} | ||
|
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,21 @@ | ||
#version 310 es | ||
#extension GL_EXT_gpu_shader5 : require | ||
|
||
precision highp float; | ||
precision highp int; | ||
|
||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; | ||
|
||
|
||
vec2 test_fma() { | ||
vec2 a = vec2(2.0, 2.0); | ||
vec2 b = vec2(0.5, 0.5); | ||
vec2 c = vec2(0.5, 0.5); | ||
return fma(a, b, c); | ||
} | ||
|
||
void main() { | ||
vec2 _e0 = test_fma(); | ||
return; | ||
} | ||
|
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,15 @@ | ||
|
||
float2 test_fma() | ||
{ | ||
float2 a = float2(2.0, 2.0); | ||
float2 b = float2(0.5, 0.5); | ||
float2 c = float2(0.5, 0.5); | ||
return mad(a, b, c); | ||
} | ||
|
||
[numthreads(1, 1, 1)] | ||
void main() | ||
{ | ||
const float2 _e0 = test_fma(); | ||
return; | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,3 @@ | ||
vertex=() | ||
fragment=() | ||
compute=(main:cs_5_1 ) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,18 @@ | ||
// language: metal1.1 | ||
#include <metal_stdlib> | ||
#include <simd/simd.h> | ||
|
||
|
||
metal::float2 test_fma( | ||
) { | ||
metal::float2 a = metal::float2(2.0, 2.0); | ||
metal::float2 b = metal::float2(0.5, 0.5); | ||
metal::float2 c = metal::float2(0.5, 0.5); | ||
return metal::fma(a, b, c); | ||
} | ||
|
||
kernel void main_( | ||
) { | ||
metal::float2 _e0 = test_fma(); | ||
return; | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,33 @@ | ||
; SPIR-V | ||
; Version: 1.1 | ||
; Generator: rspirv | ||
; Bound: 20 | ||
OpCapability Shader | ||
%1 = OpExtInstImport "GLSL.std.450" | ||
OpMemoryModel Logical GLSL450 | ||
OpEntryPoint GLCompute %16 "main" | ||
OpExecutionMode %16 LocalSize 1 1 1 | ||
%2 = OpTypeVoid | ||
%4 = OpTypeFloat 32 | ||
%3 = OpConstant %4 2.0 | ||
%5 = OpConstant %4 0.5 | ||
%6 = OpTypeVector %4 2 | ||
%9 = OpTypeFunction %6 | ||
%17 = OpTypeFunction %2 | ||
%8 = OpFunction %6 None %9 | ||
%7 = OpLabel | ||
OpBranch %10 | ||
%10 = OpLabel | ||
%11 = OpCompositeConstruct %6 %3 %3 | ||
%12 = OpCompositeConstruct %6 %5 %5 | ||
%13 = OpCompositeConstruct %6 %5 %5 | ||
%14 = OpExtInst %6 %1 Fma %11 %12 %13 | ||
OpReturnValue %14 | ||
OpFunctionEnd | ||
%16 = OpFunction %2 None %17 | ||
%15 = OpLabel | ||
OpBranch %18 | ||
%18 = OpLabel | ||
%19 = OpFunctionCall %6 %8 | ||
OpReturn | ||
OpFunctionEnd |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,12 @@ | ||
fn test_fma() -> vec2<f32> { | ||
let a = vec2<f32>(2.0, 2.0); | ||
let b = vec2<f32>(0.5, 0.5); | ||
let c = vec2<f32>(0.5, 0.5); | ||
return fma(a, b, c); | ||
} | ||
|
||
[[stage(compute), workgroup_size(1, 1, 1)]] | ||
fn main() { | ||
let _e0 = test_fma(); | ||
return; | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters