CUDA-1.7ΒΆ

Be consistent with execution space and kernel function specifiers across all declarations

Required inputs: IR

CUDA 1.7 [specifiers.consistency] Be consistent with execution space and kernel function specifiers across all declarations

All declarations of a function should have the same execution space (e.g. __host__ and __device__) and kernel (e.g. __global__) specifiers.

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

Differing specifiers on different declarations of the same function can make it unclear which specifiers are applicable to the function and which are disregarded. No diagnostic is guaranteed for mismatches of execution space and kernel specifiers across different declarations. This can lead to subtle violations of the sharing guidelines, which can lead to undefined behavior.

Example 1 (Bad)
# include "testTerminate.h"

__device__ void foo();

__host__ void foo(){} // ill-formed, execution space specifiers don't match previous declaration.

__host__ __device__ void bar();

void bar() {} // ill-formed

template <typename F>

__global__ void invoke(F f) {
  f();
}

int main() {
  invoke<<<1,1>>>(foo); // This required to get any hits on issues.
  testTerminate(cudaGetLastError());
  testTerminate(cudaDeviceSynchronize());
}
Example 2 (Bad)
__global__ void foo();

void foo() { } // ill-formed, execution space specifiers don't match previous declaration.
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

execution_space_mismatch

Mismatch in execution space specifier (__host__/__device__)

None

False

implicit_vs_explicit_host

__host__ specifier is explicit in one declaration, but implicit in the other

None

False

kernel_mismatch

Mismatch in kernel specifier (__global__)

None

False

Options