openACC：内存管理答案

【问题标题】：openACC: memory managementopenACC：内存管理
【发布时间】：2013-04-27 07:05:48
【问题描述】：

我正在使用 caps openacc 编译器。我想知道我可以自己管理内存吗？

例如，带有 CUDA 的常规 openacc 代码是：

 #pragma acc kernels copyin(a,b) copy(c)
  for (i = 0; i < SIZE; ++i)
    for (j = 0; j < SIZE; ++j)
      for (k = 0; k < SIZE; ++k)
        c[i][j] += a[i][k] * b[k][j];

我想这样改变

//allocation
cudaMalloc((void**)&a, num_bytes);
cudaMalloc((void**)&b, num_bytes);
cudaMalloc((void**)&c, num_bytes);

//transfer-in
cudaMemcpy(hostA, a, num_bytes, cudaMemcpyHostToDevice);
cudaMemcpy(hostB, b, num_bytes, cudaMemcpyHostToDevice);

//computation
//i think it will be generated as codelet by CAPS openACC compiler.
#pragma acc kernels
  for (i = 0; i < SIZE; ++i)
    for (j = 0; j < SIZE; ++j)
      for (k = 0; k < SIZE; ++k)
        c[i][j] += a[i][k] * b[k][j];

cudaMemcpy(c, hostC, num_bytes, cudaMemcpyDeviceToHost);
cudaFree(&a);cudaFree(&b);cudaFree(&c);

【问题讨论】：

标签： cuda openacc

【解决方案1】：

是的，您可以自己分配内存。在您的示例中，应该可以使用device_ptr pragma 来实现这一点，例如：

cudaMalloc((void**)&a, num_bytes);
cudaMalloc((void**)&b, num_bytes);
cudaMalloc((void**)&c, num_bytes);

cudaMemcpy(hostA, a, num_bytes, cudaMemcpyHostToDevice);
cudaMemcpy(hostB, b, num_bytes, cudaMemcpyHostToDevice);

#pragma acc data deviceptr(a, b, c)
#pragma acc kernels
  for (i = 0; i < SIZE; ++i)
    for (j = 0; j < SIZE; ++j)
      for (k = 0; k < SIZE; ++k)
        c[i][j] += a[i][k] * b[k][j];

cudaMemcpy(c, hostC, num_bytes, cudaMemcpyDeviceToHost);
cudaFree(a);cudaFree(b);cudaFree(c);

[免责声明：在浏览器中编写，从未编译或测试，使用风险自负]

这应该声明a、b 和c 是编译器预先存在的分配。如果您愿意，您还应该能够使用 OpenACC acc_malloc 例程来代替 cudaMalloc 分配内存。

感谢@user2054656 指出我在此答案的第一个版本中对device_resident 的错误使用。

【讨论】：

【解决方案2】：

我同意 talonmies 的帖子，但您应该使用 data pragma 的 deviceptr 子句。我这样说是因为device_resident 确实请求 OpenACC 实现分配内存，而 device_ptr 没有。您不需要分配内存，因为它已经由用户分配cudaMalloc()

【讨论】：

感谢您的指正，当我写下我的答案并得到device_resident 的意图时，我是在火车上凭记忆发帖的...