CUDASafety-3.1ΒΆ

Never invoke a kernel in a kernel launch configuration

Required inputs: IR

CUDA SAFETY 3.1 [safety.kernel.nested_configurations] Never invoke a kernel in a kernel launch configuration

No kernel should be launched while evaluating a kernel launch configuration.

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

Launching a kernel while evaluating a kernel launch configuration could lead to unexpected data structure conflicts.

Example 1 (Bad)
__global__ void kernel(){}

int foo() {
  kernel<<<1,1>>>(); // May overwrite any kernel setup state that has been saved before this invocation is encountered.
  return 1;
}

int bar() {
  kernel<<<foo(), 1>>>(); // invokes a kernel inside of foo()
  return 1;
}
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

nested_kernel_call

This kernel invocation happens during the evaluation of the launch configuration of another kernel invocation

None

False

Options