CUDA-1.8ΒΆ

Kernel parameters passed by value should be trivial

Required inputs: IR

CUDA 1.8 [kernel-launch.parameter.value.trivial] Kernel parameters passed by value should be trivial

Any parameter passed to a kernel launch by value should have a trivial type.

Scope: Host, Device.
Audience: CUDA C++.
Category: Required.
Hardware Applicability: All Compute Capabilities.
Rationale

Kernel parameters passed by value are transferred to the device by the equivalent of memcpy. Thus, if a kernel parameter has a non-trivial copy constructor or default constructor, it will not be invoked, and the contents of the class will instead be effectively memcpy'd; this is undefined behavior.

Example 1 (Bad)
# include <memory>
# include <algorithm>
# include <cstring>
# include <cassert>
# include "testTerminate.h"
# include "deviceDynamicArray.h"

__global__ void assert_non_zero(device_dynamic_array<int32_t> u) {
  if (u[threadIdx.x] == 0) {
    *(int*)0 = 0;
  }
}

int main() {
  constexpr int32_t n = 128;
  device_dynamic_array<int32_t> u(n);
  std::fill(u.begin(), u.end(), 1);
  // When `u` is passed by value, a temporary copy of `u` is created in this
  // host-thread and the representation of the temporary copy is `memcpy`ed into
  // the kernel parameters. The temporary copy is destroyed at the end of the
  // evaluation of the kernel launch; because kernel launches are asynchronous,
  // the kernel has not necessarily been run, so the destruction of the temporary
  // copy of `u` will race with the execution of the kernel.
  assert_non_zero<<<1, n>>>(u);
  // This call succeeds, because the kernel launch did not have a synchronous
  // error.
  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

cuda_kernel_launch_parameter_value_trivial

Kernel parameters passed by value should be trivial

None

False

Options