1

cudaArrayテクスチャ参照/オブジェクトでラップせずにデバイスからaの値を読み取る方法はありますか？私が見たすべての例は、cudaArrayテクスチャを作成するためだけに使用されています。それがそれらを使用できる唯一の方法ですか、または私は次のようなことをすることができます：

__global__ kernel(cudaArray *arr, ...) {
    float x = tex1D<float>(arr, ...);
    ...
}

cudaArray *arr;
cudaMallocArray(&arr, ...);
cudaMemcpyToArray(arr, ...);
kernel<<<...>>>(arr, ...);

それで、基本的に、そこの代わりに何をすべきtex1Dですか？また、これが可能であれば、これを行うことでパフォーマンス上の利点があると誰かが考えているかどうか知りたいと思いますが、私自身のテストも実行して確認します。

ありがとう！

4

2 に答える 2

7

cudaArrayは、テクスチャリングまたはサーフェスメモリの目的で定義されています。ここに示されているように：

CUDA配列は、テクスチャフェッチ用に最適化された不透明なメモリレイアウトです。それらは1次元、2次元、または3次元であり、要素で構成されます。各要素には、符号付きまたは符号なしの8、16、または32ビット整数、16ビットフロート、または32ビットフロートの1、2、または4つのコンポーネントがあります。CUDA配列には、テクスチャメモリで説明されているテクスチャフェッチ、またはサーフェスメモリで説明されているサーフェスの読み取りと書き込みを介してカーネルからのみアクセスできます。

したがって、実際には、cudaArrayのデータにアクセスするには、カーネルでテクスチャ関数またはサーフェス関数のいずれかを使用する必要があります。

テクスチャリングの使用に関連するパフォーマンス上の利点はいくつかあります。テクスチャリングは、補間を意味する場合があります（つまり、浮動小数点座標を使用してテクスチャから読み取る）。この種のデータ補間を必要とするアプリケーションは、GPUのテクスチャユニット内のHW補間エンジンの恩恵を受ける可能性があります。

別の利点は、おそらく任意のGPUコードでテクスチャリングを使用するために最も重要なことですが、グローバルメモリに格納されているテクスチャをバックアップするテクスチャキャッシュです。テクスチャリングは読み取り専用の操作ですが、読み取り専用のデータの配列がある場合、テクスチャキャッシュは、データにすばやくアクセスする能力を向上または拡張する可能性があります。これは通常、テクスチャリングメカニズムに格納されているデータにアクセスしている関数にデータの局所性/データの再利用が必要であることを意味します。取得されたテクスチャデータはL1キャッシュ内の何も中断しないため、通常、この種のデータのセグメンテーション/最適化は、データキャッシングに関するより大きな戦略の一部になります。L1キャッシュに他の要求がなかった場合、テクスチャメカニズム/キャッシュは、すでにL1にある場合よりもデータへの高速アクセスを提供しません。

于 2013-02-18T00:23:06.757 に答える
5

RobertCrovellaはすでにあなたの質問に答えています。次のユーザーにとって、テクスチャと表面の2つのソリューションの実例があると便利だと思います。

#include <stdio.h>
#include <thrust\device_vector.h>

// --- 2D float texture
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;

// --- 2D surface memory
surface<void, 2> surf2D;

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, 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);
    }
}

/*************************************/
/* cudaArray PRINTOUT TEXTURE KERNEL */
/*************************************/
__global__ void cudaArrayPrintoutTexture(int width, int height)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    printf("Thread index: (%i, %i); cudaArray = %f\n", x, y, tex2D(texRef, x / (float)width + 0.5f, y / (float)height + 0.5f));
}

/*************************************/
/* cudaArray PRINTOUT TEXTURE KERNEL */
/*************************************/
__global__ void cudaArrayPrintoutSurface(int width, int height)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    float temp;

    surf2Dread(&temp, surf2D, x * 4, y);

    printf("Thread index: (%i, %i); cudaArray = %f\n", x, y, temp);
}

/********/
/* MAIN */
/********/
void main()
{
    int width = 3, height = 3;

    thrust::host_vector<float> h_data(width*height, 3.f);

    // --- Allocate CUDA array in device memory
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);

    cudaArray* cuArray;

    /*******************/
    /* TEXTURE BINDING */
    /*******************/
    gpuErrchk(cudaMallocArray(&cuArray, &channelDesc, width, height));

    // --- Copy to host data to device memory
    gpuErrchk(cudaMemcpyToArray(cuArray, 0, 0, thrust::raw_pointer_cast(h_data.data()), width*height*sizeof(float), cudaMemcpyHostToDevice));

    // --- Set texture parameters
    texRef.addressMode[0] = cudaAddressModeWrap;
    texRef.addressMode[1] = cudaAddressModeWrap;
    texRef.filterMode = cudaFilterModeLinear;
    texRef.normalized = true;

    // --- Bind the array to the texture reference
    gpuErrchk(cudaBindTextureToArray(texRef, cuArray, channelDesc));

    // --- Invoking printout kernel
    dim3 dimBlock(3, 3);
    dim3 dimGrid(1, 1);
    cudaArrayPrintoutTexture<<<dimGrid, dimBlock>>>(width, height);

    gpuErrchk(cudaUnbindTexture(texRef));

    gpuErrchk(cudaFreeArray(cuArray));

    /******************/
    /* SURFACE MEMORY */
    /******************/
    gpuErrchk(cudaMallocArray(&cuArray, &channelDesc, width, height, cudaArraySurfaceLoadStore));

    // --- Copy to host data to device memory
    gpuErrchk(cudaMemcpyToArray(cuArray, 0, 0, thrust::raw_pointer_cast(h_data.data()), width*height*sizeof(float), cudaMemcpyHostToDevice));

    gpuErrchk(cudaBindSurfaceToArray(surf2D, cuArray));

    cudaArrayPrintoutSurface<<<dimGrid, dimBlock>>>(width, height);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaFreeArray(cuArray));
}
于 2014-07-08T17:17:25.220 に答える