CUDA-3.2

Use function pointers only on the host or device where their address was taken

Required inputs: IR, StaticSemanticAnalysis

CUDA 3.2 [share.pointer.function] Use function pointers only on the host or device where their address was taken

Function pointers:

• Should only be ODR-used on the host or device its address was taken.
• Should not be copied between host and device or between devices.

This applies to all function pointers, regardless of execution space specifiers.

Scope: Host, Device.

Audience: CUDA C++, CUDA Libraries.

Category: Advisory.

Hardware Applicability: All Compute Capabilities.

Rationale

A heterogeneous function (__host__ __device__) is compiled for the host and device architectures and thus has a different address in the host program and in each device programs. A device function (__device__) may have multiple addresses. Any ODR-use of a function pointer on a host or device where it was not created has undefined behavior.

Example 1 (Bad)

# include <iostream>
# include <cassert>
# include "testTerminate.h"

void print(int32_t i) {
  std::cout << i << std::endl;
}

using element_function = void(*)(int32_t);

__global__ void for_each_thread(element_function ef) {
  auto const idx = blockIdx.x * blockDim.x + threadIdx.x;
  (*ef)(idx);
}

int main() {
  constexpr int32_t n = 32;
  // We incorrectly take the address of an implicit `__host__` function and
  // pass it to our kernel; the kernel will fail when it tries to call the
  // function pointer.
  for_each_thread<<<1, n>>>(&print);
  cudaError_t const error0 = cudaGetLastError();
  testTerminate(error0);
  cudaError_t const error1 = cudaDeviceSynchronize();
  assert(cudaSuccess == error1);
}

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
device_func_indirectly_called_from_host	This call is in host code but might call a device function	None	False
function_pointer_copy_across_execution_spaces	Function pointers should not be copied between host and device or between devices	None	False
host_func_indirectly_called_from_device	This call is in device code but might call a host function	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

cuda_memcpy_functions

cuda_memcpy_functions

Type: set[bauhaus.analysis.config.FunctionName]

Default: {'cudaMemcpy', 'cudaMemcpy2D', 'cudaMemcpy2DArrayToArray', 'cudaMemcpy2DAsync', 'cudaMemcpy2DFromArray', 'cudaMemcpy2DFromArrayAsync', 'cudaMemcpy2DToArray', 'cudaMemcpy2DToArrayAsync', 'cudaMemcpy3DPeer', 'cudaMemcpy3DPeerAsync', 'cudaMemcpyAsync', 'cudaMemcpyFromSymbol', 'cudaMemcpyFromSymbolAsync', 'cudaMemcpyPeer', 'cudaMemcpyPeerAsync', 'cudaMemcpyToSymbol', 'cudaMemcpyToSymbolAsync'}

Names of functions copying data between host and device or between devices. Pointer parameters to these functions are checked for not pointing to a function.