Slang CUDA Target Support

Making it easier to work with shaders

Slang CUDA Target Support

Slang has preliminary support for producing CUDA source, and PTX binaries using NVRTC.

NOTE! NVRTC is only available for 64-bit operating systems. On Windows Visual Studio make sure you are compiling for ‘x64’ and/or use 64 bit Slang binaries.

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.
  • When using a TextureArray.Sample (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 appropriate launches
  • CUDA ‘surfaces’ are used for textures which are read/write (aka RWTexture).

The following are a work in progress or not implemented but are planned to be so in the future

  • Some resource types remain unsupported, and not all methods on all types are supported

How it works

For producing PTX binaries Slang uses NVRTC. NVRTC dll/shared library has to be available to Slang (for example in the appropriate PATH for example) for it to be able to produce PTX.

The NVRTC compiler can be accessed directly via the pass through mechanism and is identified by the enum value 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. 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.

Locating NVRTC

Finding NVRTC can require some nuance if a specific version is required. On the command line the -nvrtc-path option can be used to set the path to NVRTC. Also spProcessCommandLineArguments/processCommandLineArguments with -nvrtc-path or setDownstreamCompilerPath with SLANG_PASS_THROUGH_NVRTC can be used to set the location and/or name of NVRTC via the API.

Important points of note are

  • The name of the shared library should not include any extension (such as .dll/.so/.dynlib) or prefix (such as lib).
  • The path also doesn’t have to be path, it can just be the shared library name. Doing so will mean it will be searched for by whatever the default mechanism is on the target.
  • If a path and/or name is specified for NVRTC - this will be the only version searched for.

If a path/name is not specified for NVRTC, Slang will attempt to load a shared library called nvrtc. For non Windows targets this should be enough to find and load the latest version.

On Windows NVRTC dlls have a name the contains the version number, for example nvrtc64_102_0.dll. This will lead to the load of just nvrtc to fail. One approach to fix this is to place the NVRTC dll and associated files in the same directory as slang-compiler.dll, and rename the main dll to nvrtc.dll. Another approach is to specify directly on the command line the name including the version, as previously discussed. For example

-nvrtc-path nvrtc64_102_0

will load NVRTC 10.2 assuming that version of the dll can be found via the normal lookup mechanism.

On Windows if NVRTC is not loadable directly as ‘nvrtc’ Slang will attempt to search for the newest version of NVRTC on your system. The places searched are…

  • The instance directory (where the slang-compiler.dll and/or program exe is)
  • The CUDA_PATH enivonment variable (if set)
  • Directories in PATH that look like a CUDA installation.

If a candidate is found via an earlier mechanism, subsequent searches are not performed. If multiple candidates are found, Slang tries the newest version first.

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;

Texture

Read only textures will be bound as the opaque CUDA type CUtexObject. This type is the combination of both a texture AND a sampler. This is somewhat different from HLSL, where there can be separate SamplerState variables. This allows access of a single texture binding with different types of sampling.

If code relies on this behavior it will be necessary to bind multiple CtexObjects with different sampler settings, accessing the same texture data.

Slang has some preliminary support for TextureSampler type - a combined Texture and SamplerState. To write Slang code that can target CUDA and other platforms using this mechanism will expose the semantics appropriately within the source.

Load on a read-only texture is supported for Texture1D, Texture2D, Texture3D, and the 1D/2D array forms (it lowers to the tex*fetch_int<T> prelude templates); the mip map selection argument is ignored, because these fetch paths read the base level only. RWTexture also allows Load on its supported dimensions.

Reading a half-typed read-only texture (e.g. Texture2D<half4>) is not supported on CUDA. Load lowers to the tex*fetch_int<T> CUDA-prelude templates and SampleLevel lowers to the tex*Lod<T> CUDA runtime built-ins; neither is defined for __half types (both are float/uint/int only), so a half texel is rejected at compile time with an error. Use a float texture instead - CUDA widens f16-format data to f32 on read for free, and you can narrow the result back to half in-shader if needed.

RWTexture

RWTexture types are converted into CUsurfObject type.

In regular CUDA it is not possible to do a format conversion on an access to a CUsurfObject. Slang does add support for hardware write conversions where they are available. To enable the feature it is necessary to attribute your RWTexture with format. For example

[format("rg16f")]
RWTexture2D<float2> rwt2D_2;

The format names used are the same as for GLSL layout format types. If no format is specified Slang will assume that the format is the same as the type specified.

Note that the format attribution is on variables/parameters/fields and not part of the type system. This means that if you have a scenario like…

[format(rg16f)]
RWTexture2d<float2> g_texture;

float2 getValue(RWTexture2D<float2> t)
{
    return t[int2(0, 0)];
}

void doThing()
{
    float2 v = getValue(g_texture);
}

Even getValue will receive t without the format attribute, and so will access it, presumably erroneously. A workaround for this specific scenario would be to attribute the parameter

float2 getValue([format("rg16f")] RWTexture2D<float2> t)
{
    return t[int2(0, 0)];
}

This will only work correctly if getValue is called with a t that has that format attribute. As it stands no checking is performed on this matching so no error or warning will be produced if there is a mismatch.

There is limited software support for doing a conversion on reading. Currently this only supports only 1D, 2D, 3D RWTexture, backed with half1, half2 or half4. For this path to work NVRTC must have the cuda_fp16.h and associated files available. Please check the section on Half Support.

If hardware read conversions are desired, this can be achieved by having a Texture that uses the surface of a RWTexture. Using the Texture not only allows hardware conversion but also filtering.

It is also worth noting that CUsurfObjects in CUDA are NOT allowed to have mip maps.

By default surface access uses cudaBoundaryModeZero, this can be replaced using the macro SLANG_CUDA_BOUNDARY_MODE in the CUDA prelude. For HW format conversions the macro SLANG_PTX_BOUNDARY_MODE. These boundary settings are in effect global for the whole of the kernel.

SLANG_CUDA_BOUNDARY_MODE can be one of

  • cudaBoundaryModeZero causes an execution trap on out-of-bounds addresses
  • cudaBoundaryModeClamp stores data at the nearest surface location (sized appropriately)
  • cudaBoundaryModeTrap drops stores to out-of-bounds addresses

SLANG_PTX_BOUNDARY_MODE can be one of trap, clamp or zero. In general it is recommended to have both set to the same type of value, for example cudaBoundaryModeZero and zero.

Sampler

Samplers are in effect ignored in CUDA output. Currently we do output a variable SamplerState, but this value is never accessed within the kernel and so can be ignored. More discussion on this behavior is in Texture section.

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 infrastructure 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".

Precompiled headers for repeated NVRTC compilation

NVRTC re-parses the prelude on every compilation. NVRTC 12.8 and later can precompile it once per compatible set of compile options and leading directives within a process and reuse it across later compilations via automatic precompiled headers (the -pch compile option), which cuts the per-shader compile time for the second and later compilations in a process. This works because the prelude reaches NVRTC as a leading #include (see above), so the whole prelude falls before NVRTC’s precompiled-header “stop point” — the first token that is not part of a preprocessing directive.

On the -target ptx path Slang drives NVRTC itself and passes -pch automatically on NVRTC 12.8+, but only when the prelude is presented as a leading #include — which is the case for the command-line tools and test infrastructure (they install the include-form prelude, as described above). A host application using the library API with the default embedded prelude, or a custom prelude whose trimmed text does not begin with #include, does not get -pch on this path, because with a verbatim-text prelude NVRTC’s stop point falls near the top of the prelude rather than after it, so the precompiled header captures too little to help; such an application should install an include-form prelude (via setLanguagePrelude) if it wants the speedup.

On the -target cuda path the application drives NVRTC, so it must opt in. Present the prelude to NVRTC as a leading #include (as above, not inlined text), pass -pch, and rely on the process-global precompiled-header heap, whose default already covers the prelude. Grow the heap only reactively: if nvrtcGetPCHCreateStatus(prog) returns NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED, read the needed size with nvrtcGetPCHHeapSizeRequired(prog, &size), set it with nvrtcSetPCHHeapSize(size), and recompile with a fresh nvrtcProgram so the enlarged heap is used. Do not call nvrtcSetPCHHeapSize before every compilation, because it frees any existing precompiled header and defeats reuse.

Per NVRTC’s documented PCH semantics, NVRTC keys the precompiled header on the compilation options and the translation unit’s leading directive text, and rebuilds it when either changes. It does not track the contents of #included files, so an application that keeps the same #include line while changing the prelude header’s contents on disk within one process must invalidate reuse itself (for example by changing the prelude text passed to NVRTC) rather than relying on same-path header edits being noticed.

Half Support

Slang supports the half/float16 types on CUDA. To do so NVRTC must have access to the cuda_fp16.h and cuda_fp16.hpp files that are typically distributed as part of the CUDA SDK. When Slang detects the use of half in source, it will define SLANG_CUDA_ENABLE_HALF when slang-cuda-prelude.h is included. This will in turn try to include cuda_fp16.h and enable extra functionality within the prelude for half support.

Slang tries several mechanisms to locate cuda_fp16.h when NVRTC is initiated. The first mechanism is to look in the include paths that are passed to Slang. If cuda_fp16.h can be found in one of these paths, no more searching will be performed.

If this fails, the path where NVRTC is located will be searched. In that path “include” and “CUDA/include” paths will be searched. This is probably most suitable for Windows based targets, where NVRTC dll is placed along with other binaries. The “CUDA/include” path is used to try and make clear in this scenario what the contained files are for.

If this fails Slang will look for the CUDA_PATH environmental variable, as is typically set during a CUDA SDK installation.

If this fails - the prelude include of cuda_fp16.h will most likely fail on NVRTC invocation.

CUDA has the __half and __half2 types defined in cuda_fp16.h. The __half2 can produce results just as quickly as doing the same operation on __half - in essence for some operations __half2 is SIMD like. Slang’s prelude takes advantage of this for 2-wide half vectors: arithmetic on half2 maps to the native __half2 intrinsics (__hadd2, __hsub2, __hmul2, __h2div, __hneg2) rather than being done component-by-component. On architectures with native packed-half support (compute capability 5.3 and above) addition, subtraction, multiplication and negation compile to a single packed instruction; on older architectures CUDA’s own intrinsics fall back to two scalar operations. Division always uses two scalar operations, as CUDA has no packed half divide. Wider half vectors (half3 and half4) are backed by their own structs and their arithmetic is performed element-wise.

Since Slang supports up to 4 wide vectors Slang has to build on CUDAs half support. The types __half3 and __half4 are implemented in slang-cuda-prelude.h for this reason. It is worth noting that __half3 is declared as three separate __half members (x, y, z) with a 4-byte alignment (__align__(4)). Because of that alignment, __half3 is actually 8 bytes, rather than the 6 bytes that might be expected.

One area where this optimization isn’t fully used is in comparisons - as in effect Slang treats all the vector/matrix half comparisons as if they are scalar. This could be perhaps be improved on in the future. Doing so would require using features that are not directly available in the CUDA headers.

Fast Math

When compiling with -fp-mode fast, Slang emits #define SLANG_CUDA_ENABLE_FAST_MATH ahead of the slang-cuda-prelude.h include. This redirects the single-precision transcendental wrappers that have a fast approximate CUDA intrinsic (sin, cos, sincos, tan, log, log2, log10, exp, pow) to their __*f forms (__sinf, __cosf, and so on), which can substantially reduce the emitted instruction count at the cost of reduced accuracy. Among the half-precision counterparts of these functions, tan and pow promote to float and route through the float wrappers, so they follow the same redirect; the remaining seven (sin, cos, sincos, log, log2, log10, exp) use native __half intrinsics (hsin, hcos, hlog, hlog2, hlog10, hexp — sincos composes hsin+hcos rather than having its own intrinsic) and are unaffected. This trade-off is why it is opt-in; without -fp-mode fast every wrapper keeps its default form.

Wrappers with no fast intrinsic (exp2, atan2, asin, acos, atan, sqrt, and the hyperbolic functions) are not redirected by this define — they keep their existing standard-library implementation — as do all double-precision (F64_*) wrappers, since CUDA provides no approximate double intrinsics.

This define is intentionally transcendental-only and narrower than nvcc’s --use_fast_math: it does not enable approximate division, sqrt, flush-to-zero, or FMA contraction. Its value is on the -target cuda path, where you compile the emitted .cu with your own toolchain — the -target ptx path is compiled by Slang’s NVRTC invocation, which already receives --use_fast_math under -fp-mode fast. The redirect is also module-global (it is a single front-matter #define), so it reflects the module’s -fp-mode and does not honor a per-function floating-point-mode override, such as the Fast override auto-diff places on generated derivative functions.

Wave Intrinsics

There is broad support for HLSL Wave intrinsics, including support for SM 6.5 intrinsics.

Most Wave intrinsics will work with vector, matrix or scalar types of typical built in types - uint, int, float, double, uint64_t, int64_t.

The support is provided via both the Slang core module as well as the Slang CUDA prelude found in ‘prelude/slang-cuda-prelude.h’. Many Wave intrinsics are not directly applicable within CUDA which supplies a more low level mechanisms. The implementation of most Wave functions work most optimally if a ‘Wave’ where all lanes are used. If all lanes from index 0 to pow2(n) -1 are used (which is also true if all lanes are used) a binary reduction is typically applied. If this is not the case the implementation fallsback on a slow path which is linear in the number of active lanes, and so is typically significantly less performant.

For more a more concrete example take

int sum = WaveActiveSum(...);

When computing the sum, if all lanes (32 on CUDA), the computation will require 5 steps to complete (2^5 = 32). If say just one lane is not being used it will take 31 steps to complete (because it is now linear in amount of lanes). So just having one lane disabled required 6 times as many steps. If lanes with 0 - 15 are active, it will take 4 steps to complete (2^4 = 16).

In the future it may be possible to improve on the performance of the ‘slow’ path, however it will always remain the most efficient generally for all of 0 to pow2(n) - 1 lanes to be active.

It is also worth noting that lane communicating intrinsics performance will be impacted by the ‘size’ of the data communicated. The size here is at a minimum the amount of built in scalar types used in the processing. The CUDA language only allows direct communication with built in scalar types.

Thus

int3 v = ...;
int3 sum = WaveActiveSum(v);

Will require 3 times as many steps as the earlier scalar example just using a single int.

WaveGetLaneIndex

WaveGetLaneIndex computes the lane index by linearizing the thread index across all dimensions and masking to the warp size:

((threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x) & SLANG_CUDA_WARP_MASK

This handles multi-dimensional thread blocks correctly as long as the linearized thread index maps to consecutive warp lanes, which is the standard CUDA thread-to-lane mapping.

Alternatively, defining SLANG_USE_ASM_LANE_ID before including the CUDA prelude switches to an inline PTX assembly implementation (mov.u32 %0, %laneid). The assembly version is always correct regardless of launch configuration, but is slower than the arithmetic default.

Unsupported Intrinsics

  • Intrinsics which only work in pixel shaders
    • QuadXXXX intrinsics

OptiX Support

Slang supports OptiX for raytracing. To compile raytracing programs, NVRTC must have access to the optix.h and dependent files that are typically distributed as part of the OptiX SDK. When Slang detects the use of raytracing in source, it will define SLANG_CUDA_ENABLE_OPTIX when slang-cuda-prelude.h is included. This will in turn try to include optix.h.

Slang tries several mechanisms to locate optix.h when NVRTC is initiated. The first mechanism is to look in the include paths that are passed to Slang. If optix.h can be found in one of these paths, no more searching will be performed.

If this fails, the default OptiX SDK install locations are searched. On Windows this is %{PROGRAMDATA}\NVIDIA Corporation\OptiX SDK X.X.X\include. On Linux this is ${HOME}/NVIDIA-OptiX-SDK-X.X.X-suffix.

If OptiX headers cannot be found, compilation will fail.

Limitations

Some features are not available because they cannot be mapped with appropriate behavior to a target. Other features are unavailable because of resources to devote to more unusual features.

  • Not all Wave intrinsics are supported
  • There is not complete support for all methods on ‘objects’ like textures etc.
  • Does not currently support combined ‘TextureSampler’. A Texture behaves equivalently to a TextureSampler and Samplers are ignored.
  • Half type is not currently supported
  • GetDimensions supports querying texture width, height, and depth. Overloads taking a mip level are only available in OptiX ray-tracing stages, where they also return the mip-level count. Array-size and sample-count output parameters are currently written as zero.

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.