Qualcomm Snapdragon Game Developer Guides

https://docs.qualcomm.com/doc/80-78185-2/
Admin
Comprehensive developer documentation for optimizing games and applications on Qualcomm Snapdragon...
Tokens:20,155
Snippets:283
Trust Score:9.9
Update:1 month ago
Show doc for...
Tokens:
Raw
### Example ADB Driver Version Output

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Sample output format returned by the SurfaceFlinger GLES query.

```text
GLES: Qualcomm, Adreno (TM) 740, OpenGL ES 3.2 V@0676.0 (GIT@6f08ddb, I5e1ee3b043, 1669189803) (Date:11/23/22)
```

--------------------------------

### Vector Addition using ISPC

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/performance_optimization.html

Example of a struct-based vector addition implemented with ISPC syntax.

```cpp
struct FloatVector
{
    float V[4];
};

FloatVector operator+(const FloatVector& A, const FloatVector& B)
{
    FloatVector R;

    for (uniform int i=0; i< 4; i++)
    {
        R.V[i] = A.V[i] + B.V[i];
    }

    return R;
}
```

--------------------------------

### Implement Vulkan Driver Version Workarounds

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Macros and conditional logic to parse and check Adreno driver versions for applying specific workarounds.

```cpp
#define VK_VERSION_MAJOR(version) ((uint32_t)(version) >> 22)

#define VK_VERSION_MINOR(version) (((uint32_t)(version) >> 12) & 0x3FFU)

#define VK_VERSION_PATCH(version) ((uint32_t)(version) & 0xFFFU)

// Workarounds for Adreno driver 676
if (VK_VERSION_MINOR(version) == 676)
{
   // Guard for known Adreno issue on patch < 17 and driver 676
   if (VK_VERSION_PATCH(version) < 17)
   {
      doWorkaround();
   }
   // After patch 17 we know the issue was addressed by the vendor
   else
   {
      doNormalFlow();
   }
}
```

--------------------------------

### Initialize and execute ray query for shadow generation

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Performs ray-occlusion testing by initializing a ray query and traversing the acceleration structure to determine if a pixel is occluded by geometry.

```GLSL
// Initialize the query object
rayQueryEXT rayQuery;
rayQueryInitializeEXT(
    rayQuery,
    accelStructure,
    gl_RayFlagsTerminateOnFirstHitEXT,
    cullMask,
    WorldPos,
    minDistance,
    DirectionToLight,
    LightDistance);

    // Traverse the query -- do once for the first hit
    while(rayQueryProceedEXT(rayQuery))
    {
        // Hit something! Logic can be added here depending on the type of intersection
    }

    // Get the last intersection information
    if(rayQueryGetIntersectionTypeEXT(rayQuery, true) != gl_RayQueryCommittedIntersectionNoneEXT)
    {
         // Got an intersection -- this pixel is occluded, so retrieve distance
         float intersectionDistance = rayQueryGetIntersectionTEXT(rayQuery, true);

         // ...
    }
```

--------------------------------

### Configure Unreal Engine for Superluminal Profiling

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/tools.html

Enable frame pointers in Unreal Engine by modifying the BuildConfiguration.xml file to ensure full trace capture.

```xml
<BuildConfiguration>
    <bOmitFramePointers>false</bOmitFramePointers>
</BuildConfiguration>
```

--------------------------------

### Structure of Arrays (SoA) implementation

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/performance_optimization.html

C++ source code demonstrating a Structure of Arrays approach to pack four vectors for optimized computation.

```cpp
struct FloatVectorSoA
{
    float X[4];
    float Y[4];
    float Z[4];
    float W[4];
};
void sumSoAFloat(FloatVectorSoA& __restrict A, FloatVectorSoA& __restrict B, FloatVectorSoA& R)
{
    for (int i=0;i<4;i++)
    {
        R.X[i] = A.X[i] + B.X[i];
        R.Y[i] = A.Y[i] + B.Y[i];
        R.Z[i] = A.Z[i] + B.Z[i];
        R.W[i] = A.W[i] + B.W[i];
    }

}
```

--------------------------------

