CUDA-3.4

Only use one __shared__ array-of-unknown-bound declaration in a kernel function

Required inputs: IR

CUDA 3.4 [shared.object.dynamic_arrays] Only use one __shared__ array-of-unknown-bound declaration in a kernel function

At most one __shared__ array-of-unknown-bound declaration shall be used during the evaluation of a kernel function.

Scope: Device.

Audience: CUDA C++.

Category: Mandatory.

Hardware Applicability: All Compute Capabilities.

Rationale

Only one dynamic shared array base address can be given to a kernel. This means that only one dynamic shared array will be properly constructed, all others have undefined behavior.

Example 1 (Bad)

extern __shared__ int x[]; /* undefined */
extern __shared__ float y[]; /* undefined */

__global__ void kernel() {
  bool cmp = (void*)x == (void*)y; // Non-compliant: the kernel odr-uses both
// 'x' and 'y'.
}

Examples 2 (Bad)

extern __shared__ int x[]; /* undefined */

__global__ void kernel() {
  float *y;
  y = (float *)&x[1]; // assumes sizeof(float) equals sizeof(int).
  bool cmp = (void*)x == (void*)y; // Compliant: the kernel derives 'y' from 'x'.
}

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_shared_object_dynamic_arrays	Only use one __shared__ array-of-unknown-bound declaration in a kernel 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

use_pointer_analysis

use_pointer_analysis : bool = True

Whether to use pointer analysis to verify indirect calls.

Note: pointer analysis can only be used if StaticSemanticAnalysis is enabled.