2

単純なボックスカー フィルタを実装しているのは、2D ローカル テクスチャとグローバル メモリ アクセスの速度の違いを評価する言い訳としてのみです。

より詳細には、.cuファイルは次のとおりです

#include <cuda.h>  
#include <cuda_runtime.h>
#include "cufft.h"
#include "Kernels_Test_Texture_Float.cuh"

#define BLOCK_SIZE_x 16
#define BLOCK_SIZE_y 16

/**********************/
/* TEST TEXTURE FLOAT */
/**********************/
extern "C" void Function_Test_Texture_Float(float* data, float* dev_result, int N1, int N2){

    size_t pitch; 
    float* data_d;
    cudaMallocPitch((void**)&data_d,&pitch, N1 * sizeof(float), N2);
    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    cudaBindTexture2D(0,&data_d_texture,data_d,&desc,N1,N2,pitch);
    cudaMemcpy2D(data_d,pitch,data,sizeof(float)*N1,sizeof(float)*N1,N2,cudaMemcpyHostToDevice);

    cudaMemset(dev_result,0,sizeof(float)*N1*N2);
    dim3 dimBlock(BLOCK_SIZE_x,BLOCK_SIZE_y); dim3 dimGrid(N1/BLOCK_SIZE_x + (N1%BLOCK_SIZE_x == 0 ? 0:1),N2/BLOCK_SIZE_x + (N2%BLOCK_SIZE_x == 0 ? 0:1));
    Kernel_Test_Texture_Float<<<dimGrid,dimBlock>>>(dev_result,N1, N2);

}

/**************/
/* TEST FLOAT */
/**************/
extern "C" void Function_Test_Float(float* data, float* dev_result2, int N1, int N2){

    float* data_d;  cudaMalloc((void**)&data_d,sizeof(float)*N1*N2);
    cudaMemcpy(data_d,data,sizeof(float)*N1*N2,cudaMemcpyHostToDevice); 

    cudaMemset(dev_result2,0,sizeof(float)*N1*N2);
    dim3 dimBlock(BLOCK_SIZE_x,BLOCK_SIZE_y); dim3 dimGrid(N1/BLOCK_SIZE_x + (N1%BLOCK_SIZE_x == 0 ? 0:1),N2/BLOCK_SIZE_x + (N2%BLOCK_SIZE_x == 0 ? 0:1));
    Kernel_Test_Float<<<dimGrid,dimBlock>>>(dev_result2,data_d,N1, N2);

}

.cuhファイルは以下

texture<float,2> data_d_texture;

/**************************/
/* 2D TEXTURE TEST KERNEL */
/**************************/
__global__ void Kernel_Test_Texture_Float(float* dev_result, int N1, int N2)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;

    float datum, accumulator=0.;

    int size_x=5;
    int size_y=5;

    if((i<(N1-size_x))&&(j<(N2-size_y)))
    {
        for (int k=0; k<size_x; k++)
        for (int l=0; l<size_y; l++){
            datum = tex2D(data_d_texture,i+k,j+l);
            accumulator = accumulator + datum;
        }
        dev_result[j*blockDim.x*gridDim.x+i] = accumulator;
    }
}

/******************/
/* 2D TEST KERNEL */
/******************/
__global__ void Kernel_Test_Float(float* dev_result2, float* data_d, int N1, int N2)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;

    float accumulator=0.;

    int size_x=5;
    int size_y=5;

    if((i<(N1-size_x))&&(j<(N2-size_y)))
    {
        for (int k=0; k<size_x; k++)
            for (int l=0; l<size_y; l++){
                accumulator = accumulator + data_d[(j+l)*blockDim.x*gridDim.x+(i+k)];
        }
        dev_result2[j*blockDim.x*gridDim.x+i] = accumulator;
    }
}

ただし、グローバル メモリ カーネルはテクスチャ メモリ カーネルよりもはるかに高速です ( 94usvs 615us- タイミングは Visual Profiler の結果です - カードは GeForce GT 540M です)。

私が行っているテクスチャ メモリの使用に何か問題がありますか?それともグローバル メモリはキャッシュされているテクスチャよりも実際に高速ですか?

コメントをお寄せいただきありがとうございます。

4

1 に答える 1

0

一般的に、グローバルメモリ読み取りアクセスはテクスチャメモリ読み取りよりも高速です。さらに、GT 540Mには計算機能2.1があるため、グローバルメモリアクセス用に構成可能なL1およびL2キャッシュ階層があります。また、このキャッシュ階層は小さく、テクスチャキャッシュよりも大きくなっています。

これらの2つの側面を考慮すると、グローバルメモリの実装がテクスチャの実装よりも高速であることは驚くべきことではありません。

デバイスのL1/L2キャッシュ階層を考慮していなかった可能性があります。問題のサイズによっては、Kernel_Test_Floatカーネルを起動する前にL1キャッシュを48KBに最大化することで、パフォーマンスを向上させることができます。

cudaThreadSetCacheConfig(cudaFuncCachePreferL1);
于 2013-01-18T12:30:30.060 に答える