ベクトルのサイズが4で割り切れない場合のcuda vector type float 4の使用

Question

私はcudaベクトル型の仕組みを理解しようとしています。 n行とm列の行列を持ち、mは4で割り切れないとします。行列は線形化され、GPUメインメモリに格納されます。 float4データ型を使用して、2番目のベクトルの最初の要素を読み取ることは可能ですか？私はそれがどのように動作するかを見るために非常に単純なカーネルを書きましたが、私が使った方法に基づいて、第2のベクトルの最初の要素にアクセスすることはできません。ここでは、コードである：コードで

#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; 
} 

5私は行列の第2行の最初の4つの要素を印刷する方法、4で割り切れないので、私は、2行5列の行列を有しますフロート4を使用している間カーネルで？データは次のようである場合：

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を使用している間、第2行の最初の4つの要素を読み取ることができる方法はありますか？

4で割り切れるように各行の末尾に0のような余分な数字を埋め込む方法がありますが、データを操作しないで解決策を探しています。

Answer 1

非常に短い答えは、あなたが想像するようにこれを行うことができないということです。アライメント要件が満たされていないため

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

、違法である（m=5とa+mがfloat4境界に正しく整列していません）：CUDAは、「正しい」ポインタ別名ことを意味するタイプのアラインメントの制限を課します。古いツールチェーン/ハードウェアでは、実行時エラーが発生します。より新しいハードウェア/ツールチェーンでは、エラーなく実行されるものにコンパイルされますが、読み取りは自動的に再調整され、結果は期待したものとは異なります。

ただし、cudaMallocPitchとcudaMemcpy2Dを使用して、デバイスにピッチ付き線形メモリを割り当て、デバイスのコピーが正しく配置されるようにデータをコピーすることができます。 （

~/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として行列の個々の行にアクセス：これはない

#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; 
} 

：あなたのような何かにあなたのコードを変更した場合私は2番目のfloat4をそれぞれの行から印刷することを選択した。これは同様に整列していない。私が紹介したクラスは、デバイス上でピッチメモリを使用するために必要なポインタ算術を単純化/隠す単なる砂糖です（cudaMallocPitchdocumentationに記載されています）。