如何让NVCC对隐式整数向下转换/截断抛出错误？

在CUDA开发中，隐式整数向下转换/截断的情况很容易出现，而且这类失误的发生率远超预期。我希望在发生此类隐式转换时，CUDA能直接抛出错误。比如下面这个缩放内核：

__global__ void foo(int *const a, const int N){
    const auto tid = threadIdx.x + blockDim.x * blockIdx.x;
    // 这里存在从int64到int32的隐式向下转换/截断——如何让它触发错误？
    a[tid] = tid * 3L;
}

如何让a[tid] = tid * 3L这一行触发错误？

在GCC中，-Wconversion选项会标记类似的隐式转换问题：

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

在Clang中，-Wshorten64-to-32选项会触发相关提示：

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

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

__global__ void foo(int *const a, const int N){
    // 错误的数据类型：希望能检测到这一点
    const int tid = threadIdx.x + blockDim.x * blockIdx.x;
    // 从int64到int32的隐式向下转换/截断——如何让它触发错误？
    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编译器基于GCC/Clang，可以通过传递宿主编译器的警告选项并将其升级为错误来实现需求：

针对基于GCC的NVCC #

添加编译选项：-Xcompiler -Wconversion,-Werror=conversion

• -Xcompiler：将后续参数传递给宿主编译器（GCC）
• -Wconversion：开启隐式转换的警告检测
• -Werror=conversion：将该类警告直接升级为编译错误

针对基于Clang的NVCC #

添加编译选项：-Xcompiler -Wshorten-64-to-32,-Werror=shorten-64-to-32

• -Wshorten-64-to-32：专门检测64位整数截断为32位的隐式转换
• -Werror=shorten-64-to-32：将该类警告升级为编译错误

效果验证 #

使用上述选项编译最小示例时，编译器会针对a[tid] = tid * 3L行抛出错误，提示从long（int64）转换到int（int32）会丢失精度，强制你显式处理转换（比如用static_cast<int>(tid * 3L)）或修正数据类型避免隐式转换。

内容的提问来源于stack exchange，提问作者Richard