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
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
|
// unit-test-glsl-compile.cpp
#include "../../source/core/slang-io.h"
#include "../../source/core/slang-process.h"
#include "../../tools/platform/performance-counter.h"
#include "slang-com-ptr.h"
#include "slang.h"
#include "unit-test/slang-unit-test.h"
#include <stdio.h>
#include <stdlib.h>
using namespace Slang;
// Test the compilation API for cross-compiling glsl source to SPIRV.
SLANG_UNIT_TEST(compileBenchmark)
{
const char* userSourceBody = R"(
// shader.slang
struct PushConstantCompute
{
uint64_t bufferAddress;
uint numVertices;
};
struct Vertex
{
float3 position;
};
[[vk::push_constant]]
ConstantBuffer<PushConstantCompute> pushConst;
[shader("compute")]
[numthreads(256, 1, 1)]
void main(uint3 threadIdx : SV_DispatchThreadID)
{
uint index = threadIdx.x;
if(index >= pushConst.numVertices)
return;
Vertex* vertices = (Vertex*)pushConst.bufferAddress;
float angle = (index + 1) * 2.3f;
float3 vertex = vertices[index].position;
float cosAngle = cos(angle);
float sinAngle = sin(angle);
float3x3 rotationMatrix = float3x3(
cosAngle, -sinAngle, 0.0,
sinAngle, cosAngle, 0.0,
0.0, 0.0, 1.0
);
float3 rotatedVertex = mul(rotationMatrix, vertex);
vertices[index].position = rotatedVertex;
}
)";
ComPtr<slang::IGlobalSession> globalSession;
SlangGlobalSessionDesc globalDesc = {};
globalDesc.enableGLSL = false;
SLANG_CHECK(slang_createGlobalSession2(&globalDesc, globalSession.writeRef()) == SLANG_OK);
slang::TargetDesc targetDesc = {};
targetDesc.format = SLANG_SPIRV;
targetDesc.profile = globalSession->findProfile("spirv_1_5");
slang::SessionDesc sessionDesc = {};
sessionDesc.targetCount = 1;
sessionDesc.targets = &targetDesc;
auto start = platform::PerformanceCounter::now();
for (int pass = 0; pass < 100; pass++)
{
ComPtr<slang::ISession> session;
SLANG_CHECK(globalSession->createSession(sessionDesc, session.writeRef()) == SLANG_OK);
ComPtr<slang::IBlob> diagnosticBlob;
auto module = session->loadModuleFromSourceString(
"m",
"m.slang",
userSourceBody,
diagnosticBlob.writeRef());
SLANG_CHECK(module != nullptr);
ComPtr<slang::IEntryPoint> entryPoint;
SlangResult res = module->findAndCheckEntryPoint(
"main",
SLANG_STAGE_COMPUTE,
entryPoint.writeRef(),
diagnosticBlob.writeRef());
SLANG_CHECK(res == SLANG_OK);
slang::IComponentType* componentTypes[2] = {module, entryPoint.get()};
ComPtr<slang::IComponentType> composedProgram;
session->createCompositeComponentType(
componentTypes,
2,
composedProgram.writeRef(),
diagnosticBlob.writeRef());
ComPtr<slang::IComponentType> linkedProgram;
composedProgram->link(linkedProgram.writeRef(), diagnosticBlob.writeRef());
ComPtr<slang::IBlob> code;
linkedProgram->getEntryPointCode(0, 0, code.writeRef(), diagnosticBlob.writeRef());
}
auto time = platform::PerformanceCounter::getElapsedTimeInSeconds(start);
getTestReporter()->addExecutionTime(time);
}
|
