CUDA-4.3ΒΆ

No virtual arguments to kernel functions

Required inputs: IR

CUDA 4.3 [kernel.virtual_kernel_parameters] No virtual arguments to kernel functions

No kernel argument or subobject of a kernel argument should have a virtual base class or virtual function.

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

The virtual tables (e.g. vtables) differ between the host and the device, since kernel arguments are captured on the host and sent to the device this can result in the wrong address being used to access a virtual function.

Example 1 (Bad)
struct S {
  virtual __host__ __device__ ~S(){};
};

__global__ void kernel(S s) {} // Non-compliant: kernel argument has a virtual
// function.

struct T : virtual S {
  __host__ __device__ T(){};
};

__global__ void kernel(T t) {} // Non-compliant: kernel argument has a virtual
// base class.
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

cuda_kernel_possible_virtual_kernel_parameters

Kernel argument or subobject of a kernel argument may have a virtual base class or virtual function

None

False

cuda_kernel_virtual_kernel_parameters

Kernel argument or subobject of a kernel argument has a virtual base class or virtual function

None

False

Options