forked from gfx-rs/wgpu
-
Notifications
You must be signed in to change notification settings - Fork 0
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Make tests properly exercise Some("") global names. (gfx-rs#1503)
The GLSL empty-global-name.frag test doesn't suffice because the GLSL front end doesn't produce the same IR as the SPIR-V included in the bug report. As far as I know, only a genuine SPIR-V input test can produce a global whose name is `Some("")`. Include the SPIR-V assembly source.
- Loading branch information
Showing
9 changed files
with
115 additions
and
28 deletions.
There are no files selected for viewing
This file was deleted.
Oops, something went wrong.
Binary file not shown.
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,44 @@ | ||
;; Make sure we handle globals whose assigned name is "". | ||
;; | ||
;; In MSL, the anonymous global sometimes ends up looking like | ||
;; | ||
;; struct Blah { int member; } ; | ||
;; | ||
;; where the null name just becomes an empty string before that last semicolon. | ||
;; This is, unfortunately, valid MSL, simply declaring the type Blah, so it will | ||
;; pass validation. However, an attempt to *use* the global will generate a | ||
;; garbage expression like ".member", so we include a function that returns the | ||
;; member's value. | ||
|
||
OpCapability Shader | ||
OpMemoryModel Logical GLSL450 | ||
OpEntryPoint GLCompute %main "main" %global | ||
OpExecutionMode %main LocalSize 64 1 1 | ||
|
||
OpName %global "" | ||
OpDecorate %block Block | ||
OpMemberDecorate %block 0 Offset 0 | ||
OpDecorate %global DescriptorSet 0 | ||
OpDecorate %global Binding 0 | ||
|
||
%void = OpTypeVoid | ||
%int = OpTypeInt 32 1 | ||
%block = OpTypeStruct %int | ||
%ptr_int = OpTypePointer StorageBuffer %int | ||
%ptr_block = OpTypePointer StorageBuffer %block | ||
%fn_void = OpTypeFunction %void | ||
%fn_int = OpTypeFunction %int | ||
%zero = OpConstant %int 0 | ||
%one = OpConstant %int 1 | ||
|
||
;; This global is said to have a name of "". | ||
%global = OpVariable %ptr_block StorageBuffer | ||
|
||
%main = OpFunction %void None %fn_void | ||
%main_prelude = OpLabel | ||
%member_ptr = OpAccessChain %ptr_int %global %zero | ||
%member_val = OpLoad %int %member_ptr | ||
%plus_one = OpIAdd %int %member_val %one | ||
OpStore %member_ptr %plus_one | ||
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,19 @@ | ||
|
||
struct type1 { | ||
int member; | ||
}; | ||
|
||
RWByteAddressBuffer global : register(u0); | ||
|
||
void function() | ||
{ | ||
int _expr8 = asint(global.Load(0)); | ||
global.Store(0, asuint((_expr8 + 1))); | ||
return; | ||
} | ||
|
||
[numthreads(64, 1, 1)] | ||
void main() | ||
{ | ||
function(); | ||
} |
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,21 @@ | ||
// language: metal1.1 | ||
#include <metal_stdlib> | ||
#include <simd/simd.h> | ||
|
||
struct type1 { | ||
int member; | ||
}; | ||
|
||
void function( | ||
device type1& global | ||
) { | ||
int _e8 = global.member; | ||
global.member = _e8 + 1; | ||
return; | ||
} | ||
|
||
kernel void main1( | ||
device type1& global [[user(fake0)]] | ||
) { | ||
function(global); | ||
} |
This file was deleted.
Oops, something went wrong.
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 @@ | ||
[[block]] | ||
struct type2 { | ||
member: i32; | ||
}; | ||
|
||
[[group(0), binding(0)]] | ||
var<storage, read_write> global: type2; | ||
|
||
fn function1() { | ||
let e8: i32 = global.member; | ||
global.member = (e8 + 1); | ||
return; | ||
} | ||
|
||
[[stage(compute), workgroup_size(64, 1, 1)]] | ||
fn main() { | ||
function1(); | ||
} |
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