0

我一直在研究cuda。

在 CUDA 编程指南中，共享内存访问时间比全局内存时间快。

所以，我编写了执行矩阵乘法的代码。

这是我的代码。version1 使用全局内存，version2 使用共享内存

我的gpu是tesla c2070 cuda sdk 4.2版

主要代码

#define Matrix_Width   9216
#define Matrix_Divide  4
#define Tile_Width     32
#define Ce_Size 4096

#if Matrix_Width == 9216
    #define Matrix_Size 9216*9216
#elif Matrix_Width == 12800
    #define Matrix_Size 12800*12800
#elif Matrix_Width == 15872
    #define Matrix_Size 15872*15872
#elif Matrix_Width == 18432
    #define Matrix_Size 18432*18432
#endif

float* H_Input1 = (float*)malloc( sizeof(float) * Matrix_Size );
float* H_Input2 = (float*)malloc( sizeof(float) * Matrix_Size );
float* H_Output = (float*)malloc( sizeof(float) * Matrix_Size );

for( int i=0 ; i < Matrix_Size ; i++ ){
    H_Input1[i] = 1.0f;
H_Input2[i] = 1.0f;
}
memset( H_Output, 0 , sizeof(float) * Matrix_Size );

float* D_Input1;
float* D_Input2;
float* D_Output;

cudaMalloc( (void**)&D_Input1, sizeof(float) * Matrix_Size );
cudaMalloc( (void**)&D_Input2, sizeof(float) * Matrix_Size );
cudaMalloc( (void**)&D_Output, sizeof(float) * Matrix_Size );
cudaMemcpy( D_Input1, H_Input1, sizeof(float) * Matrix_Size, cudaMemcpyHostToDevice );
cudaMemcpy( D_Input2, H_Input2, sizeof(float) * Matrix_Size, cudaMemcpyHostToDevice );
cudaMemcpy( D_Output, H_Output, sizeof(float) * Matrix_Size, cudaMemcpyHostToDevice );


event_pair Event;
start_timer( &Event );
dim3  dimGrid( Matrix_Width/Matrix_Divide/Tile_Width, Matrix_Width/Matrix_Divide/Tile_Width, 1 );
dim3 dimBlock( Tile_Width, Tile_Width, 1 );

kernel_global<< dimGrid, dimBlock>>>( D_Input1, D_Input2, D_Output );
stop_timer( &Event, "1GB mMemory Test\n" );
cudaMemcpy( H_Output, D_Output, sizeof(float) * Matrix_Size, cudaMemcpyDeviceToHost );

内核版本1

__global__ void kernel_global( float* Input1, float* Input2, float* Output ){

for( int i = 0 ; i < Matrix_Divide ; i++ ){
for( int j = 0 ; j < Matrix_Divide ; j++ ){

float Sum = 0;
int Row = (i * (Matrix_Width/Matrix_Divide)) + (blockIdx.y * Tile_Width) + threadIdx.y; 
int Col = (j * (Matrix_Width/Matrix_Divide)) + (blockIdx.x * Tile_Width) + threadIdx.x;

    for( int k = 0 ; k < Matrix_Width ; k++ ){
        Sum += Input1[ Row * Matrix_Width + k ] * Input2[ k * Matrix_Width + Col ];
    }
    Output[ Row*Matrix_Width+Col] = Sum;
    }
    }
}

内核版本2

    __global__ void kernel_shared( float* Input1, float* Input2, float* Output ){

    __shared__ float Input1_s[Tile_Width][Tile_Width];
    __shared__ float Input2_s[Tile_Width][Tile_Width];

    int Bx = blockIdx.x;
    int By = blockIdx.y;
    int Tx = threadIdx.x;
    int Ty = threadIdx.y;

    for( int i = 0 ; i < Matrix_Divide ; i++ ){
        for( int j = 0 ; j < Matrix_Divide ; j++ ){

            float Sum = 0;
            int Row = (i * (Matrix_Width/Matrix_Divide)) + (By * Tile_Width) + Ty;
            int Col = (j * (Matrix_Width/Matrix_Divide)) + (Bx * Tile_Width) + Tx;


            for( int m = 0 ; m < Matrix_Width/Tile_Width ; m++ ){

                Input1_s[Ty][Tx] = Input1[ Row * Matrix_Width + ( m * Tile_Width + Tx ) ];
                Input2_s[Ty][Tx] = Input2[ ( m * Tile_Width + Ty ) * Matrix_Width + Col ];
                __syncthreads();


                for( int k = 0 ; k < Tile_Width; k++ ){
                    Sum += Input1_s[Ty][k] * Input2_s[k][Tx];
                }
                __syncthreads();
            }
            Output[ Row*Matrix_Width+Col] = Sum;
        }
    }
}

此代码制作了 Width = 9216 的矩阵

它不能一次计算。因为块的最大计数是 65535 和线程 1024

所以我使用 4 划分矩阵宽度，所以矩阵划分 16 个块。

一次可以计算 1 个块。

所以我使用loopcount为16的循环（i * j = 16）

并且一个块被划分块和线程..（ tile_width = 32 ）

测试结果好奇怪。

版本 1 耗时 90 秒

版本 2 耗时 130 秒

我无法理解这个结果

我认为共享内存元素在瓷砖中被重新使用......

为什么版本 1 比版本 2 快？

此致！！

4

1 回答 1

3

Tesla C2070 是一款计算能力 2.0 设备，可缓存全局内存访问。因此，在这两种情况下，内部循环都在片上存储器中运行（除了第一次迭代）。

唯一的区别是第二个版本中将数据加载到共享内存中的额外代码需要额外的时间，而在第一个版本中，这项工作是由缓存逻辑在硬件中完成的。

只有在计算能力为 1.x 的设备上，或者如果您比硬件缓存的 LRU（最近最少使用）逻辑更清楚哪些数据正在被重用，手动缓存数据到共享内存中才值得。

于 2012-09-21T09:42:42.640 回答