CUDASafety-1.4

Device code should not contain unsupported types

Required inputs: IR

CUDA SAFETY 1.4 [safety.unsupported_type] Device code should not contain unsupported types

The long double type shall not be used in device code.

Scope: Device.
Audience: CUDA C++.
Category: Mandatory.
Hardware Applicability: All Compute Capabilities.
Rationale

The long double type is not supported in the CUDA safety subset.

Example 1 (Bad)
template<typename T>

__global__ void kernel(long double) { // non-compliant: use of long double in
// kernel arguments.
  unsigned i = sizeof(long double); // non-compliant: use of long double in
// a constant expression.
}

int main() {
  kernel<long double><<<1,1>>>(0); // non-compliant: instantiation of
// a device function is dependent on
// long double.
}
Excerpt from NVIDIA CUDA C++ Guidelines for robust and safety-critical programming, Version 3.0.1, Copyright (C) 2018-2023 NVIDIA Corporation.

Possible Messages

Key

Text

Severity

Disabled

unsupported_type_in_device_code

Do not use unsupported types in device code

None

False

Options

forbidden_types

forbidden_types : set[bauhaus.ir.LIR_Class_Name] = {'C_Long_Double_Type'}

Forbidden types, given as LIR class names.