admin管理员组

文章数量:1391964

In CUDA it's both easy to implicitly downcast/truncate integers and surprisingly common for programmers to do so.

I would like CUDA to raise an error when implicit downcasting occurs.

Consider this scaling kernel:

__global__ void foo(int *const a, const int N){
    const auto tid = threadIdx.x + blockDim.x * blockIdx.x;
    // Implicit downcast/truncation from int64 to int32 - how can we get an error?
    a[tid] = tid * 3L;
}

How can I get an error for a[tid] = tid * 3L?

In GCC -Wconverion would flag a similar situation:

<source>:27:9: warning: conversion from ‘int64_t’ {aka ‘long int’} to ‘int’ may change value [-Wconversion]
   27 |     int y = x;
      |         ^

and in clang -Wshorten64-to-32 would raise:

<source>:6:11: warning: implicit conversion loses integer precision: 'int64_t' (aka 'long') to 'int' [-Wshorten-64-to-32]
    6 |     int y=x;
      |         ~ ^

Minimum working example

#include <cstdint>
#include <iostream>
#include <limits>
#include <vector>

__global__ void foo(int *const a, const int N){
    // Wrong data type: it would be nice to detect this
    const int tid = threadIdx.x + blockDim.x * blockIdx.x;
    // Implicit downcast/truncation from int64 to int32 - how can we get an error?
    a[tid] = tid * 3L;
}

int main(){
    constexpr int N = 10000;
    constexpr int threads = 128;
    int *a_d;
    cudaMalloc(&a_d, N * sizeof(int));
    foo<<<(N + threads - 1) / 128, threads>>>(a_d, N);

    std::vector<int> v(N);
    cudaMemcpy(v.data(), a_d, N * sizeof(int), cudaMemcpyDeviceToHost);
    for(int i=0;i<10;i++){
        std::cout<<v.at(i)<<std::endl;
    }

    return 0;
}

In CUDA it's both easy to implicitly downcast/truncate integers and surprisingly common for programmers to do so.

I would like CUDA to raise an error when implicit downcasting occurs.

Consider this scaling kernel:

__global__ void foo(int *const a, const int N){
    const auto tid = threadIdx.x + blockDim.x * blockIdx.x;
    // Implicit downcast/truncation from int64 to int32 - how can we get an error?
    a[tid] = tid * 3L;
}

How can I get an error for a[tid] = tid * 3L?

In GCC -Wconverion would flag a similar situation:

<source>:27:9: warning: conversion from ‘int64_t’ {aka ‘long int’} to ‘int’ may change value [-Wconversion]
   27 |     int y = x;
      |         ^

and in clang -Wshorten64-to-32 would raise:

<source>:6:11: warning: implicit conversion loses integer precision: 'int64_t' (aka 'long') to 'int' [-Wshorten-64-to-32]
    6 |     int y=x;
      |         ~ ^

Minimum working example

#include <cstdint>
#include <iostream>
#include <limits>
#include <vector>

__global__ void foo(int *const a, const int N){
    // Wrong data type: it would be nice to detect this
    const int tid = threadIdx.x + blockDim.x * blockIdx.x;
    // Implicit downcast/truncation from int64 to int32 - how can we get an error?
    a[tid] = tid * 3L;
}

int main(){
    constexpr int N = 10000;
    constexpr int threads = 128;
    int *a_d;
    cudaMalloc(&a_d, N * sizeof(int));
    foo<<<(N + threads - 1) / 128, threads>>>(a_d, N);

    std::vector<int> v(N);
    cudaMemcpy(v.data(), a_d, N * sizeof(int), cudaMemcpyDeviceToHost);
    for(int i=0;i<10;i++){
        std::cout<<v.at(i)<<std::endl;
    }

    return 0;
}

• cuda
• nvcc

Share Improve this question Follow edited Mar 14 at 4:32 Johan 76.9k2727 gold badges200200 silver badges340340 bronze badges asked Mar 14 at 1:00 RichardRichard 61.9k3939 gold badges196196 silver badges275275 bronze badges 3

• 1 @EternalDreamer: auto is almost always the better choice and the purpose of 3L is to demonstrate the implicit downcast I'm trying to solve: the number stands in for more complex operations that don't need to be in a MWE. – Richard Commented Mar 14 at 8:04
• You could try to make your CUDA code compatible with clang and use that for static analysis (or even clang-tidy). – paleonix Commented Mar 14 at 10:19
• Excellent question, dunno why ppl downvote it. I would really like to have more control over the warning level in CUDA code, because the current state is dismal, but alas. – Johan Commented Mar 15 at 17:08

Add a comment  | 

1 Answer 1

Sorted by: Reset to default 4

One possible approach: make the body of your kernel code a __device__ __host__ function, call that from your kernel, and use -Xcompiler -Wconversion :

