cuda中的内存碎片

Question

我遇到了无法理解的内存分配问题。我正在尝试在 GPU 中分配一个 char 数组（我猜这可能是内存碎片问题）。

这是我的代码，

#include<stdio.h>
#include<stdlib.h>
#include<string.h>
#include<cuda.h>

inline void gpuAssert(cudaError_t code, char *file, int line, 
                 int abort=1)
{  
   if (code != cudaSuccess) {
      printf("GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

__global__ void calc(char *k,char *i)
{
       *i=*k;
}
int main()
{
        char *dev_o=0;
        char *i;
        i = (char*)malloc(10*sizeof(char));

        cudaMalloc((void**)&dev_o,10*sizeof(char)); //Line 31


        calc<<<1,1>>>("arun",dev_o);

        gpuErrchk(cudaMemcpy(&i,dev_o,10*sizeof(char),cudaMemcpyDeviceToHost));

        cudaFree(dev_o);

        printf("string : %s \n",i);

        return 0;
}

但我得到的输出是，

GPUassert: 内存不足 sample2.cu 31

在同样的情况下，我尝试在 GPU 中分配整数并且它工作正常。

我的 GPU 设备信息显示为，

--- General Information for device 0 ---
Name:GeForce GTX 460 SE
Compute capability:2.1
Clock rate:1296000
Device copy overlap:Enabled
Kernel execition timeout :Enabled
--- Memory Information for device 0 ---
Total global mem:1073283072
Total constant Mem:65536
Max mem pitch:2147483647
Texture Alignment:512
--- MP Information for device 0 ---
Multiprocessor count:6
Shared mem per mp:49152
Registers per mp:32768
Threads in warp:32
Max threads per block:1024
Max thread dimensions:(1024, 1024, 64)
Max grid dimensions:(65535, 65535, 65535)

谁能告诉我问题出在哪里以及如何克服它？

Answer 1

您的代码中有几处错误。

1. cudaMemcpy(&i, ...) 应该是 cudaMemcpy(i, ...)。

2. 按照this post 中的说明检查内核调用的返回错误。如果您不这样做，该错误似乎稍后会出现在您的代码中。

   gpuErrchk(cudaPeekAtLastError());
   gpuErrchk(cudaDeviceSynchronize());
   

3. 你的内核的char *k 参数是一个主机指针。在调用内核之前，您应该创建另一个设备数组并将数据复制到设备。
4. 您也没有在calc() 内核中的线程上执行任何并行工作，因为您没有使用线程索引threadIdx.x。不过，这可能是为了测试。

如果您解决这些问题，您会得到以下结果：

#include<stdio.h>
#include<stdlib.h>
#include<string.h>
#include<cuda.h>

inline void gpuAssert(cudaError_t code, char *file, int line, 
                 int abort=1)
{  
   if (code != cudaSuccess) {
      printf("GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

__global__ void calc(char* k, char *i)
{
       i[threadIdx.x] = k[threadIdx.x];
}

int main()
{
        const char* msg = "arun";

        char *dev_i, *dev_k;
        char *i, *k;

        k = (char*)malloc(10*sizeof(char));
        i = (char*)malloc(10*sizeof(char));

        sprintf(k, msg);

        cudaMalloc((void**)&dev_i, 10*sizeof(char));
        cudaMalloc((void**)&dev_k, 10*sizeof(char));

        gpuErrchk(cudaMemcpy(dev_k, k, 10*sizeof(char), cudaMemcpyHostToDevice));

        calc<<<1,5>>>(dev_k, dev_i);

        gpuErrchk(cudaPeekAtLastError());
        // Synchronization will be done in the next synchronous cudaMemCpy call, else
        // you would need cudaDeviceSynchronize() to detect execution errors.
        //gpuErrchk(cudaDeviceSynchronize());

        gpuErrchk(cudaMemcpy(i, dev_i, 10*sizeof(char), cudaMemcpyDeviceToHost));

        printf("string : %s\n", i);

        cudaFree(dev_i);
        cudaFree(dev_k);
        free(i);
        free(k);

        return 0;
}