### Perform arithmetic operations with memory

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/architecture.html

Demonstrates loading, adding, and storing values using 64-bit and 32-bit registers.

```ARM Assembly
ldr         x8, [x8]
ldr         x9, [x9]
add         x8, x8, x9
str         x8, [x10]
```

```ARM Assembly
ldr         w8, [x8]
ldr         w9, [x9]
add         w8, w8, w9
str         w8, [x10]
```

--------------------------------

### Optimize non-opaque traversal with rayQueryProceedEXT

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Avoid using a while loop for rayQueryProceedEXT when using termination flags like gl_RayFlagsTerminateOnFirstHitEXT to simplify branch generation.

```glsl
rayQueryEXT rayQuery;
rayQueryInitializeEXT(rayQuery, rayTraceAS, gl_RayFlagsTerminateOnFirstHitEXT | gl_RayFlagsCullOpaqueEXT, cullMask, WorldPos, minDistance, DirectionToLight, LightDistance);

// Traverse the query. No need for a while(), since we want first hit or
// non-opaque intersections only
rayQueryProceedEXT(rayQuery));

// Determine if the shadow query collided
if(rayQueryGetIntersectionTypeEXT(rayQuery, true) != gl_RayQueryCommittedIntersectionNoneEXT)
{
    // Got an intersection == Shadow, do something
}
```

--------------------------------

### Loop Unrolling for GPR Optimization

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Demonstrates replacing a loop with explicit unrolled statements to potentially reduce GPR usage.

```GLSL
for (i = 0; i < 4; ++i) {
    diffuse += ComputeDiffuseContribution(normal, light[i]);
}
```

```GLSL
diffuse += ComputeDiffuseContribution(normal, light[0]);
diffuse += ComputeDiffuseContribution(normal, light[1]);
diffuse += ComputeDiffuseContribution(normal, light[2]);
diffuse += ComputeDiffuseContribution(normal, light[3]);
```

--------------------------------

### Vector Addition using NEON Intrinsics

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/performance_optimization.html

Demonstrates adding two vectors of four 32-bit integers using NEON intrinsic functions.

```cpp
uint32x4_t sumVector(uint32_t* A, uint32_t* B)
{
    uint32x4_t temp = vld1q_u32(A);
    uint32x4_t temp1 = vld1q_u32(B);
    uint32x4_t vec128 = vaddq_u32(temp, temp1);

    return vec128;
}
```

--------------------------------

### Query Driver Version via ADB

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Command to retrieve GLES driver information from the device shell.

```bash
adb shell dumpsys SurfaceFlinger | grep GLES
```

--------------------------------

### Configure Compute Shader Local Size

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Ensure the local workgroup size product is at least 64 to prevent the driver from forcing a command buffer flush and causing CPU-GPU stalls.

```GLSL
layout(local_size_x = 32, local_size_y = 1, local_size_z = 1) in; //workgroup size = local_size_x*local_size_y*local_size_z = 32x1x1 = 32 -- make sure this equation produces at least 64 to avoid a Cpu-wait-on-Gpu
```

--------------------------------

### Query Processor Count and Speed with Windows API

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/performance_optimization.html

Uses a chain of Windows APIs including GetActiveProcessorCount, GetLogicalProcessorInformationEx, and GetLogicalProcessorInformation to determine processor topology and performance metrics.