# cat t360.cu
__host__ __device__ void f(int *d, long long *l){

  *d = *l;
}

__global__ void k(int *d, long long *l){
  f(d, l);
}

int main(){
  int *d = NULL;
  long long *l = NULL;
  k<<<1,1>>>(d, l);
  cudaDeviceSynchronize();
}

# nvcc -o t360 t360.cu
# nvcc -o t360 t360.cu -Xcompiler -Wconversion
t360.cu: In function ‘void f(int*, long long int*)’:
t360.cu:3:9: warning: conversion from ‘long long int’ to ‘int’ may change value [-Wconversion]
    3 |   *d = *l;
      |        ~^~
#

Things will be more complicated if you have code in your kernel that cannot be run in host code, such as CUDA intrinsics like warp shuffle. It would still be possible to use this method, but you would have to chop up your kernel around these points, probably getting messy in some cases.

With a bit of internet wayback research effort, you can also use the method described here. The specific warning we are looking for is impl_narrowing_64_bit_int, however its also going to flag things in CUDA header files that nvcc automatically includes:

# nvcc -o t360 t360.cu -Xcudafe="--diag_warn=impl_narrowing_64_bit_int"
/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(895): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<long int>(min(static_cast<int>(a), static_cast<int>(b)));
                                         ^

Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(895): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<long int>(min(static_cast<int>(a), static_cast<int>(b)));
                                                              ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(912): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(912): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(929): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(929): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(946): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(946): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1019): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<long int>(max(static_cast<int>(a), static_cast<int>(b)));
                                         ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1019): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<long int>(max(static_cast<int>(a), static_cast<int>(b)));
                                                              ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1036): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1036): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1053): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1053): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1070): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1070): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

t360.cu(3): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
    *d = *l;
       ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(895): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<long int>(min(static_cast<int>(a), static_cast<int>(b)));
                                         ^

Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(895): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<long int>(min(static_cast<int>(a), static_cast<int>(b)));
                                                              ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(912): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(912): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(929): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(929): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(946): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(946): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umin(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1019): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<long int>(max(static_cast<int>(a), static_cast<int>(b)));
                                         ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1019): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<long int>(max(static_cast<int>(a), static_cast<int>(b)));
                                                              ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1036): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1036): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1053): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1053): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1070): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                   ^

/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/math_functions.hpp(1070): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
      retval = static_cast<unsigned long int>(umax(static_cast<unsigned int>(a), static_cast<unsigned int>(b)));
                                                                                 ^

t360.cu(3): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
    *d = *l;
       ^

#

so that may be an issue.

Finally, you can probably get what you want with

#pragma nv_diag_warning 1373

(prior to CUDA 12.0 use diag_warning instead of nv_diag_warning)

like so:

# cat t360.cu
#pragma nv_diag_warning 1373

__host__ __device__ void f(int *d, long long *l){

  *d = *l;
}

__global__ void k(int *d, long long *l){
  f(d, l);
}

int main(){
  int *d = NULL;
  long long *l = NULL;
  k<<<1,1>>>(d, l);
  cudaDeviceSynchronize();
}

# nvcc -o t360 t360.cu
t360.cu(5): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
    *d = *l;
       ^

Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

t360.cu(5): warning #1373-D: implicit conversion of a 64-bit integral type to a smaller integral type (potential portability problem)
    *d = *l;
       ^

Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

#

Sorry, I completely missed that the question was asking to convert a warning to an error. I believe a possible method to do that is to use for example:

#pragma nv_diag_error 1373

instead of:

#pragma nv_diag_warning 1373

Also, to generalize things a bit, the 1373 and the impl_narrowing_64_bit_int are related/equivalent forms of the diag ID. One can be used in place of the other. The general process I used was to inspect the list of errors from the internet wayback results above, then test the diag IDs (text form) that looked likely, approximately in the sequence I have shown in this answer. Once you find one that does what you want, the actual warning output will indicate the ID (1373 in this case). That can be used subsequently.

Additionally, once a diagnostic ID of interest is identified, nvcc has controls to handle it in various ways:

for the diagnostic ID indicated by XXX:

#pragma nv_diag_suppress XXX - do not issue a warning or error

#pragma nv_diag_warning XXX - issue a warning

#pragma nv_diag_error XXX - issue (or promote to) an error

#pragma nv_diag_default XXX - restore the default behavior (as to issue a warning, error, or silence) for subsequent source code

#pragma nv_diag_once XXX - issue the current behavior for the ID only once in subsequent source code

As indicated, the above pragmas are expected to influence compiler messaging behavior for subsequent code, after the pragma.

The CLI way to set an error is:

nvcc test.cu  -Xcudafe="--diag_error=impl_narrowing_64_bit_int"

本文标签： cudaHow can I get NVCC to error upon implicit integer downcastingtruncationStack Overflow