global memoryからshared memoryへの大量の読み込み（補足）

昨日の説明は文章だけでわかりにくかったから、矩形領域をglobal memoryからshared memoryへ読み込むコードと結果の追加。
uint4で読むときには一列をhalf warp(16thread)の１回の読み込みで済ましている。(今回は横幅の最大を64要素に制限しているので)
intで読むときは1warp (32thread)が一列の読み込みを担当する。
この時に横幅が32要素以上あるときにはループが必要になる。

uint4の方が速い場合があるという自分の主張のためにintで読み込む関数は、殆どチューニングしていない。
"常にintが速い派"の人がint側をチューニングすると結果は変わるかもしれない。

#include 
#include 

#define TEST_TIMES (1024*64)
#define MAX_WIDTH  64
#define MAX_HEIGHT 48

__shared__ int buffer[MAX_WIDTH * MAX_HEIGHT];

/* uint4型で読み込む */
__global__ void load_uint4(int *in, int w, int h)
{
    int i, j;

    for (i=0; i> 4;
        int x = tid & 15;
        int h_ofs = blockDim.x >> 4;

        if (x * 4 < w){
            for (; line < h; line += h_ofs){
                int *a = in + w * line + 4 * x;
                int *b = buffer + w * line + 4 * x;
                *(uint4*)b = *(uint4*)a;
            }
        }
        __syncthreads();
    }
}

/* int型で読みこむ */
__global__ void loat_int(int *in, int w, int h)
{
    int i, j;

    for (i=0; i> 5;
        int x = tid & 31;
        int x_loop = (w+31) >> 5;
        int y_ofs = (blockDim.x >> 5);

        for (;line < h; line += y_ofs){
            int *a = in + w * line + x;
            int *b = buffer + w * line + x;

            for (j = 0; j < x_loop; j++){
                if (x + j * 32 < w){
                    int v;
                    v = a[j * 32];
                    b[j * 32] = v;
                }
            }
        }
        __syncthreads();
    }
}

int main()
{
    void *p;
    cudaEvent_t s, e;
    int wtbl = {32, 48, 64};
    int htbl = {4, 5, 8, 9, 16, 17, 32, 48};
    int i, j;

    cudaEventCreate(&s);
    cudaEventCreate(&e);

    cudaMalloc(&p, MAX_WIDTH*MAX_HEIGHT*sizeof(int));

    float v1, v2;
    for (i = 0;i < sizeof(htbl)/sizeof(int); i++){
        for (j = 0; j < sizeof(wtbl)/sizeof(int); j++){
            cudaEventRecord(s, 0);
            load_uint4<<<1,64>>>((int*)p, wtbl[j], htbl[i]);
            cudaEventRecord(e, 0);

            cudaEventSynchronize(e);
            cudaEventElapsedTime(&v1, s, e);

            cudaEventRecord(s, 0);
            load_int<<<1,64>>>((int*)p, wtbl[j], htbl[i]);
            cudaEventRecord(e, 0);

            cudaEventSynchronize(e);
            cudaEventElapsedTime(&v2, s, e);

            printf("[%2d x %2d] uint4=%f int=%f\n", wtbl[j], htbl[i], v1, v2);
        }
    }
}

上のコードを実行するとこうなる。
[]の中は読み込む領域の幅ｘ高さ。
実行にかかった時間の単位はms

[32 x 4] uint4=66.250656 int=119.388702
[48 x 4] uint4=73.464035 int=192.734009
[64 x 4] uint4=78.324257 int=192.731491
[32 x 5] uint4=104.425728 int=171.324326
[48 x 5] uint4=112.621086 int=278.363129
[64 x 5] uint4=123.332573 int=279.045105
[32 x 8] uint4=108.605698 int=223.904190
[48 x 8] uint4=119.630241 int=365.088013
[64 x 8] uint4=132.729858 int=367.008209
[32 x 9] uint4=144.567932 int=272.575806
[48 x 9] uint4=157.935013 int=458.384735
[64 x 9] uint4=175.682175 int=455.767151
[32 x 16] uint4=188.546112 int=428.510803
[48 x 16] uint4=209.417053 int=725.599609
[64 x 16] uint4=232.570526 int=713.375244
[32 x 17] uint4=226.004807 int=484.262482
[48 x 17] uint4=247.822525 int=812.188965
[64 x 17] uint4=270.601349 int=806.044556
[32 x 32] uint4=349.365509 int=845.310425
[48 x 32] uint4=385.352203 int=1414.444824
[64 x 32] uint4=429.602112 int=1444.995605
[32 x 48] uint4=508.608246 int=1259.764771
[48 x 48] uint4=557.349609 int=2130.817871
[64 x 48] uint4=624.313965 int=2143.508789