当你推回一个位于堆栈上的值，但不是在你退缩的时候堆栈位置的地址- 从 for 循环的一次迭代到下一次迭代不会改变：
for (auto &input : {a, b}) {
CUdeviceptr ptr; // a stack variable
cuMemAlloc(&ptr, input.size()*sizeof(float));
cuMemcpyHtoD(ptr, input.data(), input.size()*sizeof(float));
kernel_arguments.push_back(reinterpret_cast<void *> (&ptr));
}
这就解释了为什么第一个和第二个参数似乎都在引用您的第二个内核输入参数（i_2、3）。
否则，当我围绕您所展示的内容构建完整代码时，我不会收到任何错误 700。
（一旦变量超出范围，压入/使用堆栈变量的地址也会导致 UB）
这是一个对分配循环进行微不足道的修改（即在每次迭代时覆盖堆栈值）的示例，以解决该问题：
$ cat vectorAddDrv.cpp
/*
* Copyright 1993-2015 NVIDIA Corporation. All rights reserved.
*
* Please refer to the NVIDIA end user license agreement (EULA) associated
* with this source code for terms and conditions that govern your use of
* this software. Any use, reproduction, disclosure, or distribution of
* this software and related documentation outside the terms of the EULA
* is strictly prohibited.
*
*/
/* Vector addition: C = A + B.
*
* This sample is a very basic sample that implements element by element
* vector addition. It is the same as the sample illustrating Chapter 3
* of the programming guide with some additions like error checking.
*
*/
// Includes
#include <stdio.h>
#include <string.h>
#include <iostream>
#include <cstring>
#include <cuda.h>
// includes, project
#include <helper_cuda_drvapi.h>
#include <helper_functions.h>
// includes, CUDA
#include <builtin_types.h>
#define _DEBUG
#include <vector>
using namespace std;
// Variables
CUdevice cuDevice;
CUcontext cuContext;
CUmodule cuModule;
CUfunction vecAdd_kernel;
// Functions
bool findModulePath(const char *, string &, char **, string &);
//define input fatbin file
#ifndef FATBIN_FILE
#define FATBIN_FILE "vectorAdd_kernel64.fatbin"
#endif
// Host code
int main(int argc, char **argv)
{
printf("Vector Addition (Driver API)
");
int N = 50000, devID = 0;
size_t size = N * sizeof(float);
// Initialize
checkCudaErrors(cuInit(0));
cuDevice = findCudaDeviceDRV(argc, (const char **)argv);
// Create context
checkCudaErrors(cuCtxCreate(&cuContext, 0, cuDevice));
// first search for the module path before we load the results
string module_path;
std::ostringstream fatbin;
if (!findFatbinPath(FATBIN_FILE, module_path, argv, fatbin))
{
exit(EXIT_FAILURE);
}
else
{
printf("> initCUDA loading module: <%s>
", module_path.c_str());
}
if (!fatbin.str().size())
{
printf("fatbin file empty. exiting..
");
exit(EXIT_FAILURE);
}
// Create module from binary file (FATBIN)
checkCudaErrors(cuModuleLoadData(&cuModule, fatbin.str().c_str()));
// Get function handle from module
checkCudaErrors(cuModuleGetFunction(&vecAdd_kernel, cuModule, "VecAdd_kernel"));
std::vector<void *> kernel_arguments;
std::vector<float> a = {2};
std::vector<float> b = {3};
for (auto &input : {a, b}) {
CUdeviceptr *ptr = new CUdeviceptr;
cuMemAlloc(ptr, input.size()*sizeof(float));
cuMemcpyHtoD(*ptr, input.data(), input.size()*sizeof(float));
kernel_arguments.push_back(reinterpret_cast<void *> (ptr));
}
std::vector<float> c(1);
for (auto &output : {c}) {
CUdeviceptr *ptr = new CUdeviceptr;
cuMemAlloc(ptr, output.size()*sizeof(float));
kernel_arguments.push_back(reinterpret_cast<void *> (ptr));
}
CUresult result = cuLaunchKernel(vecAdd_kernel, 1, 1, 1,
1024, 1, 1, 0, NULL,
kernel_arguments.data(), NULL);
const char *error;
cuGetErrorString(result, &error);
std::cout << result << " " << error << std::endl;
checkCudaErrors(cuCtxSynchronize());
cuGetErrorString(result, &error);
std::cout << result << " " << error << std::endl;
# if 0
// Allocate input vectors h_A and h_B in host memory
h_A = (float *)malloc(size);
h_B = (float *)malloc(size);
h_C = (float *)malloc(size);
// Initialize input vectors
RandomInit(h_A, N);
RandomInit(h_B, N);
// Allocate vectors in device memory
checkCudaErrors(cuMemAlloc(&d_A, size));
checkCudaErrors(cuMemAlloc(&d_B, size));
checkCudaErrors(cuMemAlloc(&d_C, size));
// Copy vectors from host memory to device memory
checkCudaErrors(cuMemcpyHtoD(d_A, h_A, size));
checkCudaErrors(cuMemcpyHtoD(d_B, h_B, size));
if (1)
{
// This is the new CUDA 4.0 API for Kernel Parameter Passing and Kernel Launch (simpler method)
// Grid/Block configuration
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
void *args[] = { &d_A, &d_B, &d_C, &N };
// Launch the CUDA kernel
checkCudaErrors(cuLaunchKernel(vecAdd_kernel, blocksPerGrid, 1, 1,
threadsPerBlock, 1, 1,
0,
NULL, args, NULL));
}
else
{
// This is the new CUDA 4.0 API for Kernel Parameter Passing and Kernel Launch (advanced method)
int offset = 0;
void *argBuffer[16];
*((CUdeviceptr *)&argBuffer[offset]) = d_A;
offset += sizeof(d_A);
*((CUdeviceptr *)&argBuffer[offset]) = d_B;
offset += sizeof(d_B);
*((CUdeviceptr *)&argBuffer[offset]) = d_C;
offset += sizeof(d_C);
*((int *)&argBuffer[offset]) = N;
offset += sizeof(N);
// Grid/Block configuration
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
// Launch the CUDA kernel
checkCudaErrors(cuLaunchKernel(vecAdd_kernel, blocksPerGrid, 1, 1,
threadsPerBlock, 1, 1,
0,
NULL, NULL, argBuffer));
}
#endif
#ifdef _DEBUG
checkCudaErrors(cuCtxSynchronize());
#endif
#if 0
// Copy result from device memory to host memory
// h_C contains the result in host memory
checkCudaErrors(cuMemcpyDtoH(h_C, d_C, size));
// Verify result
int i;
for (i = 0; i < N; ++i)
{
float sum = h_A[i] + h_B[i];
if (fabs(h_C[i] - sum) > 1e-7f)
{
break;
}
}
#endif
exit(EXIT_SUCCESS);
}
$ nvcc -I/usr/local/cuda/samples/common/inc -o test vectorAddDrv.cpp -lcuda
$ compute-sanitizer ./test
========= COMPUTE-SANITIZER
Vector Addition (Driver API)
> Using CUDA Device [0]: Tesla V100-PCIE-32GB
> findModulePath found file at <./vectorAdd_kernel64.fatbin>
> initCUDA loading module: <./vectorAdd_kernel64.fatbin>
0 no error
index 0
0x7f8023c00000
2.000000
0x7f8023c00200
3.000000
0 no error
========= ERROR SUMMARY: 0 errors
$