```cpp
#include <Windows.h>
#include <iostream>

extern "C" {
#include <Powrprof.h>
}

#include <vector>

#pragma comment(lib, "Powrprof.lib")

typedef struct _PROCESSOR_POWER_INFORMATION {
    ULONG  Number;
    ULONG  MaxMhz;
    ULONG  CurrentMhz;
    ULONG  MhzLimit;
    ULONG  MaxIdleState;
    ULONG  CurrentIdleState;
} PROCESSOR_POWER_INFORMATION, * PPROCESSOR_POWER_INFORMATION;

typedef BOOL(WINAPI* LPFN_GLPI)(PSYSTEM_LOGICAL_PROCESSOR_INFORMATION, PDWORD);
typedef BOOL(WINAPI* LPFN_GLPIEX)(LOGICAL_PROCESSOR_RELATIONSHIP, PSYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX, PDWORD);
typedef DWORD(WINAPI* LPFN_GAPC)(WORD);

#define ALL_PROCESSOR_GROUPS 0xffff

// Helper function to count set bits in the processor mask.
DWORD CountSetBits(ULONG_PTR bitMask)
{
    DWORD LSHIFT = sizeof(ULONG_PTR) * 8 - 1;
    DWORD bitSetCount = 0;
    ULONG_PTR bitTest = (ULONG_PTR)1 << LSHIFT;
    DWORD i;

    for (i = 0; i <= LSHIFT; ++i)
    {
        bitSetCount += ((bitMask & bitTest) ? 1 : 0);
        bitTest /= 2;
    }

    return bitSetCount;
}

DWORD GetInstalledProcessorCount()
{
    // on Windows 7 and later, use GetActiveProcessorCount() ...

    LPFN_GAPC gapc = (LPFN_GAPC)GetProcAddress(GetModuleHandle(TEXT("kernel32")), "GetActiveProcessorCount");
    if (gapc)
        return gapc(ALL_PROCESSOR_GROUPS);

    // on Vista and later, try GetLogicalProcessorInformationEx() next ...

    LPFN_GLPIEX glpiex = (LPFN_GLPIEX)GetProcAddress(GetModuleHandle(TEXT("kernel32")), "GetLogicalProcessorInformationEx");
    if (glpiex)
    {
        std::vector<BYTE> buffer;
        PSYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX info = NULL;
        DWORD bufsize = 0;

        // not using RelationGroup because it does not return accurate info under WOW64...
        while (!glpiex(RelationProcessorCore, info, &bufsize))
        {
            if (GetLastError() != ERROR_INSUFFICIENT_BUFFER)
                return 0;

            buffer.resize(bufsize);
            info = (PSYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX)&buffer[0];
        }

        DWORD logicalProcessorCount = 0;

        while (bufsize >= sizeof(SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX))
        {
            // for RelationProcessorCore, info->Processor.GroupCount is always 1...
            logicalProcessorCount += CountSetBits(info->Processor.GroupMask[0].Mask);
            bufsize -= info->Size;
            info = (PSYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX)(((LPBYTE)info) + info->Size);
        }

        return logicalProcessorCount;
    }

    // on XP and later, try GetLogicalProcessorInformation() next...

    LPFN_GLPI glpi = (LPFN_GLPI)GetProcAddress(GetModuleHandle(TEXT("kernel32")), "GetLogicalProcessorInformation");
    if (glpi)
    {
        std::vector<BYTE> buffer;
        PSYSTEM_LOGICAL_PROCESSOR_INFORMATION info = NULL;
        DWORD bufsize = 0;

        while (!glpi(info, &bufsize))
        {
            if (GetLastError() != ERROR_INSUFFICIENT_BUFFER)
                return 0;

            buffer.resize(bufsize);
            info = (PSYSTEM_LOGICAL_PROCESSOR_INFORMATION)&buffer[0];
        }

        DWORD logicalProcessorCount = 0;

        while (bufsize >= sizeof(SYSTEM_LOGICAL_PROCESSOR_INFORMATION))
        {
            if (info->Relationship == RelationProcessorCore)
            {
                // A hyperthreaded core supplies more than one logical processor.
                logicalProcessorCount += CountSetBits(info->ProcessorMask);
            }

            bufsize -= sizeof(SYSTEM_LOGICAL_PROCESSOR_INFORMATION);
            ++info;
        }

        return logicalProcessorCount;
    }

    // fallback to GetSystemInfo() last ...

    SYSTEM_INFO si = { 0 };
    GetSystemInfo(&si);

    return si.dwNumberOfProcessors;
}

int main(int argc, char* argv[])
{
    DWORD dwNumProcessors = GetInstalledProcessorCount();

    std::vector<PROCESSOR_POWER_INFORMATION> a(dwNumProcessors);
    DWORD dwSize = sizeof(PROCESSOR_POWER_INFORMATION) * dwNumProcessors;
    std::cout << "Number of processors: " << dwNumProcessors << std::endl;
    CallNtPowerInformation(ProcessorInformation, NULL, 0, &a[0], dwSize);
    std::cout << "Processor\tMHz\tMaxMHz\tMHz Limit\tIdle State\tMax Idle State" << std::endl;
    for (int i = 0; i < dwNumProcessors; i++)
    {
        std::cout << a[i].Number << "\t\t";
        std::cout << a[i].CurrentMhz << '\t';
        std::cout << a[i].MaxMhz << '\t';
        std::cout << a[i].MhzLimit << "\t\t\t";
        std::cout << a[i].CurrentIdleState << "\t\t";
        std::cout << a[i].MaxIdleState << "\t" << std::endl;
    }



    a.clear();

    system("pause");
    return 0;
}
```

