CUDASafety-1.3¶
Device code should not contain unsupported type operators
Required inputs: IR
CUDA SAFETY 1.3 [safety.unsupported_type_operators] Device code should not contain unsupported type operators
Device code should not contain uses of dynamic_cast and typeid.
Scope: Device.
Audience: CUDA C++.
Category: Mandatory.
Hardware Applicability: All Compute Capabilities.
Rationale
Dynamic_cast and typeid are not support in the safety subset.
Example 1 (Bad)
# include <typeinfo> template<typename T, typename S> __global__ void kernel(T* t) { auto x = typeid(t); // Non-compliant: use of typeid operator in device code. auto p = dynamic_cast<S*>(t); // Non-compliant: use of dynamic_cast operator // in device code. }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_operator_in_device_code |
Do not use unsupported type operators in device code |
None |
False |
Options¶
This rule shares the following common options: exclude_in_macros, exclude_messages_in_system_headers, excludes, extend_exclude_to_macro_invocations, includes, justification_checker, languages, post_processing, provider, report_at, severity
The following places define options that affect this rule: Stylechecks, Analysis-GlobalOptions
forbidden_type_operators¶
forbidden_type_operators : set[bauhaus.ir.PIR_Class_Name] = {'Dynamic_Cast', 'Typeid_Operator'}