CUDA:处理不同大小的数组

CUDA: working with arrays of different sizes

本文关键字:数组 处理 CUDA      更新时间:2023-10-16

在这个例子中,我试图使用来自10x9数组的值创建一个10x8数组。它看起来像我访问内存不正确,但我不确定我的错误在哪里。

c++中的代码类似于

for (int h = 0; h < height; h++){
    for (int i = 0; i < (width-2); i++)
        dd[h*(width-2)+i] = hi[h*(width-1)+i] + hi[h*(width-1)+i+1];
}

这是我在CUDA中尝试的:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdint.h>
#include <iostream>
#define TILE_WIDTH 4
using namespace std;
__global__ void cudaOffsetArray(int height, int width, float *HI, float *DD){
    int             x                   =   blockIdx.x * blockDim.x + threadIdx.x; // Col // width
    int             y                   =   blockIdx.y * blockDim.y + threadIdx.y; // Row // height
    int             grid_width          =   gridDim.x  * blockDim.x;
  //int             index               =   y * grid_width + x;
    if ((x < (width - 2)) && (y < (height)))
        DD[y * (grid_width - 2) + x] = (HI[y * (grid_width - 1) + x] + HI[y * (grid_width - 1) + x + 1]);
}
int main(){
    int height  = 10;
    int width   = 10;
    float *HI = new float [height * (width - 1)];
    for (int i = 0; i < height; i++){
        for (int j = 0; j < (width - 1); j++)
            HI[i * (width - 1) + j] = 1;
    }
    float   *gpu_HI;
    float   *gpu_DD;
    cudaMalloc((void **)&gpu_HI, (height * (width - 1) * sizeof(float)));
    cudaMalloc((void **)&gpu_DD, (height * (width - 2) * sizeof(float)));
    cudaMemcpy(gpu_HI, HI, (height * (width - 1) * sizeof(float)), cudaMemcpyHostToDevice);
    dim3            dimGrid((width - 1) / TILE_WIDTH + 1, (height - 1)/TILE_WIDTH + 1, 1);
    dim3            dimBlock(TILE_WIDTH, TILE_WIDTH, 1);
    cudaOffsetArray<<<dimGrid,dimBlock>>>(height, width, gpu_HI, gpu_DD);
    float *result   = new float[height * (width - 2)];
    cudaMemcpy(result, gpu_DD,  (height * (width - 2) * sizeof(float)), cudaMemcpyDeviceToHost);
    for (int i = 0; i < height; i++){
        for (int j = 0; j < (width - 2); j++)
            cout << result[i * (width - 2) + j] << " ";
        cout << endl;
    }
    cudaFree(gpu_HI);
    cudaFree(gpu_DD);
    delete[] result;
    delete[] HI;
    system("pause");
}

我也在全局函数中尝试过:

if ((x < (width - 2)) && (y < (height)))
    DD[y * (grid_width - 2) + (blockIdx.x - 2) * blockDim.x + threadIdx.x] = 
        (HI[y * (grid_width - 1) + (blockIdx.x - 1) * blockDim.x + threadIdx.x] + 
         HI[y * (grid_width - 1) + (blockIdx.x - 1) * blockDim.x + threadIdx.x + 1]);

要"修复"你的代码,将内核中这一行grid_width的用法改为width:

    DD[y * (grid_width - 2) + x] = (HI[y * (grid_width - 1) + x] + HI[y * (grid_width - 1) + x + 1]);

:

    DD[y * (width - 2) + x] = (HI[y * (width - 1) + x] + HI[y * (width - 1) + x + 1]);

解释:

你的grid_width:

dim3            dimGrid((width * 2 - 1) / TILE_WIDTH + 1, (height - 1)/TILE_WIDTH + 1, 1);
dim3            dimBlock(TILE_WIDTH, TILE_WIDTH, 1);

实际上并不对应于您的数组大小(10x10、10x9或10x8)。我不确定为什么要在x维中启动2*width线程,但这意味着线程数组比数据数组大得多。

所以当你在内核中使用grid_width时:

    DD[y * (grid_width - 2) + x] = (HI[y * (grid_width - 1) + x] + HI[y * (grid_width - 1) + x + 1]);

索引将是一个问题。如果您将上面的grid_width的每个实例更改为width(对应于数据数组的实际宽度),我认为您将获得更好的索引。通常,启动"额外线程"不是问题,因为在内核中有线程检查行:

if ((x < (width - 2)) && (y < (height)))

但是当你启动额外的线程时,它使你的网格更大,所以你不能使用网格尺寸来正确地索引到你的数据数组