来自2D数组CUDA的2D纹理

Question

我正在尝试将一个Nx3数组传递给内核，并像在纹理内存中一样从它读取，然后写入第二个数组。下面是我用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提供的音调参数有关，但使用较小的值会产生无效参数错误。

提前感谢！

Answer 1

在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(浮点数)。

Answer 2

我可以给你一个完整的解决方案，但你可能学不到:D，所以这里有一些建议，也许你可以自己解决剩下的问题。

提示1。

当使用cudaBindTexture2D时，它需要一个偏移量和间距。这两个参数都有特定的硬件相关对齐限制。如果使用cudaMalloc(..)，则偏移量保证为0。使用cudaMallocPitch(..)检索音高。您还需要确保您的主机内存以相同的方式倾斜，否则您的memcpy将无法按预期工作。

提示2。

了解二维索引。当访问W[i][j]中的元素时，您需要知道元素W[i][j+1]是内存中的下一个元素，而不是W[i+1][j]。

提示3。

使用一维数组并自己计算2D索引。这会给你更好的控制。