cudaMalloc 和 cudaMemcpy 不能处理内核调用答案

【问题标题】：cudaMalloc and cudaMemcpy not working on kernel callcudaMalloc 和 cudaMemcpy 不能处理内核调用
【发布时间】：2016-06-03 19:03:22
【问题描述】：

我已经初始化了一个数组，我试图在内核调用的每个线程中使用它（每个线程使用数组的不同部分，因此没有依赖关系）。我使用cudaMalloc 创建数组并在设备上保存内存，然后使用cudaMemcpy 将数组从主机复制到设备。

我将cudaMalloc 返回的指针传递给内核调用以供每个线程使用。

int SIZE = 100;

int* data = new int[SIZE];
int* d_data = 0;

cutilSafeCall( cudaMalloc(&d_data, SIZE * sizeof(int)) );
for (int i = 0; i < SIZE; i++)
    data[i] = i;

cutilSafeCall( cudaMemcpy(d_data, data, SIZE * sizeof(int), cudaMemcpyHostToDevice) );

此代码取自 here。用于内核调用。

kernel<<<blocks, threads>>> (results, d_data);

我使用结构Result 跟踪每个线程的结果。下一个代码可以正常工作。

__global__ void mainKernel(Result res[], int* data){
   int x = data[0];
}

但是当我将该值分配给res:

__global__ void mainKernel(Result res[], int* data){
   int threadId = (blockIdx.x * blockDim.x) + threadIdx.x;

   int x = data[0];

   res[threadId].x = x;
}

出现错误：

文件中的 cudaSafeCall() 运行时 API 错误，第 355 行：遇到非法内存访问。

任何涉及使用该指针的操作都会出现相同的错误

__global__ void mainKernel(Result res[], int* data){
   int threadId = (blockIdx.x * blockDim.x) + threadIdx.x;

   int x = data[0];

   if (x > 10)
      res[threadId].x = 5;
}

res的定义没有问题。将任何其他值分配给 res[threadId].x 不会给我任何错误。

这是运行 cuda-memcheck 的输出：

========= 无效的 __global__ 读取大小为 4
========= 在 mainKernel 中的 0x00000150（结果*，int*）
========= 块 (49,0,0) 中的线程 (86,0,0)
========= 地址 0x13024c0000 超出范围
========= 在内核启动时将主机回溯保存到驱动程序入口点
========= 主机框架：/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x150d6d]
========= 主机帧：./out [0x2cc4b]
========= 主机帧：./out [0x46c23]
========= 主机帧：./out [0x3e37]
========= 主机帧：./out [0x3ca1]
========= 主机帧：./out [0x3cd6]
========= 主机帧：./out [0x39e9]
========= 主机框架：/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xf5) [0x21ec5]
========= 主机帧：./out [0x31b9]

编辑：

这是完整代码的示例：

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <string.h>
#include <iostream>
#include <assert.h>

typedef struct    
{
   int x,y,z;
} Result;

__global__ void mainKernel(Result  pResults[], int* dataimage)
{

   int threadId = (blockIdx.x * blockDim.x) + threadIdx.x;

   int xVal = dataimage[0];
   if (xVal > 10)
       pResults[threadId].x = 5;

}

int main (int argc, char** argv)
{

   int NUM_THREADS = 5*5;

   int SIZE = 100;

   int* data = new int[SIZE];
   int* d_data = 0;

   cutilSafeCall( cudaMalloc(&d_data, SIZE * sizeof(int)) );
   for (int i = 0; i < SIZE; i++)
       data[i] = i;

   cutilSafeCall( cudaMemcpy(d_data, data, SIZE * sizeof(int), cudaMemcpyHostToDevice) );

   unsigned int GPU_ID = 1;  // not actually :-)
   // unsigned int GPU_ID =  cutGetMaxGflopsDeviceId() ;
   cudaSetDevice(GPU_ID); 

   Result * results_GPU = 0;
   cutilSafeCall( cudaMalloc( &results_GPU,  NUM_THREADS * sizeof(Result)) );

   Result * results_CPU = 0;
   cutilSafeCall( cudaMallocHost( &results_CPU, NUM_THREADS * sizeof(Result)) );

   mainKernel<<<5,5>>> ( results_GPU, d_data );

   cudaThreadSynchronize(); 

   cutilSafeCall( cudaMemcpy(results_CPU, results_GPU, NUM_THREADS * sizeof(Result),cudaMemcpyDeviceToHost) );

   cutilSafeCall(cudaFree(results_GPU));
   cutilSafeCall(cudaFreeHost(results_CPU));
   cudaThreadExit();

 } // ()

