0

我有以下内容：

struct LR { double eps_dielect; 
        double sgm_conductivity; 
        double eno_ns_surfref;
        double frq_mhz; 
        double conf; 
        double rel;
        double erp;
        int radio_climate;  
        int pol;
        float antenna_pattern[361][1001];
          } LR;

我需要将 LR.antenna_pattern 传递给一个函数，在 CUDA 设备中分配内存然后复制它。float** 类型应该代表 LR.antenna_pattern[361][1001] 就好了，但我不知道如何实例化 float** 变量，以便它是指向 LR.antenna_pattern 的指针

我尝试 float** antennaPattern = (void**)&LR.antenna_pattern 但它不起作用。如何创建指向 LR.antenna_pattern 的指针？

4

1 回答 1

1

一种方法是展平您的 2D 数组并以 1D 方式处理它，并使用指针算法来处理行和列维度。首先在您的结构定义中，将天线模式元素替换为：

struct LR { 
.
.
float *antenna_pattern;
} LR;

然后你需要做一个主机端 malloc 来分配空间：

#define COL 1001
#define ROW 361
#define DSIZE (ROW*COL)

LR.antenna_pattern = (float *)malloc(DSIZE*sizeof(float));

还有一个设备端 cuda malloc：

float *d_antenna_pattern;
cudaMalloc((void **) &d_antenna_pattern, DSIZE*sizeof(float));

到设备的副本如下所示：

cudaMemcpy(d_antenna_pattern, LR.antenna_pattern, DSIZE*sizeof(float), cudaMemcpyHostToDevice);

当您想引用这些数组时，您必须执行指针运算，例如：

float my_val_xy = ap[(x*COL)+y];  // to access element at [x][y] on the device

float my_val_xy = LR.antenna_pattern[(x*COL)+y]; // on the host

如果你想在整个过程中保持二维数组下标，你可以使用适当的 typedef 来做到这一点。例如，请参阅我对此问题的回答中的第一个代码示例。为了把它画出来，你需要从 typedef 开始：

#define COL 1001
#define ROW 361
#define DSIZE (ROW*COL)

typedef float aParray[COL];

并修改您的结构定义：

struct LR { 
.
.
aParray *antenna_pattern;
} LR;

主机端 malloc 看起来像：

LR.antenna_pattern = (aParray *)malloc(DSIZE*sizeof(float));

设备端 cuda malloc 看起来像：

aParray *d_antenna_pattern;
cudaMalloc((void **) &d_antenna_pattern, DSIZE*sizeof(float));

到设备的副本如下所示：

cudaMemcpy(d_antenna_pattern, LR.antenna_pattern, DSIZE*sizeof(float), cudaMemcpyHostToDevice);

设备内核定义将需要一个函数参数，例如：

__global__ void myKernel(float ap[][COL]) {

然后在内核内部，您可以访问 x,y 处的元素，如下所示：

float my_val_xy = ap[x][y];

现在回答一个后续问题，询问如果 LR 无法更改该怎么办，这里有一个完整的示例代码，它结合了其中一些想法而不修改 LR 结构：

#include<stdio.h>

// for cuda error checking
#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            return 1; \
        } \
    } while (0)



struct LR {
  int foo;
  float antenna_pattern[361][1001];
  } LR;

__global__ void mykernel(float ap[][1001]){

  int tid = threadIdx.x + (blockDim.x*blockIdx.x);
  float myval = 0.0;

  if (tid == 0){
    for (int i=0; i<361; i++)
      for (int j=0; j<1001; j++)
         ap[i][j] = myval++;
  }
}


int main(){

  typedef float aParray[1001];

  aParray *d_antenna_pattern;
  cudaMalloc((void **) &d_antenna_pattern, (361*1001)*sizeof(float));
  cudaCheckErrors("cudaMalloc fail");
  float *my_ap_ptr;
  my_ap_ptr = &(LR.antenna_pattern[0][0]);

  for (int i=0; i< 361; i++)
    for (int j=0; j<1001; j++)
      LR.antenna_pattern[i][j] = 0.0;
  cudaMemcpy(d_antenna_pattern, my_ap_ptr, (361*1001)*sizeof(float), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy fail");
  mykernel<<<1,1>>>(d_antenna_pattern);
  cudaCheckErrors("Kernel fail");

  cudaMemcpy(my_ap_ptr, d_antenna_pattern, (361*1001)*sizeof(float), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy 2 fail");
  float myval = 0.0;
  for (int i=0; i<361; i++)
    for (int j=0; j<1001; j++)
      if (LR.antenna_pattern[i][j] != myval++) {printf("mismatch at offset x: %d y: %d actual: %f expected: %f\n", i, j, LR.antenna_pattern[i][j], --myval); return 1;}
  printf("Results match!\n");
  return 0;
}

如果您更喜欢使用扁平化方法，请将d_antenna_pattern定义替换为：

float *d_antenna_pattern;

并将核函数参数对应更改为：

__global__ void mykernel(float *ap){

然后使用内核中的指针算术方法访问：

ap[(i*1001)+j] = myval++;
于 2012-11-27T19:39:00.383 回答