--------------------------------

### Pack Scalar Constants into Vectors

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Group scalar constants into vectors to improve hardware fetch effectiveness and enable efficient 'mad' instructions.

```GLSL
float scale, bias;
vec4 a = Pos * scale + bias;
```

```GLSL
vec2 scaleNbias;
vec4 a = Pos * scaleNbias.x + scaleNbias.y;
```

--------------------------------

### Set EGLSurface format for HDR

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Configures the EGL surface to use 10-bit color depth per channel with 2-bit alpha.

```cpp
EGLConfig EGLConfigList [1]; int ConfigAttributes [] = {
    EGL_RED_SIZE,10
    EGL_GREEN_SIZE,10
    EGL_BLUE_SIZE,10
    EGL_ALPHA_SIZE,2
    EGL_COLOR_COMPONENT_TYPE_EXT,
EGL_COLOR_COMPONENT_TYPE_FIXED_EXT, EGL_NONE};

eglChooseConfig (eglDisplay,ConfigAttributes,EGLConfigList,1, eglNumConfigs);
```

--------------------------------

### Set EGLSurface color space

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Applies the BT2020 PQ color space to the window surface during creation.

```cpp
EGLint attribs[] = {EGL_GL_COLORSPACE_KHR,EGL_GL_COLORSPACE_BT2020_PQ_EXT,EGL_NONE };
EGLSurface eglSurface=eglCreateWidowSurface(eglDisplay,eglConfigParam, InWindow, attribs);
```

--------------------------------

### Set SMPTE 2086 metadata

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Updates the EGL surface with display chromaticity and luminance metadata using scaling factors.

```cpp
EGLint SurfaceAttribs [] = {
EGL_SMPTE2086_DISPLAY_PRIMARY_RX_EXT,       EGL_SMPTE2086_DISPLAY_PRIMARY_RY_EXT,
    EGL_SMPTE2086_DISPLAY_PRIMARY_GX_EXT,
    EGL_SMPTE2086_DISPLAY_PRIMARY_GY_EXT,
    EGL_SMPTE2086_DISPLAY_PRIMARY_BX_EXT,
    EGL_SMPTE2086_DISPLAY_PRIMARY_BY_EXT,
    EGL_SMPTE2086_WHITE_POINT_X_EXT,
    EGL_SMPTE2086_WHITE_POINT_Y_EXT,
    EGL_SMPTE2086_MAX_LUMINANCE_EXT,
    EGL_SMPTE2086_MIN_LUMINANCE_EXT
};
static const DisplayChromacities DisplayChromacityList [] = { {{0.70800f, 0.29200f, 0.17000f, 0.79700f, 0.13100f, 0.04600f, 0.31270f, 0.32900f}}, // DG_Rec2020 };
for (uint32_t i = 0; i < 8; i++)
{
    eglSurfaceAttrib(PImplData->eglDisplay,eglSurface, SurfaceAttribs[i],EGLint(DisplayChromacityList[0].ChromaVals[i]* EGL_METADATA_SCALING_EXT));
}
```

--------------------------------

### SoA vector addition assembly

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/performance_optimization.html

Assembly output for the SoA implementation. MSVC fails to vectorize due to perceived loop-carried data dependencies.

