当矢量大小不可分为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 %dn", 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列的矩阵，因为使用float 4时，我如何在矩阵中的第二行的前四个元素打印出4个？如果数据像：

2 3 4 5 9

4 2 5 9 1

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

我知道一种可能的方法是在每行末尾填充额外的数字，以使其可将其除以4，但是我正在寻找无需操纵数据的解决方案。

Answer 1

非常简短的答案是，您无法像您想象的那样执行此操作。CUDA对类型施加对齐限制，这意味着"正确"的指针别名：

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

是非法的，因为不满足对齐要求（具有m=5的a+m不能正确对齐与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 %dn", 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 %fn", 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文档中进行了描述。