CUDA矩阵向量乘代码中atomicAdd函数参数匹配错误求助

问题描述 #

我用CUDA实现带共享内存的矩阵向量乘法并做性能分析，在Visual Studio 2022运行代码时，出现错误："no instance of overloaded function "atomicAdd" matches the argument list"。完整代码如下：

#ifndef __CUDACC__  
#define __CUDACC__
#endif
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <cuda.h>
#include "cuda_runtime.h"
#include <iostream>

#define TILE_SIZE 16
#define BLOCK_SIZE 256

__global__ void matvec(double* A, double* B, double* C, int n)
{
    __shared__ double s_A[TILE_SIZE][TILE_SIZE];
    __shared__ double s_B[TILE_SIZE];

    int bx = blockIdx.x;
    int tx = threadIdx.x;
    int i = bx * blockDim.x + tx;

    if (i < n) {
        s_B[tx] = B[i];
        for (int j = 0; j < n; j += TILE_SIZE) {
            s_A[tx][j + threadIdx.y] = A[(i * n) + j + threadIdx.y];
        }
    }
    __syncthreads();

    if (i < n) {
        double tmp = 0.0;
        for (int j = 0; j < n; j += TILE_SIZE) {
            tmp += s_A[threadIdx.x][j + threadIdx.y] * s_B[j + threadIdx.y];
        }
        atomicAdd(&C[i], tmp);
    }
}

int main()
{
    int n = 5000;
    double* A, * B, * C;
    double* d_A, * d_B, * d_C;

    A = (double*)malloc(n * n * sizeof(double));
    B = (double*)malloc(n * sizeof(double));
    C = (double*)malloc(n * sizeof(double));

    for (int i = 0; i < n * n; i++) {
        A[i] = rand() / (double)RAND_MAX;
    }
    for (int i = 0; i < n; i++) {
        B[i] = rand() / (double)RAND_MAX;
        C[i] = 0.0;
    }

    cudaMalloc((void**)&d_A, n * n * sizeof(double));
    cudaMalloc((void**)&d_B, n * sizeof(double));
    cudaMalloc((void**)&d_C, n * sizeof(double));

    cudaMemcpy(d_A, A, n * n * sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, n * sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_C, C, n * sizeof(double), cudaMemcpyHostToDevice);

    dim3 dimBlock(TILE_SIZE, BLOCK_SIZE / TILE_SIZE, 1);
    dim3 dimGrid((n + BLOCK_SIZE - 1) / BLOCK_SIZE, 1, 1);

    matvec <<<dimGrid, dimBlock>>> (d_A, d_B, d_C, n);

    cudaMemcpy(C, d_C, n * sizeof(double), cudaMemcpyDeviceToHost);

    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    free(A);
    free(B);
    free(C);

    return 0;
}

错误含义 #

这个错误的核心原因是：CUDA原生的atomicAdd函数默认不支持double类型的原子操作。只有GPU计算能力达到6.0及以上（Pascal架构及之后），且开启对应编译选项时，才能使用double类型的原子加法。

解决方法 #

方法1：启用double类型原子操作的编译选项 #

在Visual Studio中配置项目：

• 右键项目 → 属性 → 配置属性 → CUDA C/C++ → Device
• 在Code Generation选项中，添加支持double原子操作的计算能力，比如compute_60,sm_60（Pascal）、compute_70,sm_70（Volta）等，根据你的GPU型号选择对应参数。
• 保存配置后重新编译。

方法2：替换为float类型（精度允许时） #

如果场景可以接受单精度浮点的精度损失，直接将所有double类型替换为float：

• 代码中所有double改为float
• rand()/(double)RAND_MAX改为rand()/(float)RAND_MAX
• 原生atomicAdd支持float类型，无需额外编译选项。

方法3：手动实现double类型原子加法（适配低算力GPU） #

如果你的GPU计算能力低于6.0，可以基于CAS（Compare-And-Swap）手动实现double原子加法：
添加以下设备函数：

__device__ void atomicAddDouble(double* address, double val)
{
    unsigned long long int* address_as_ull = (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
    } while (assumed != old);
}

然后将原代码中的atomicAdd(&C[i], tmp);替换为atomicAddDouble(&C[i], tmp);

额外代码问题修正 #

原代码还有逻辑错误，会导致运行崩溃或结果错误：

• threadIdx.y越界访问：block配置为dim3 dimBlock(TILE_SIZE, BLOCK_SIZE / TILE_SIZE, 1)，threadIdx.y范围是0到BLOCK_SIZE/TILE_SIZE-1，但代码中直接用threadIdx.y访问s_A[tx][j + threadIdx.y]，会导致共享内存越界。
• 共享内存加载逻辑错误：矩阵向量乘法的共享内存使用方式有误，且原代码中atomicAdd是多余的——每个结果元素应由单个线程负责计算，不需要原子操作。

修正后的核心kernel示例（配合方法1的编译选项）：

__global__ void matvec(double* A, double* B, double* C, int n)
{
    __shared__ double s_A[TILE_SIZE][TILE_SIZE];
    __shared__ double s_B[TILE_SIZE];

    int row = blockIdx.x * blockDim.x + threadIdx.x;
    double sum = 0.0;

    for (int col_block = 0; col_block < n; col_block += TILE_SIZE) {
        // 加载向量块到共享内存
        if (col_block + threadIdx.x < n) {
            s_B[threadIdx.x] = B[col_block + threadIdx.x];
        } else {
            s_B[threadIdx.x] = 0.0;
        }
        // 加载矩阵的行块到共享内存
        if (row < n && col_block + threadIdx.y < n) {
            s_A[threadIdx.y][threadIdx.x] = A[row * n + col_block + threadIdx.y];
        } else {
            s_A[threadIdx.y][threadIdx.x] = 0.0;
        }
        __syncthreads();

        // 计算当前tile内的乘积和
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += s_A[threadIdx.x][k] * s_B[k];
        }
        __syncthreads();
    }

    if (row < n) {
        C[row] = sum; // 每个row由单个线程负责，无需atomicAdd
    }
}

内容的提问来源于stack exchange，提问作者Anthony J. B.