CUDA-1.12ΒΆ

Constant memory objects shall not be modified from the device program

Required inputs: IR, StaticSemanticAnalysis

CUDA 1.12 [memory.constant_writes] Constant memory objects shall not be modified from the device program

An object in constant memory shall not be modified from the device program.

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

An CUDA constant memory has special semantics allowing both the hardware and the software to make assumptions about how constant memory may or may not change. Changing constant memory from the device code may violate assumptions and result in unexpected behavior.

Example 1 (Bad)
__device__ void foo(int* rx) {
    *rx = 123;
}

__global__ void test() {
    __constant__ static int x;
    foo((int*)&x);
    //int& rx = x;
    //rx = 123; // Non-compliant: assignment to a __constant__ memory location.
}
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

constant_modification

An object in constant memory shall not be modified from the device program

None

False

Options