```assembly
ldp     q24, q29, [x0]
ldp     q28, q25, [x1]
ldp     q30, q31, [x0, 32]
ldp     q26, q27, [x1, 32]
fadd    v28.4s, v28.4s, v24.4s
fadd    v29.4s, v29.4s, v25.4s
fadd    v30.4s, v30.4s, v26.4s
fadd    v31.4s, v31.4s, v27.4s
stp     q28, q29, [x2]
stp     q30, q31, [x2, 32]
ret
```

```assembly
ldp     q0, q3, [x1]
ldp     q1, q2, [x0]
fadd    v0.4s, v1.4s, v0.4s
ldp     q1, q5, [x0, #32]
fadd    v2.4s, v2.4s, v3.4s
ldp     q4, q3, [x1, #32]
fadd    v1.4s, v1.4s, v4.4s
fadd    v3.4s, v5.4s, v3.4s
stp     q0, q2, [x2]
stp     q1, q3, [x2, #32]
ret
```

```assembly
 add         x9,x2,#0x20
 add         x8,x1,#0x10
 sub         x13,x0,x1
 sub         x12,x2,x1
 sub         x11,x0,x2
 mov         x10,#4
|$LL22@S|
 ldur        s17,[x8,#-0x10]
 sub         x10,x10,#1
 ldr         s16,[x0],#4
 fadd        s16,s17,s16
 ldr         s17,[x13,x8]
 stur        s16,[x9,#-0x20]
 ldr         s16,[x8]
 fadd        s16,s17,s16
 ldr         s17,[x11,x9]
 str         s16,[x12,x8]
 ldr         s16,[x8,#0x10]
 fadd        s16,s17,s16
 ldr         s17,[x8,#0x20]
 add         x8,x8,#4
 str         s16,[x9],#4
 ldr         s16,[x0,#0x2C]
 fadd        s16,s17,s16
 str         s16,[x9,#0xC]
 cbnz        x10,|$LL22@sumS|
 ret
```

--------------------------------

### Add immediate value

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/architecture.html

Adds an immediate constant value to a register using the # prefix.

```ARM Assembly
add         x4, x3, #17
```

--------------------------------

### Generated NEON Assembly for Vector Addition

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/performance_optimization.html

Assembly output generated by ISPC targeting ARM64 NEON instructions.

```asm
ldp     q0, q1, [x0]
ldp     q2, q3, [x1]
ldp     q4, q5, [x0, #32]
ldp     q6, q7, [x1, #32]
fadd    v0.4s, v0.4s, v2.4s
fadd    v1.4s, v1.4s, v3.4s
fadd    v2.4s, v4.4s, v6.4s
fadd    v3.4s, v5.4s, v7.4s
ret
```

--------------------------------

### Assembly Output for NEON Intrinsics

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/performance_optimization.html

Comparison of generated assembly code for the NEON intrinsic vector addition across different compilers.

```asm
ldr     q0, [x0]
ldr     q31, [x1]
add     v0.4s, v0.4s, v31.4s
ret
```

```asm
ldr     q0, [x0]
ldr     q1, [x1]
add     v0.4s, v1.4s, v0.4s
ret
```

```asm
ldr         q17,[x1]
ldr         q16,[x0]
add         v0.4s,v16.4s,v17.4s
ret
```

--------------------------------

### Enable Early Fragment Tests in GLSL

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/mobile_best_practices.html

Use this qualifier in your GLSL shader to activate EarlyFragmentTest, which enables fast-z mode for depth-only renderpasses using indirect draws and SSBOs.

```GLSL
layout(early_fragment_tests) in;
```

--------------------------------

### Branch instructions

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/architecture.html

Performs unconditional and conditional branching based on register comparisons.

```ARM Assembly
b           label
```

```ARM Assembly
ldr         w8, [x9]
cmp         w8, #10
ble         label
```

--------------------------------

### Load data from memory

Source: https://docs.qualcomm.com/doc/80-78185-2/topic/architecture.html

Loads data from a memory address stored in a register into a destination register.

```ARM Assembly
ldr         x1, [x7]
```
Tokens:
Raw