【问题讨论】：

问题显然出在results，但您已经设法完全省略了所有显示您如何定义和分配它的代码。您能否编辑您的问题以包含其他人可以编译和运行的简短、完整的代码？没有它，你的问题很难得到任何答案
要清楚，确切的代码会产生您发布的错误？你有多少个 CUDA GPU？
是的（您必须包括 CudaSafeCall）。我有两个 GPU。我删除了用于选择 GPU (cudaSetDevice) 的行，它不会引发任何错误。发生了什么事？

标签： memory cuda

【解决方案1】：

你的问题在于这个调用序列：

  cutilSafeCall( cudaMalloc(&d_data, SIZE * sizeof(int)) );
   for (int i = 0; i < SIZE; i++)
       data[i] = i;

   cutilSafeCall( cudaMemcpy(d_data, data, SIZE * sizeof(int), cudaMemcpyHostToDevice) );

   unsigned int GPU_ID = 1;
   cudaSetDevice(GPU_ID); 

   Result * results_GPU = 0;
   cutilSafeCall( cudaMalloc( &results_GPU,  NUM_THREADS * sizeof(Result)) );

   Result * results_CPU = 0;
   cutilSafeCall( cudaMallocHost( &results_CPU, NUM_THREADS * sizeof(Result)) );

   mainKernel<<<5,5>>> ( results_GPU, d_data );

实际发生的情况是您正在分配 d_data 并在不同的 GPU 上运行内核，而 d_data 在您启动内核的 GPU 上无效。

详细地说，因为您在cudaSetDevice 之前为d_data 调用cudaMalloc，所以您在默认设备上分配d_data，然后显式分配results_GPU 并运行内核在设备 1 上。显然设备 1 和默认设备不是同一个 GPU（设备的枚举通常在运行时 API 中从 0 开始）。

如果你像这样更改代码：

   unsigned int GPU_ID = 1;
   cutilSafeCall(cudaSetDevice(GPU_ID)); 

   cutilSafeCall( cudaMalloc(&d_data, SIZE * sizeof(int)) );
   for (int i = 0; i < SIZE; i++)
       data[i] = i;

   cutilSafeCall( cudaMemcpy(d_data, data, SIZE * sizeof(int), cudaMemcpyHostToDevice) );

   Result * results_GPU = 0;
   cutilSafeCall( cudaMalloc( &results_GPU,  NUM_THREADS * sizeof(Result)) );

   Result * results_CPU = 0;
   cutilSafeCall( cudaMallocHost( &results_CPU, NUM_THREADS * sizeof(Result)) );

   mainKernel<<<5,5>>> ( results_GPU, d_data );

即在进行任何分配之前选择非默认设备，问题应该会消失。您的非常简单的内核不会发生这种情况的原因：

__global__ void mainKernel(Result res[], int* data){
   int x = data[0];
}

很简单，CUDA 编译器默认执行非常激进的优化，并且因为data[0] 的读取结果实际上并没有被使用，所以整个读取可以被优化掉，你只剩下一个空的存根内核不做任何事情。只有当从内存加载的结果用于内存写入时，代码才会在编译期间被优化掉。如果你好奇，你可以通过反汇编编译器发出的代码来确认这一点。

请注意，有一些方法可以通过点对点访问在支持它的多 GPU 系统上工作，但必须在代码中明确配置才能使用该工具。

【讨论】：