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] = 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));
cudaMemcpy(d_x, x, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, N * sizeof(float), cudaMemcpyHostToDevice);
cudaEventRecord(start);
saxpy << <(N + 511) / 512, 512 >> >(N, 2.0f, d_x, d_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);
CUDA中内存访问越界问题在核函数中对非法的内存访问不一定会报错，那我们加CHECK（error）会帮我们报错吗？经实验发现，有的非法访问很致命，函数不会运行，这种错误，可以CHECK到，但有些非法内存访问不致命，可能是访问到其他设备内存上了，这时CHECK不会报错，但是会出现逻辑上的错误，因为访问的值是错误的。 所以，我们在做内存访问时，一定要自己做好逻辑上的判断，保证内存访问不错位，不错...
报错 RuntimeError: CUDA error: an illegal memory access was encountered
CUDA kernel errors might be asynchronously reported at some other API call,
so the stacktrace below might be ...
先说一下在网上看到的问题：
第一种可能你的程序涉及到并行计算，但你只有一张卡，因此只要将程序涉及到并行计算的部分改成单卡即可
找找有没有 torch.nn.DataParallel()
第二种一部分数据或者模型在cpu上，另外一部分在gpu上。
PS：第二种和第一种可以强行 os.environ["CUDA_VISIBLE_DEVICES"]为指定ID
第三种CUDNN版本不对？感觉这种可能比较小
然后说一下我是怎么解决的：
最关键的，也是我遇到的问题，这个错误没有表明和显存溢出存在着联系，因为显存溢出
Recently, I encountered “an illegal memory access was encountered” error during CUDA programming. such as:
cudaSafeCall(cudaMemcpy(h_res, d_res, gParams.n*sizeof(uint32), cudaMemcpyDeviceToHost));
Because
用pytorch在多卡训练transformers的时候出现了以下问题：
RuntimeError: CUDA error: an illegal memory access was encountered
terminate called after throwing an instance of 'c10::Error'
what(): CUDA error: an illegal memory access was encountered
Exception raised from create
文章目录系列文章目录前言一、内存访问模式之全局内存读取1. 内存访问模式基础知识1.1 对齐和合并访问二、全局内存读取示例1.代码实现2.NVIDIA Visual Profiler运行结果分析总结参考资料
之前在复习现代C++的新特性，没有继续CUDA C编程的学习，今天开始继续之前的学习，这里跟大家分享内存访问模式中全局内存读取的知识。
一、内存访问模式之全局内存读取
1. 内存访问模式基础知识
我们所编写的GPU程序容易受到内存带宽的限制，因此，最大限度利用全局内存
CUDA 是 Nvidia 提供的 GPU 编程框架，它提供了一组 API，使得程序员可以通过 C/C++ 程序来编写并行程序，以利用 GPU 的海量计算能力。在 CUDA 中，每一个 GPU 都有多级存储器，包括共享存储器、常量存储器、纹理存储器、全局存储器和本地存储器。
共享存储器是 GPU 中最快的存储器类型，它在 GPU 内部的多个线程之间共享，用于线程之间进行数据交换。常量存储器用于存储只读常量数据，纹理存储器用于存储纹理数据，全局存储器用于存储 GPU 上所有线程可以访问的数据，本地存储器用于存储每个线程块中线程私有的数据。
在 CUDA 程序设计中，应该尽量使用共享存储器和常量存储器，因为它们的访问速度比全局存储器和纹理存储器要快。另外，在编写并行程序时应该注意避免不必要的数据传输，并尽量减少对全局存储器和纹理存储器的访问。
