parallel-processing - 当向量大小不能被 4 整除时，使用 cuda 向量类型 float 4

Question

我试图了解 cuda 向量类型的工作原理。假设我有一个 n 行 m 列的矩阵，并且 m 不能被 4 整除。矩阵被线性化并存储在 GPU 主内存中。是否可以使用 float4 数据类型并读取第二个向量的第一个元素？我编写了一个非常简单的内核来看看它是如何工作的，但根据我使用的方式，我无法访问第二个向量的第一个元素。这是代码:

#include<iostream>
#include <ctime>
#include<stdio.h>
#include<stdlib.h>
#include<math.h>
using namespace std;
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

__global__ void ker(float * a,int n, int m)
{
        float4 f;
        f=reinterpret_cast<float4*>(a)[1];
        printf("%f %f %f %f,",f.x,f.y,f.z,f.w);

}
int main()
{
        int n=2,m=5;
        float *a=new float[n*m];
        for(int i=0;i<n;i++)
        {
                for(int j=0;j<m;j++)
                        {
                                a[i*m+j]=rand()%10;
                                cout<<a[i*m+j]<<" ";

                        }
                cout<<"\n";
        }
        float * dev_a;
        cudaMalloc ((void**)&dev_a,sizeof(float)*m*n);
        gpuErrchk(cudaMemcpy(dev_a, a, sizeof(float) * m* n, cudaMemcpyHostToDevice));
        ker<<<1,1>>>(dev_a,n,m);
        gpuErrchk( cudaPeekAtLastError());
        cudaFree(dev_a);
        delete []a;
        return 0;
}

在代码中，我有一个 2 行 5 列的矩阵，由于 5 不能被 4 整除，如何在使用 float 4 时在内核中打印矩阵第二行的前四个元素？如果数据是这样的:

2 3 4 5 9

4 2 5 9 1

f=reinterpret_cast<float4*>(a)[1];读取数据 block 9 4 2 5和f=reinterpret_cast<float4*>(a)[2];读 9 1 0 0这不是我想要的( 4 2 5 9 )。有什么方法可以在使用 float4 时读取第二行的前四个元素吗？

我知道一种可能的方法是在每行末尾填充额外的数字(例如 0)，使其可被 4 整除，但我正在寻找一种无需操作数据的解决方案。

Answer 1

非常简短的答案是，您无法按照您的想象做到这一点。 CUDA 对类型施加了对齐限制，这意味着“正确的”指针别名:

f = *reinterpret_cast<float4*>(a+m);

是非法的，因为不满足对齐要求(a+m 和 m=5 未正确对齐到 float4边界)。在较旧的工具链/硬件上，这会产生运行时错误。在较新的硬件/工具链上，它将编译成运行时不会出错的东西，但读取会自动重新对齐，结果不是您所期望的。

但是，您可以使用 cudaMallocPitch 和 cudaMemcpy2D 在设备上分配倾斜线性内存，并复制您拥有的数据，以便设备副本正确对齐，你正在尝试做的事情将会奏效。如果您将代码更改为:

#include <iostream>
#include <ctime>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
using namespace std;
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
    if (code != cudaSuccess) {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

template<typename T, typename T0>
struct pitchedpointer
{
    char *p;
    size_t pitch;

    __host__ __device__ pitchedpointer() {};
    __host__ __device__
    pitchedpointer(T0* _p, size_t _pitch) : p(reinterpret_cast<char*>(_p)), pitch(_pitch) {};
    __device__ __host__
    T& operator()(size_t i, size_t j) {
        T* v = reinterpret_cast<T*>(p + i*pitch);        
        return v[j]; 
    }
    __device__ __host__
    const T& operator()(size_t i, size_t j) const {
        T* v = reinterpret_cast<T*>(p + i*pitch);        
        return v[j]; 
    }
};

__global__ void ker(float * a, int m, int n, size_t pitch)
{
    int row = threadIdx.x;
    pitchedpointer<float4,float> p(a, pitch);
    float4 f = p(row,1);
    printf("%d: %f %f %f %f\n", row, f.x, f.y, f.z, f.w);
}
int main()
{
    int n=4,m=9;
    float *a=new float[n*m];
    for(int i=0;i<n;i++)
    {
        for(int j=0;j<m;j++)
        {
            a[i*m+j]=rand()%10;
            cout << a[i*m+j] << " ";
        }
        cout << endl;
    }
    float * dev_a;
    size_t pitch;
    int m4 = 1 + (m-1)/4;
    gpuErrchk( cudaMallocPitch((void**)&dev_a, &pitch, sizeof(float4)*m4, n) );
    gpuErrchk( cudaMemcpy2D(dev_a, pitch, a, sizeof(float)*m, sizeof(float)*m, n, cudaMemcpyHostToDevice) );
    ker<<<1,n>>>(dev_a, m, n, pitch);
    gpuErrchk( cudaPeekAtLastError() );
    gpuErrchk( cudaDeviceSynchronize() );
    cudaFree(dev_a);
    delete []a;
    cudaDeviceReset();
    return 0;
}

它的作用是:

~/SO$ nvcc -arch=sm_52 -std=c++11 float4align.cu 
~/SO$ ./a.out 
3 6 7 5 3 5 6 2 9 
1 2 7 0 9 3 6 0 6 
2 6 1 8 7 9 2 0 2 
3 7 5 9 2 2 8 9 7 
0: 3.000000 5.000000 6.000000 2.000000
1: 9.000000 3.000000 6.000000 0.000000
2: 7.000000 9.000000 2.000000 0.000000
3: 2.000000 2.000000 8.000000 9.000000

如您所见，它以 float4 形式正确访问矩阵的各个行，而不会违反对齐要求(我选择打印每行的第二个 float4 ，这同样是未对准的)。我引入的类只是一些糖，用于简化/隐藏在设备上使用倾斜内存所需的指针算术，这在 cudaMallocPitch documentation 中进行了描述。 。