CUDADirective-1.3ΒΆ

Assign an initial value to device-thread-block storage duration objects prior to their use

Required inputs: IR, StaticSemanticAnalysis

CUDA DIRECTIVE 1.3 [initiation.object.device-thread-block] Assign an initial value to device-thread-block storage duration objects prior to their use

A device-thread-block storage object (e.g. __shared___ variable) should be initialized before its value is read.

Scope: Device.
Audience: CUDA C++.
Hardware Applicability: All Compute Capabilities.
Rationale

Variables with device-thread-block storage duration (e.g. __shared__ variables) are uninitialized. This is a consequence of the lifetime of device-thread-block storage duration variables, which ends at the end of device-thread-block. Use of an uninitialized value is undefined behavior; thus, the a device-thread-block storage duration object must be initialized before its value is read.

Example 1 (Bad)
# include <cuda.h>

__global__ void bad_copy(int *d, int n)
{
  __shared__ int s[64];
  int t = threadIdx.x;
  int tr = n-t-1;
  d[t] = s[tr]; // s is ubd
}
Example 2 (Bad)
# include <cuda.h>

__global__ void bad_copy(int *d, int n)
{
  extern __shared__ int s[];
  int t = threadIdx.x;
  int tr = n-t-1;
  d[t] = s[tr]; // non-conformant: s has not been defined.
}
Example 3 (Good)
# include <cuda.h>

__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[64];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}
Example 4 (Good)
# include <cuda.h>

__global__ void bad_copy(int *d, int n)
{
  extern __shared__ int s[];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}
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

possible_uninitialized_read

Possible read access of uninitialized __shared__ variable

None

False

possible_unsynced_read

Possible read access of unsynchronized __shared__ variable

None

False

Options