CUDA-3.5

Do not create aliases between restricted pointer arguments

Required inputs: IR, StaticSemanticAnalysis

CUDA 3.5 [restricted_pointers] Do not create aliases between restricted pointer arguments

Do not create aliases when sending addresses into functions via restricted pointer arguments.

Scope: Device.

Audience: CUDA C++.

Category: Required.

Hardware Applicability: All Compute Capabilities.

Rationale

The __restrict__ keyword gives the compiler leeway to perform optimizations that are only valid if the restricted pointer arguments are not aliases when they enter a function.

Example 1 (Bad)

__device__ void foo(int& a , int * __restrict__ b) {
  a = *b;
}

__device__ void bar(int a) {
  foo(a, &a); // non-compliant: the same object is accessed by a reference and
// a restricted pointer.
}

Example 2 (Good)

__device__ void foo(int& a , int * b) {
  a = *b;
}

__device__ void bar(int a) {
  foo(a, &a);
}

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
restrict_pointer_assignment_undefined_behavior	Avoid undefined behavior, assignment between restrict-qualified pointers in the same scope.	None	False
restrict_undefined_behavior	Avoid undefined behavior, assignment to overlapping objects.	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

functions

functions

Type: dict[bauhaus.analysis.config.FileGlobPattern, dict[str, typing.Tuple[int, int, int]]]

Default:

{
   '': {
      '__builtin___memcpy_chk': (0, 1, 2)
   },
   'string.h': {
      'memcpy': (0, 1, 2),
      'memcpy_s': (0, 2, 3),
      'strncat': (0, 1, 2),
      'strtok_s': (0, 2, 3)
   }
}

Mapping from header filename globbings to mappings from function names to tuples (destination, source, size) containing argument positions.