Spectral LLVM diagnostics

C++ family

This section describes new or enhanced diagnostics that apply to all C++ family code.

Improved diagnostic readability

• Print template <char... Cs> as a string rather than individual characters. This makes the diagnostic more readable.

CUDA

This section describes new or enhanced diagnostics that apply to CUDA code and SCALE code.

-fcuda-remarks

Enables a set of diagnostics that are useful for optimizing GPU programs. It turns on remarks to tell the user about:

• Empty kernels: Highlights kernels that were generated but that are empty. Such kernels can be generated accidentally by templated code.
• Kernels that use local memory: Highlights kernels that use local memory, and analyses the memory use to explain to the user why.

-Wcuda-attribute-ignored

Warns when attributes that apply to CUDA code, such as kernels, are ignored. For example where __host__ or __device__ is applied to a defaulted method (since this is automatically the case).

-Wcuda-external-static-member-variable

C++17 static inline member variables. This allows a static member variable to be defined (not just declared) in a class body. This avoids the need for a separate definition outside the class. For GPU code, this allows, for example:

class C
{
    static inline __shared__ int sharedCount;
};

In GPU code, it is rare to reference a variable from another translation unit. This means that a device-side static member variable that is neither defined in the same translation unit as the class, nor marked inline is probably a mistake. This warning alerts the programmer to such a mistake.

-Wcuda-function-arguments

This diagnostic prints warnings about bad arguments to certain CUDA language functions and intrinsics. This section documents the functions for which warnings can be generated.

__shfl and friends

The diagnostic checks that the warp size is in range and is a power of two. It supports the following functions:

• __shfl
• __shfl_down
• __shfl_up
• __shfl_xor
• __shfl_sync
• __shfl_down_sync
• __shfl_up_sync
• __shfl_xor_sync
• __nvvm_shfl_down_i32
• __nvvm_shfl_up_i32
• __nvvm_shfl_bfly_i32
• __nvvm_shfl_idx_i32
• __nvvm_shfl_down_f32
• __nvvm_shfl_up_f32
• __nvvm_shfl_bfly_f32
• __nvvm_shfl_idx_f32
• __nvvm_shfl_sync_down_i32
• __nvvm_shfl_sync_up_i32
• __nvvm_shfl_sync_bfly_i32
• __nvvm_shfl_sync_idx_i32
• __nvvm_shfl_sync_down_f32
• __nvvm_shfl_sync_up_f32
• __nvvm_shfl_sync_bfly_f32
• __nvvm_shfl_sync_idx_f32
• __nvvm_shfl_sync_down_pred_i32
• __nvvm_shfl_sync_up_pred_i32
• __nvvm_shfl_sync_bfly_pred_i32
• __nvvm_shfl_sync_idx_pred_i32