SPIR-V to ISPC: Convert GPU Compute to the CPU

By Jonathan Kennedy, Published: 06/01/2018, Last Updated: 06/05/2018

There is a growing trend within the games industry to move compute work to the graphics processing unit (GPU) resulting in engines and/or studios developing large portfolios of GPU compute shaders for many different compute tasks. However, there are times when it would be convenient to run those compute shaders on the CPU without having to re-invest in developing C/C++ variants of them. There are many reasons for doing this, including simple experimentation and debug, utilizing spare CPU cycles and encouraging CPU-based content scaling, CPU-based interaction with other CPU side game assets, for deterministic consistency of results, and so on.

To help address this opportunity while also utilizing the single instruction, multiple data (SIMD) vector units built into modern CPU cores, we have started developing a prototype translator, based on the open source Khronos* SPIRV-Cross project, that will take Standard Portable Intermediate Representation (SPIR-V1) as input and output kernels from Intel® SPMD Program Compiler (Intel® SPC), commonly known as ISPC. As shown in the usage sample below, ispc.exe takes C-style kernels and generates highly vectorized CPU object files targeting multiple ISAs such as Intel® Streaming SIMD Extensions (Intel® SSE), Intel® Advanced Vector Extensions 2 (Intel® AVX2), and Intel® Advanced Vector Extensions 512 (Intel® AVX-512).

The project should be considered as a starting point for conversion rather than a fully featured and performant solution. The project currently supports a subset of the standard SPIR-V intrinsics, built-ins and types, but it was designed to utilize a core performance feature of ispc.exe, which is the notion of uniform (scalar) or varying (vector) variables. This allows optimizations such as avoiding expensive and divergent vector branches if the test condition is scalar.

The code has been tested on a handful of shaders, such as the compute examples in Sascha Willems’ Vulkan* repository and the particle system compute shaders in Microsoft DirectX* 12 MiniEngine sample.

The code can be downloaded from our GitHub* repository and has currently only been tested on Windows* systems. The GitHub Readme contains more detailed documentation about the implementation and supported features.


glslangValidator.exe -H -V -o test.spv test.comp

spirv-cross.exe --ispc --output test.ispc test.spv

ispc.exe -O2 test.ispc -o test.ispc.obj -h test.ispc.h --target=avx2 --opt=fast-math

Example API Usage

ispc::raytracing_get_workgroup_size(workgroupSize[0], workgroupSize[1], workgroupSize[2]);

int32_t dispatch[3] = { textureComputeTarget.width / workgroupSize[0], textureComputeTarget.height / workgroupSize[1], 1 };
int32_t dispatch_count = dispatch[0] * dispatch[1];

concurrency::parallel_for<uint32_t>(0, dispatch_count, [&](uint32_t dispatchID)
    int32_t workgroupID[3] = { dispatchID % dispatch[0], dispatchID / dispatch[1], 0 };
    ispc::raytracing_dispatch_single(workgroupID, dispatch, planeCount, *pPlanes, sphereCount, *pSpheres, *pUBO, resultImage);

This project is open sourced under the original SPIRV-Cross Apache 2.0 license and we welcome any comments and contributions.

Further information on using ISPC in games can be found in the article, Use the Intel® SPMD Program Compiler for CPU Vectorization in Games.


  1. SPIR-V is the default shader language for Vulkan and can be generated from OpenGL* Shading Language (GLSL) and High-Level Shading Language (HLSL) shaders by tools such as the glslangValidator and shaderc compilers.

Product and Performance Information


Intel's compilers may or may not optimize to the same degree for non-Intel microprocessors for optimizations that are not unique to Intel microprocessors. These optimizations include SSE2, SSE3, and SSSE3 instruction sets and other optimizations. Intel does not guarantee the availability, functionality, or effectiveness of any optimization on microprocessors not manufactured by Intel. Microprocessor-dependent optimizations in this product are intended for use with Intel microprocessors. Certain optimizations not specific to Intel microarchitecture are reserved for Intel microprocessors. Please refer to the applicable product User and Reference Guides for more information regarding the specific instruction sets covered by this notice.

Notice revision #20110804