CUDASafety-2.1ΒΆ

Do not use incorrect deduced execution spaces

Required inputs: IR

CUDA SAFETY 2.1 [execution_space.deduced] Do not use incorrect deduced execution spaces

Use of deduced execution spaces should be done correctly.

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

The use of deduced execution spaces may result in capturing function pointer on one device and invocation on a different device.

Example 1 (Bad)
struct S { int i; };

__device__ void g(S a, S b) {
  S& (S::*mfptr)(const S&) = &S::operator=; // address of host operator= may be captured here.
  (b.*mfptr)(a); // non-compliant: implicit execution space is used for indirect
// call instead of direct invocation.
}
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

incorrect_deduced_host_exec_space

Target function has deduced host execution space, but is used in device code

None

False

Options