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
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
|
Slang CUDA Target Support
=========================
Slang has preliminary support for producing CUDA source, and PTX binaries using nvrtc.
# Features
* Can compile Slang source into CUDA source code
* Supports compute style shaders
* Supports a 'bindless' CPU like model
* Can compile CUDA source to PTX through 'pass through' mechansism
# Limitations
These limitations apply to Slang transpiling to CUDA.
* Only supports the 'texture object' style binding (The texture object API is only supported on devices of compute capability 3.0 or higher. )
* Samplers are not separate objects in CUDA - they are combined into a single 'TextureObject'. So samplers are effectively ignored on CUDA targets.
* Whilst there is tex1Dfetch there are no equivalents for higher dimensions - so such accesses are not currently supported
* When using a TextureArray (layered texture in CUDA) - the index will be treated as an int, as this is all CUDA allows
* Care must be used in using `WaveGetLaneIndex` wave intrinsic - it will only give the right results for appopriate launches
The following are a work in progress or not implmented but are planned to be so in the future
* Resource types including surfaces
# How it works
For producing PTX binaries Slang uses nvrtc. Nvrtc dll/shared library has to be available to Slang (in the appropriate PATH for example) for it to be able to produce PTX. The nvrtc compiler can be accessed directly through
```
SLANG_PASS_THROUGH_NVRTC,
```
Much like other targets that use downstream compilers Slang can be used to compile CUDA source directly to PTX via the pass through mechansism. That the Slang command line options will broadly be mapped down to the appropriate options for the nvrtc compilation. In the API the `SlangCompileTarget` for CUDA is `SLANG_CUDA_SOURCE` and for PTX is `SLANG_PTX`. These can also be specified on the Slang command line as `-target cuda` and `-target ptx`.
Binding
=======
Say we have some Slang source like the following:
```
struct Thing { int a; int b; }
Texture2D<float> tex;
SamplerState sampler;
RWStructuredBuffer<int> outputBuffer;
ConstantBuffer<Thing> thing3;
[numthreads(4, 1, 1)]
void computeMain(
uint3 dispatchThreadID : SV_DispatchThreadID,
uniform Thing thing,
uniform Thing thing2)
{
// ...
}
```
This will be turned into a CUDA entry point with
```
struct UniformEntryPointParams
{
Thing thing;
Thing thing2;
};
struct UniformState
{
CUtexObject tex; // This is the combination of a texture and a sampler(!)
SamplerState sampler; // This variable exists within the layout, but it's value is not used.
RWStructuredBuffer<int32_t> outputBuffer; // This is implemented as a template in the CUDA prelude. It's just a pointer, and a size
Thing* thing3; // Constant buffers map to pointers
};
// [numthreads(4, 1, 1)]
extern "C" __global__ void computeMain(UniformEntryPointParams* params, UniformState* uniformState)
```
With CUDA - the caller specifies how threading is broken up, so `[numthreads]` is available through reflection, and in a comment in output source code but does not produce varying code.
The UniformState and UniformEntryPointParams struct typically vary by shader. UniformState holds 'normal' bindings, whereas UniformEntryPointParams hold the uniform entry point parameters. Where specific bindings or parameters are located can be determined by reflection. The structures for the example above would be something like the following...
`StructuredBuffer<T>`,`RWStructuredBuffer<T>` become
```
T* data;
size_t count;
```
`ByteAddressBuffer`, `RWByteAddressBuffer` become
```
uint32_t* data;
size_t sizeInBytes;
```
## Unsized arrays
Unsized arrays can be used, which are indicated by an array with no size as in `[]`. For example
```
RWStructuredBuffer<int> arrayOfArrays[];
```
With normal 'sized' arrays, the elements are just stored contiguously within wherever they are defined. With an unsized array they map to `Array<T>` which is...
```
T* data;
size_t count;
```
Note that there is no method in the shader source to get the `count`, even though on the CUDA target it is stored and easily available. This is because of the behavior on GPU targets
* That the count has to be stored elsewhere (unlike with CUDA)
* On some GPU targets there is no bounds checking - accessing outside the bound values can cause *undefined behavior*
* The elements may be laid out *contiguously* on GPU
In practice this means if you want to access the `count` in shader code it will need to be passed by another mechanism - such as within a constant buffer. It is possible in the future support may be added to allow direct access of `count` work across targets transparently.
## Prelude
For CUDA the code to support the code generated by Slang is partly defined within the 'prelude'. The prelude is inserted text placed before the generated CUDA source code. For the Slang command line tools as well as the test infrastructure, the prelude functionality is achieved through a `#include` in the prelude text of the `prelude/slang-cuda-prelude.h` specified with an absolute path. Doing so means other files the `slang-cuda-prelude.h` might need can be specified relatively, and include paths for the backend compiler do not need to be modified.
The prelude needs to define
* 'Built in' types (vector, matrix, 'object'-like Texture, SamplerState etc)
* Scalar intrinsic function implementations
* Compiler based definations/tweaks
For a client application - as long as the requirements of the generated code are met, the prelude can be implemented by whatever mechanism is appropriate for the client. For example the implementation could be replaced with another implementation, or the prelude could contain all of the required text for compilation. Setting the prelude text can be achieved with the method on the global session...
```
/** Set the 'prelude' for generated code for a 'downstream compiler'.
@param passThrough The downstream compiler for generated code that will have the prelude applied to it.
@param preludeText The text added pre-pended verbatim before the generated source
That for pass-through usage, prelude is not pre-pended, preludes are for code generation only.
*/
void setDownstreamCompilerPrelude(SlangPassThrough passThrough, const char* preludeText);
```
The code that sets up the prelude for the test infrastucture and command line usage can be found in ```TestToolUtil::setSessionDefaultPrelude```. Essentially this determines what the absolute path is to `slang-cpp-prelude.h` is and then just makes the prelude `#include "the absolute path"`.
Limitations
===========
## WaveGetLaneIndex
This defaults to `threadIdx.x & SLANG_CUDA_WARP_MASK`. Depending on how the kernel is launched this could be incorrect.
There other ways to get lane index, for example using inline assembly. This mechanism though is apparently slower than the simple method used here.
There is potential to calculate the lane id using the [numthreads] markup in Slang/HLSL, but that also requires some assumptions of how that maps to a lane index.
Language aspects
================
# Arrays passed by Value
Slang follows the HLSL convention that arrays are passed by value. This is in contrast with CUDA where arrays follow C++ conventions and are passed by reference. To make generated CUDA follow this convention an array is turned into a 'FixedArray' struct type.
To get something more similar to CUDA/C++ operation the array can be marked in out or inout to make it passed by reference.
|
