Alastair Donaldson f7e73d4ee3 Add tests derived from VK-GL-CTS
This adds SPIR-V assembly and WGSL tests derived from VK-GL-CTS commit
571256871c2e2f03995373e1e4a02958d8cd8cf5. The following procedure was
followed:

- Those .amber files in VK-GL-CTS wholly owned by Google were
  identified

- All GLSL and SPIR-V shaders were extracted from the Amber files and
  converted into SPIR-V binaries

- The compact-ids pass of spirv-opt was applied to each binary

- Duplicate binaries were removed

- spirv-opt -O was used to obtain an optimized version of each remaining
  binary, with duplicates discarded

- Binaries that failed validation using spirv-val with target
  environment SPIR-V 1.3 were discarded

- Those binaries that tint could not successfully convert into WGSL were
  put aside for further investigation

- SPIR-V assembly versions of the remaining binaries are included in
  this CL

- test-runner with -generate-expected and -generate-skip was used to
  generate expected .spvasm, .msl, .hlsl and .wgsl outputs for these
  SPIR-V assembly tests

- Each successfully-generated .expected.wgsl is included in this CL
  again, as a WGLSL test

- test-runner with -generate-expected and -generate-skip was used again,
  to generate expected outputs for these WGSL tests

Change-Id: Ibe9baf2729cf97e0b633db9a426f53362a5de540
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58842
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
2021-07-23 13:10:12 +00:00

58 lines
1.3 KiB
Plaintext

#include <metal_stdlib>
using namespace metal;
struct doesNotMatter {
/* 0x0000 */ int global_seed;
/* 0x0004 */ int data[1];
};
struct buf1 {
/* 0x0000 */ packed_float2 injectionSwitch;
};
void main_1(constant buf1& x_10, device doesNotMatter& x_7, thread uint3* const tint_symbol_2) {
int lid = 0;
int val = 0;
int i = 0;
uint const x_40 = (*(tint_symbol_2)).x;
lid = as_type<int>(x_40);
int const x_43 = x_7.global_seed;
val = x_43;
i = 0;
while (true) {
int const x_48 = i;
if ((x_48 < 2)) {
} else {
break;
}
int const x_51 = lid;
if ((x_51 > 0)) {
int const x_55 = lid;
int const x_58 = x_7.data[(x_55 - 1)];
int const x_59 = val;
val = (x_59 + x_58);
float const x_62 = x_10.injectionSwitch.x;
if ((x_62 > 100.0f)) {
break;
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
{
int const x_66 = i;
i = (x_66 + 1);
}
}
int const x_68 = lid;
if ((x_68 == 0)) {
x_7.data[0] = 42;
}
return;
}
kernel void tint_symbol(uint3 gl_LocalInvocationID_param [[thread_position_in_threadgroup]], constant buf1& x_10 [[buffer(1)]], device doesNotMatter& x_7 [[buffer(0)]]) {
thread uint3 tint_symbol_3 = 0u;
tint_symbol_3 = gl_LocalInvocationID_param;
main_1(x_10, x_7, &(tint_symbol_3));
return;
}