1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
|
//TEST:SIMPLE(filecheck=CHECK_GLSL): -stage compute -entry computeMain -allow-glsl -target glsl
//TEST:SIMPLE(filecheck=CHECK_SPV): -stage compute -entry computeMain -allow-glsl -target spirv -emit-spirv-directly
//TEST:SIMPLE(filecheck=CHECK_HLSL): -stage compute -entry computeMain -target hlsl -allow-glsl
//TEST:SIMPLE(filecheck=CHECK_CUDA): -stage compute -entry computeMain -target cuda -allow-glsl
//TEST(compute, vulkan):COMPARE_COMPUTE(filecheck-buffer=BUF):-vk -compute -entry computeMain -allow-glsl
//TEST(compute, vulkan):COMPARE_COMPUTE(filecheck-buffer=BUF):-vk -compute -entry computeMain -allow-glsl -emit-spirv-directly
#version 450
//TEST_INPUT:ubuffer(data=[0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15], stride=4):name=inputBuffer
layout(binding = 0) buffer MyBlockNameIn
{
uvec4 a[4];
} inputBuffer;
//TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name=outputBuffer
layout(binding = 1) buffer MyBlockName
{
uvec4 a;
uvec4 b;
uvec3 c;
} outputBuffer;
layout(local_size_x = 1) in;
void computeMain()
{
{
f16vec4 m0 = f16vec4(inputBuffer.a[0]);
f16vec4 m1 = f16vec4(inputBuffer.a[1]);
f16vec4 m2 = f16vec4(inputBuffer.a[2]);
f16vec4 m3 = f16vec4(inputBuffer.a[3]);
f16mat4 m = f16mat4(m0, m1, m2, m3);
f16vec4 res = m * f16vec4(1, 1, 1, 1);
outputBuffer.a = uvec4(res);
// BUF: 18
// BUF-NEXT: 1C
// BUF-NEXT: 20
// BUF-NEXT: 24
}
{
f16vec4 m0 = f16vec4(inputBuffer.a[0]);
f16vec4 m1 = f16vec4(inputBuffer.a[1]);
f16vec4 m2 = f16vec4(inputBuffer.a[2]);
f16mat3x4 m = f16mat3x4(m0, m1, m2);
f16vec4 res = m * f16vec3(1, 1, 1);
outputBuffer.b = uvec4(res);
// BUF-NEXT: C
// BUF-NEXT: F
// BUF-NEXT: 12
// BUF-NEXT: 15
}
{
f16vec3 m0 = f16vec3(inputBuffer.a[0].xyz);
f16vec3 m1 = f16vec3(inputBuffer.a[1].xyz);
f16mat2x3 m = f16mat2x3(m0, m1);
f16vec3 res = m * f16vec2(1, 1);
outputBuffer.c = uvec3(res);
// BUF-NEXT: 4
// BUF-NEXT: 6
// BUF-NEXT: 8
}
// CHECK_GLSL: f16mat4x4
// CHECK_GLSL: f16vec4
// CHECK_GLSL: f16mat3x4
// CHECK_GLSL: f16mat2x3
// CHECK_GLSL: f16vec3
// CHECK_GLSL: f16vec2
// CHECK_SPV: OpTypeFloat 16
// CHECK_HLSL: vector<half,4>
// CHECK_HLSL: matrix<half,int(4),int(4)>
// CHECK_HLSL: matrix<half,int(3),int(4)>
// CHECK_HLSL: matrix<half,int(2),int(3)>
// CHECK_CUDA: __half4
// CHECK_CUDA: makeMatrix<__half, 4, 4>
// CHECK_CUDA: makeMatrix<__half, 3, 4>
// CHECK_CUDA: makeMatrix<__half, 2, 3>
}
|
