CUDA程序触发illegal memory access错误，如何修复？

问题描述 #

已将每个block最大线程数设为1024、每个维度最大grid block数设为62500并硬编码到源码中，运行程序时出现以下错误：

user0@server:~/$ ./test.bin

threads_per_block 32 * 8 * 4 = 1024

blocks_per_grid 62500 * 62500 * 62500 = 1798993472
CUDA Runtime Error at: test.cu:160
an illegal memory access was encountered cudaDeviceSynchronize()
CUDA Runtime Error at: test.cu:161
an illegal memory access was encountered cudaMemcpy(host_c, device_c, sizeof(t) * length, cudaMemcpyDeviceToHost)
File written successfully.

错误原因分析 #

1. 核函数索引逻辑错误：核函数中错误地将一维数组的总长度length赋值给三维维度dimx、dimy、dimz，导致计算出的index = z * dimx * dimy + y * dimx + x远大于数组实际长度（比如当length=1e6时，三维坐标的index会达到1e18级别），直接触发非法内存访问。
2. Grid维度设置完全不合理：硬编码的blocks_per_grid = dim3(62500, 62500, 62500)会创建远超需求的线程数，不仅浪费资源，还会导致大量线程计算出越界的索引值。
3. 线程索引与数组维度不匹配：实际处理的是一维数组，却使用三维线程索引来计算位置，逻辑上完全不匹配。

修复方案 #

1. 修正核函数的索引逻辑 #

将核函数改为一维线程索引模式，直接对应一维数组的位置：

__global__ void kernel_func(float *arr1, float *arr2, float *outp, int length) 
{
    // 计算当前线程的一维全局索引
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    // 确保索引不越界
    if (index < length) 
    {
        outp[index] = arr1[index] + arr2[index];
    }
}

2. 合理计算Grid和Block维度 #

根据数组总长度、每个Block的线程数，计算需要的Block数量，避免硬编码过大的Grid维度：

int max_threads_per_block = 1024;
threads_per_block = dim3(max_threads_per_block);
// 计算需要的Block数，向上取整
int blocks_per_grid = (length + max_threads_per_block - 1) / max_threads_per_block;
// 确保Block数不超过设备支持的最大Grid维度（这里用62500作为上限）
blocks_per_grid = min(blocks_per_grid, 62500);

3. 修正Kernel调用参数 #

使用修正后的Grid和Block维度调用核函数：

kernel_func<<<blocks_per_grid, threads_per_block>>>(device_a, device_b, device_c, length);

修改后的完整代码 #

#include <iostream>
#include <cstdio>
#include <cstdlib>
#include <ctime>
#include <string>
#include <cuda.h>
#include <fstream>
#include <iomanip>
#include <string>
#include <algorithm> // 用于min函数

typedef float t;

const int min_ = 0;
const int max_ = 10;

void init_rand()
{
    srand(time(0));
}

float rand_float(float min_, float max_)
{
    float randomFloat = (max_ - min_) * ((float)rand() / RAND_MAX) + min_;
    return randomFloat;
}

int rand_int(int min_, int max_) 
{
    return rand() % (max_ - min_ + 1) + min_;
}

void write_output_to_file(t* host_a, t* host_b, t* host_c, std::string fileName, int length) 
{
    std::ofstream outputFile;
    outputFile.open(fileName);

    if(outputFile.is_open()) 
    {
        for(int i = 0; i < length; i++) 
        {
            outputFile << std::setw(10) << std::left << host_a[i] << std::setw(10) << std::left << host_b[i] << std::setw(10) << std::left << host_c[i] << "\n";
        }
        outputFile.close();
        std::cout << "File written successfully.\n";
    } 
    else 
    {
        std::cerr << "Error opening file.\n";
    }
}

void print_dim3(std::string text, dim3 data)
{
    std::cout << "\n";
    std::cout << text << " ";
    std::cout << data.x <<" * "<< data.y<<" * "<< data.z << " = ";
    std::cout << data.x * data.y * data.z;
    std::cout << "\n";
}

#define CHECK_CUDA_ERROR(val) check((val), #val, __FILE__, __LINE__)
template <typename T>
void check(T err, const char* const func, const char* const file, const int line)
{
    if (err != cudaSuccess)
    {
        std::cerr << "CUDA Runtime Error at: " << file << ":" << line
                  << std::endl;
        std::cerr << cudaGetErrorString(err) << " " << func << std::endl;
    }
}

#define CHECK_LAST_CUDA_ERROR() checkLast(__FILE__, __LINE__)
void checkLast(const char* const file, const int line)
{
    cudaError_t err{cudaGetLastError()};
    if (err != cudaSuccess)
    {
        std::cerr << "CUDA Runtime Error at: " << file << ":" << line
                  << std::endl;
        std::cerr << cudaGetErrorString(err) << std::endl;
    }
}

__global__ void kernel_func(float *arr1, float *arr2, float *outp, int length) 
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < length) 
    {
        outp[index] = arr1[index] + arr2[index];
    }
}

int main()
{
    t * device_a;
    t * device_b;
    t * device_c;
    t * host_a;
    t * host_b;
    t * host_c;
    int length;
    dim3 threads_per_block;
    int blocks_per_grid;
    
    length = 1000000;
    host_a = (t *) malloc(sizeof(t) * length);
    host_b = (t *) malloc(sizeof(t) * length);
    host_c = (t *) malloc(sizeof(t) * length);
    
    if (host_a == nullptr || host_b == nullptr || host_c == nullptr)
    {
        std::cerr << "Error: Memory allocation for host arrays failed." << std::endl;
        exit(1);
    }

    CHECK_CUDA_ERROR(cudaMalloc((void**) &device_a, sizeof(t) * length));
    CHECK_CUDA_ERROR(cudaMalloc((void**) &device_b, sizeof(t) * length));
    CHECK_CUDA_ERROR(cudaMalloc((void**) &device_c, sizeof(t) * length));

    init_rand(); // 初始化随机数生成器
    for (int i = 0; i < length ; ++i) 
    {
        host_a[i] = rand_int(min_, max_);
        host_b[i] = rand_int(min_, max_);
        host_c[i] = 0;
    }

    int max_threads_per_block = 1024;
    int max_block_per_grid_per_dim = 62500;
    threads_per_block = dim3(max_threads_per_block);
    blocks_per_grid = (length + max_threads_per_block - 1) / max_threads_per_block;
    blocks_per_grid = std::min(blocks_per_grid, max_block_per_grid_per_dim);

    print_dim3("threads_per_block", threads_per_block);
    std::cout << "\nblocks_per_grid = " << blocks_per_grid << "\n";

    CHECK_CUDA_ERROR(cudaMemcpy(device_a, host_a, sizeof(t) * length, cudaMemcpyHostToDevice));
    CHECK_CUDA_ERROR(cudaMemcpy(device_b, host_b, sizeof(t) * length, cudaMemcpyHostToDevice));
    
    kernel_func<<<blocks_per_grid, threads_per_block>>>(device_a, device_b, device_c, length);
    
    CHECK_LAST_CUDA_ERROR();
    
    CHECK_CUDA_ERROR(cudaDeviceSynchronize());  
    CHECK_CUDA_ERROR(cudaMemcpy(host_c, device_c, sizeof(t) * length, cudaMemcpyDeviceToHost));
    
    write_output_to_file(host_a, host_b, host_c, "output.txt", length);
    
    cudaFree(device_a);
    cudaFree(device_b);
    cudaFree(device_c);
    free(host_a);
    free(host_b);
    free(host_c);
}

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