1

Nx3配列をカーネルに渡し、テクスチャメモリのようにカーネルから読み取り、2番目の配列に書き込もうとしています。これがN=8の私の簡略化されたコードです：

#include <cstdio>
#include "handle.h"
using namespace std;

texture<float,2> tex_w;

__global__ void kernel(int imax, float(*w)[3], float (*f)[3])
{
  int i = threadIdx.x;
  int j = threadIdx.y;

  if(i<imax)
      f[i][j] = tex2D(tex_w, i, j);
}

void print_to_stdio(int imax, float (*w)[3])
{
  for (int i=0; i<imax; i++)
    {
      printf("%2d  %3.6f\t  %3.6f\t %3.6f\n",i, w[i][0], w[i][1], w[i][2]);
    }
}

int main(void)
{
  int imax = 8;
  float (*w)[3];
  float (*d_w)[3], (*d_f)[3];
  dim3 grid(imax,3);

  w = (float (*)[3])malloc(imax*3*sizeof(float));

  for(int i=0; i<imax; i++)
    {
      for(int j=0; j<3; j++)
        {
          w[i][j] = i + 0.01f*j;
        }
    }

  cudaMalloc( (void**) &d_w, 3*imax*sizeof(float) );
  cudaMalloc( (void**) &d_f, 3*imax*sizeof(float) );

  cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
  HANDLE_ERROR( cudaBindTexture2D(NULL, tex_w, d_w, desc, imax, 3, sizeof(float)*imax ) );

  cudaMemcpy(d_w, w, 3*imax*sizeof(float), cudaMemcpyHostToDevice);

  // just use threads for simplicity                                                                  
  kernel<<<1,grid>>>(imax, d_w, d_f);

  cudaMemcpy(w, d_f, 3*imax*sizeof(float), cudaMemcpyDeviceToHost);

  cudaUnbindTexture(tex_w);
  cudaFree(d_w);
  cudaFree(d_f);

  print_to_stdio(imax, w);

  free(w);
  return 0;
}

このコードを実行すると、次のようになります。 

0  0.000000   0.010000   0.020000
1  1.000000   1.010000   1.020000
2  2.000000   2.010000   2.020000
3  3.000000   3.010000   3.020000
4  4.000000   4.010000   4.020000
5  5.000000   5.010000   5.020000
6  6.000000   6.010000   6.020000
7  7.000000   7.010000   7.020000

しかし、代わりに私は得る： 

0  0.000000   2.020000   5.010000
1  0.010000   3.000000   5.020000
2  0.020000   3.010000   6.000000
3  1.000000   3.020000   6.010000
4  1.010000   4.000000   6.020000
5  1.020000   4.010000   7.000000
6  2.000000   4.020000   7.010000
7  2.010000   5.000000   7.020000

これは、cudaBindTexture2Dに指定したピッチパラメータと関係があると思いますが、小さい値を使用すると、無効な引数エラーが発生します。

前もって感謝します！

4

2 に答える 2

4

brano の回答とピッチの仕組みを詳しく調べた後、私は自分の質問に答えます。変更されたコードは次のとおりです。

#include <cstdio>
#include <iostream>
#include "handle.cu"

using namespace std;

texture<float,2,cudaReadModeElementType> tex_w;

__global__ void kernel(int imax, float (*f)[3])
{
  int i = threadIdx.x;
  int j = threadIdx.y;
  // width = 3, height = imax                                                                         
  // but we have imax threads in x, 3 in y                                                            
  // therefore height corresponds to x threads (i)                                                    
  // and width corresponds to y threads (j)                                                           
  if(i<imax)
    {
      // linear filtering looks between indices                                                       
      f[i][j] = tex2D(tex_w, j+0.5f, i+0.5f);
    }
}

void print_to_stdio(int imax, float (*w)[3])
{
  for (int i=0; i<imax; i++)
    {
      printf("%2d  %3.3f  %3.3f  %3.3f\n",i, w[i][0], w[i][1], w[i][2]);
    }
  printf("\n");
}

int main(void)
{
  int imax = 8;
  float (*w)[3];
  float (*d_f)[3], *d_w;
  dim3 grid(imax,3);

  w = (float (*)[3])malloc(imax*3*sizeof(float));

  for(int i=0; i<imax; i++)
    {
      for(int j=0; j<3; j++)
        {
          w[i][j] = i + 0.01f*j;
        }
    }

  print_to_stdio(imax, w);

  size_t pitch;
  HANDLE_ERROR( cudaMallocPitch((void**)&d_w, &pitch, 3*sizeof(float), imax) );

  HANDLE_ERROR( cudaMemcpy2D(d_w,             // device destination                                   
                             pitch,           // device pitch (calculated above)                      
                             w,               // src on host                                          
                             3*sizeof(float), // pitch on src (no padding so just width of row)       
                             3*sizeof(float), // width of data in bytes                               
                             imax,            // height of data                                       
                             cudaMemcpyHostToDevice) );

  HANDLE_ERROR( cudaBindTexture2D(NULL, tex_w, d_w, tex_w.channelDesc, 3, imax, pitch) );

  tex_w.normalized = false;  // don't use normalized values                                           
  tex_w.filterMode = cudaFilterModeLinear;
  tex_w.addressMode[0] = cudaAddressModeClamp; // don't wrap around indices                           
  tex_w.addressMode[1] = cudaAddressModeClamp;

  // d_f will have result array                                                                       
  cudaMalloc( &d_f, 3*imax*sizeof(float) );

  // just use threads for simplicity                                                                  
  kernel<<<1,grid>>>(imax, d_f);

  cudaMemcpy(w, d_f, 3*imax*sizeof(float), cudaMemcpyDeviceToHost);

  cudaUnbindTexture(tex_w);
  cudaFree(d_w);
  cudaFree(d_f);

  print_to_stdio(imax, w);

  free(w);
  return 0;
}

memcpy() を使用してホスト マシンでピッチを処理する代わりに、memcpy2D() を使用すると、デバイス データとホスト データの両方のピッチ引数が受け入れられます。ホストで単純に割り当てられたデータを使用しているため、ピッチは単純に行幅、または 3*sizeof(float) になると理解しています。

于 2012-08-13T19:10:21.460 に答える
2

私はあなたに完全な解決策を与えることができますが、あなたは学ぶことができないかもしれません :D 、代わりにいくつかのヒントがあり、残りは自分で修正できるかもしれません.

ヒント 1.
使用cudaBindTexture2Dする場合は、オフセットとピッチを要求します。両方のパラメーターには、ハードウェアに依存する特定のアラインメント制限があります。を使用すると、オフセットは 0 であることが保証されますcudaMalloc(..)。を使用してピッチを取得しcudaMallocPitch(..)ます。また、ホスト メモリが同じようにピッチングされていることを確認する必要があります。そうしmemcpyないと、期待どおりに動作しません。

ヒント 2.
2D でのインデックス作成を理解する。要素にアクセスするときは、要素がメモリ内の次の要素であり、 NOT であるW[i][j]ことを知る必要があります。W[i][j+1]W[i+1][j]

ヒント 3.
1D 配列を使用して、2D インデックスを自分で計算します。これにより、より適切に制御できます。

于 2012-08-13T10:34:53.130 に答える