1

我有一个用 C 语言编写的低级 for 循环，一位朋友建议我用 CUDA 编写。我已经设置了我的 CUDA 环境并一直在查看文档，但我仍然在为语法问题苦苦挣扎，这已经超过 2 周了。谁能帮我吗？这在 CUDA 中会是什么样子？

float* red = new float [N];
float* green = new float [N];
float* blue = new float [N];

for (int y = 0; y < h; y++)
{
    // Get row ptr from the color image
    const unsigned char* src = rowptr<unsigned char>(color, 0, y, w);

    // Get row ptrs for the destination channel features
    float* rptr = rowptr<float>(red, 0, y, w);
    float* gptr = rowptr<float>(green, 0, y, w);
    float* bptr = rowptr<float>(blue, 0, y, w);

    for (int x = 0; x < w; x++)
    {
        *rptr++ = (float)*src++;
        *gptr++ = (float)*src++;
        *bptr++ = (float)*src++;
    }
}
4

2 回答 2

1

这是一些示例代码。我不知道它是否真的能回答你的问题。可能您需要了解更多关于 CUDA 的信息。如果您能抽出时间，从nvidia 网络研讨会页面参加本次网络研讨会和本次网络研讨会将花费 2 个小时。cuda C 程序员手册也是一个很好的可读参考。

#include <stdio.h>

#define N      256
#define NUMROW   N
#define NUMCOL   N
#define PIXSIZE  3
#define REDOFF   0
#define GREENOFF 1
#define BLUEOFF  2
#define nTPB    16
#define GRNVAL   5
#define REDVAL   7
#define BLUVAL   9

#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"); \
            exit(1); \
        } \
    } while (0)

__global__ void kern(const unsigned numrow, const unsigned numcol, const unsigned char* src, float* rptr, float* gptr, float* bptr){

  unsigned idx = threadIdx.x + (blockDim.x*blockIdx.x);
  unsigned idy = threadIdx.y + (blockDim.y*blockIdx.y);
  if ((idx < numcol) && (idy < numrow)){

    rptr[(idy*numcol)+idx] = (float)src[(((idy*numcol)+idx)*PIXSIZE)+REDOFF];
    gptr[(idy*numcol)+idx] = (float)src[(((idy*numcol)+idx)*PIXSIZE)+GREENOFF];
    bptr[(idy*numcol)+idx] = (float)src[(((idy*numcol)+idx)*PIXSIZE)+BLUEOFF];
    }
}

int main (){

  float *h_red, *h_green, *h_blue;
  float *d_red, *d_green, *d_blue;
  unsigned char *h_img, *d_img;

  if ((h_img =(unsigned char*)malloc(NUMROW*NUMCOL*PIXSIZE*sizeof(unsigned char))) == 0) {printf("malloc fail\n"); return 1;}
  if ((h_red =(float*)malloc(NUMROW*NUMCOL*sizeof(float))) == 0) {printf("malloc fail\n"); return 1;}
  if ((h_green =(float*)malloc(NUMROW*NUMCOL*sizeof(float))) == 0) {printf("malloc fail\n"); return 1;}
  if ((h_blue =(float*)malloc(NUMROW*NUMCOL*sizeof(float))) == 0) {printf("malloc fail\n"); return 1;}

  cudaMalloc((void **)&d_img, (NUMROW*NUMCOL*PIXSIZE)*sizeof(unsigned char));
  cudaCheckErrors("cudaMalloc1 fail");
  cudaMalloc((void **)&d_red, (NUMROW*NUMCOL)*sizeof(float));
  cudaCheckErrors("cudaMalloc2 fail");
  cudaMalloc((void **)&d_green, (NUMROW*NUMCOL)*sizeof(float));
  cudaCheckErrors("cudaMalloc3 fail");
  cudaMalloc((void **)&d_blue, (NUMROW*NUMCOL)*sizeof(float));
  cudaCheckErrors("cudaMalloc4 fail");

  for (int i=0; i<NUMROW*NUMCOL; i++){
    h_img[(i*PIXSIZE)+ REDOFF]   = REDVAL;
    h_img[(i*PIXSIZE)+ GREENOFF] = GRNVAL;
    h_img[(i*PIXSIZE)+ BLUEOFF]  = BLUVAL;
    }

  cudaMemcpy(d_img, h_img, (NUMROW*NUMCOL*PIXSIZE)*sizeof(unsigned char), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy1 fail");

  dim3 block(nTPB, nTPB);
  dim3 grid(((NUMCOL+nTPB-1)/nTPB),((NUMROW+nTPB-1)/nTPB));
  kern<<<grid,block>>>(NUMROW, NUMCOL, d_img, d_red, d_green, d_blue);
  cudaMemcpy(h_red, d_red, (NUMROW*NUMCOL)*sizeof(float), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy2 fail");
  cudaMemcpy(h_green, d_green, (NUMROW*NUMCOL)*sizeof(float), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy3 fail");
  cudaMemcpy(h_blue, d_blue, (NUMROW*NUMCOL)*sizeof(float), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy4 fail");

  for (int i=0; i<(NUMROW*NUMCOL); i++){
    if (h_red[i] != REDVAL) {printf("Red mismatch at offset %d\n", i); return 1;}
    if (h_green[i] != GRNVAL) {printf("Green mismatch at offset %d\n", i); return 1;}
    if (h_blue[i] != BLUVAL) {printf("Blue mismatch at offset %d\n", i); return 1;}
    }
  printf("Success!\n");
  return 0;
}

为了回答评论中提出的问题，这里有一个修改过的内核，它展示了如何使用评论中定义的 rowptr<> 模板。只需将上面的内核代码替换为：

template <typename T> T* rowptr(T* start, int x, int y, int w) __device__ __host__ { return start + y*w + x; }

__global__ void kern(const unsigned numrow, const unsigned numcol, unsigned char* isrc, float* rptr, float* gptr, float* bptr){


  unsigned idx = threadIdx.x + (blockDim.x*blockIdx.x);
  unsigned idy = threadIdx.y + (blockDim.y*blockIdx.y);
  if ((idx < numcol) && (idy < numrow)){
    unsigned char *src = rowptr<unsigned char>(isrc, (idx*PIXSIZE), idy, (numcol*PIXSIZE));

    rptr[(idy*numcol)+idx] = (float)*src++;
    gptr[(idy*numcol)+idx] = (float)*src++;
    bptr[(idy*numcol)+idx] = (float)*src;
    }
}
于 2012-10-24T19:26:29.150 回答
0

由于此代码是一个简单的 for 循环，而不是编写 CUDA，另一种选择是检查 openACC。这样，您只需向现有代码添加指令。

于 2012-10-24T05:08:47.090 回答