CUDA中内存访问越界问题

最新推荐文章于 2025-05-24 21:00:28 发布

fb_help

最新推荐文章于 2025-05-24 21:00:28 发布

阅读量5.2k

点赞数

CC 4.0 BY-SA版权

分类专栏： CUDA 文章标签：笔记

本文链接：https://blog.youkuaiyun.com/fb_help/article/details/79599991

CUDA 专栏收录该内容

51 篇文章

订阅专栏

本文探讨了CUDA中核函数的非法内存访问问题，并通过实验发现并非所有非法访问都会导致错误。文章强调了确保内存访问正确性的必要性，并提供了一个具体的代码示例来说明这一现象。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

CUDA中内存访问越界问题

在核函数中对非法的内存访问不一定会报错，那我们加CHECK（error）会帮我们报错吗？经实验发现，有的非法访问很致命，函数不会运行，这种错误，可以CHECK到，但有些非法内存访问不致命，可能是访问到其他设备内存上了，这时CHECK不会报错，但是会出现逻辑上的错误，因为访问的值是错误的。
所以，我们在做内存访问时，一定要自己做好逻辑上的判断，保证内存访问不错位，不错误！

当x先申请设备内存的时候，y+1报错，x+1不报错，但属于逻辑错位。
当y先申请设备内存时，y+1逻辑错位，x+1报错。

例子：

#include <iostream>
#include <stdio.h>
#define SIZE 5
#define CHECK(res) { if(res != cudaSuccess){printf("Error ：%s:%d , ", __FILE__,__LINE__);   \
printf("code : %d , reason : %s \n", res,cudaGetErrorString(res));exit(-1);}}


/*
*   在内存交换的情况下，CPU 0.09 GPU 0.25
*/
#include <stdio.h>
#include <memory>
#include <math.h>

__global__
void saxpy(int n, float a, float *x, float *y)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;

    if(i<4)
    {
        printf("x: %f  y :%f\n",x[i+2],y[i+1]);     
    }

    y[i] = x[i] + y[i+1];
    if(y[i]!=3){
        printf("asd:x: %f  y :%f\n",x[i],y[i]);     
    }
}
void add1(int n, float *x, float *y)
{
    int i = 0;
    for (i; i < n; i++)
    {
        //y[i] = sqrt(x[i] * y[i] * 5 * 5 * 95);
        y[i] = x[i] + y[i];
    }

}
int main(void)
{
    int N =1<<20;
    float *x, *y, *d_x, *d_y;
    x = (float*)malloc(N * sizeof(float));
    y = (float*)malloc(N * sizeof(float));

    cudaMalloc(&d_y, N * sizeof(float));
    cudaMalloc(&d_x, N * sizeof(float));    




    for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaMemcpy(d_x, x, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, y, N * sizeof(float), cudaMemcpyHostToDevice);

    cudaEventRecord(start);

    // Perform SAXPY on 1M elements
    saxpy << <(N + 511) / 512, 512 >> >(N, 2.0f, d_x, d_y);
    //add1(N,x,y);

    CHECK(cudaDeviceSynchronize());

    cudaEventRecord(stop);

    cudaMemcpy(y, d_y, N * sizeof(float), cudaMemcpyDeviceToHost);

    cudaEventSynchronize(stop);
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);

    float maxError = 0.0f;
    for (int i = 0; i < N; i++) {
        maxError = max(maxError, abs(y[i] - 3.0f));
    }

    printf("Max error: %f \n", maxError);
    printf("Effective Bandwidth (GB/s): %f \n",  milliseconds);

    CHECK(cudaGetLastError());

    cudaFree(d_x);
    cudaFree(d_y);
    free(x);
    free(y);
}