/docs/MyDocs

#include <stdlib.h>

#include <stdio.h>

#include <string.h>

 

#include <cutil.h>

 

 

typedef float data_type;

#define matrix_width 4096

#define matrix_height 4096

#define size (matrix_width * matrix_height)

 

#define iters 1

 

#include <test_kernel.cu>

 

 

 

void test() {

    int i,j;

    data_type *a, *b;

 

    unsigned int timer;

    cutCreateTimer(&timer);

 

    data_type *d_a, *d_b, *d_res;

 

/*

    CUDA_SAFE_CALL(cudaMalloc( (void**) &d_a, size*sizeof(data_type)));

    CUDA_SAFE_CALL(cudaMalloc( (void**) &d_b, size*sizeof(data_type)));

    CUDA_SAFE_CALL(cudaMalloc( (void**) &d_res, size*sizeof(data_type)));

 

    CUDA_SAFE_CALL(cudaMemcpy(d_a, a, size*sizeof(data_type),  cudaMemcpyHostToDevice));

    CUDA_SAFE_CALL(cudaMemcpy(d_b, b, size*sizeof(data_type),  cudaMemcpyHostToDevice));

 

    vecAdd<<<1,size>>>(d_res, d_a, d_b);

 

    CUDA_SAFE_CALL(cudaMemcpy(res, d_res, size*sizeof(data_type),cudaMemcpyDeviceToHost));

*/

 

    int block_size = 512;

 

        /* 

        transpose2:

        should be symmetrical, otherwise (if 16x8) after first step of 

        transpose we will have 16 threads to iterate other 8 elements),

        using other dimmension of threads will destroy memory access */

        

    int block_width = 16;

    int block_height = 16;

    int real_block_size = block_width * block_height;

    

    int wblocks = matrix_width / block_width;

    int real_width = wblocks * block_width;

    if (real_width < matrix_width) {

        wblocks++;

        real_width += block_width;

    }

 

    int hblocks = matrix_height / block_height;

    int real_height = hblocks * block_height;

    if (real_height < matrix_height) {

        hblocks++;

        real_height += block_height;

    }

 

    int real_size = real_width * real_height;

 

    dim3 block_dim(block_width, block_height, 1);

    dim3 grid_dim(real_width / block_width, real_height / block_height, 1);

    

    a = (float*) malloc(real_size*sizeof(data_type));

    b = (float*) malloc(real_size*sizeof(data_type));

    

 

 

 

    for (i = 0; i < real_width; i++) {

        for (j = 0; j < real_height; j++) {

            a[i + j*real_width] = i + 0.0001 * j;

            b[i] = 0;

        }

    }

 

 

    CUDA_SAFE_CALL(cudaMalloc( (void**) &d_a, real_size*sizeof(data_type)));

    CUDA_SAFE_CALL(cudaMalloc( (void**) &d_b, real_size*sizeof(data_type)));

 

    CUDA_SAFE_CALL(cudaMemcpy(d_a, a, real_size*sizeof(data_type),  cudaMemcpyHostToDevice));

 

 

    transpose3<<<grid_dim, block_dim, (1 + block_width) * block_height*sizeof(data_type)>>>(d_b, d_a, real_width, real_height);

    cudaThreadSynchronize();

 

    cutResetTimer(timer);

    cutStartTimer(timer);

    for (i=0;i<iters;++i) {

        transpose3<<<grid_dim, block_dim/*, (1 + block_width) * block_height*sizeof(data_type)*/>>>(d_b, d_a, real_width, real_height);

    }

    cudaThreadSynchronize();

    cutStopTimer(timer);

    float time3 = cutGetTimerValue(timer);

 

    transpose2<<<grid_dim, block_dim, (1 + block_width) * block_height*sizeof(data_type)>>>(d_b, d_a, real_width, real_height);

    cudaThreadSynchronize();

 

    cutResetTimer(timer);

    cutStartTimer(timer);

    for (i=0;i<iters;++i) {

        transpose2<<<grid_dim, block_dim, (1 + block_width) * block_height*sizeof(data_type)>>>(d_b, d_a, real_width, real_height);

    }

    cudaThreadSynchronize();

    cutStopTimer(timer);

    float time2 = cutGetTimerValue(timer);

 

 

    transpose1<<<grid_dim, block_dim, (1 + block_width) * block_height*sizeof(data_type)>>>(d_b, d_a, real_width, real_height);

    cudaThreadSynchronize();

 

    cutResetTimer(timer);

    cutStartTimer(timer);

    for (i=0;i<iters;++i) {

        transpose1<<<grid_dim, block_dim, (1 + block_width) * block_height*sizeof(data_type)>>>(d_b, d_a, real_width, real_height);

    }

    cudaThreadSynchronize();

    cutStopTimer(timer);

    float time1 = cutGetTimerValue(timer);

 

 

 

 

 

    CUDA_SAFE_CALL(cudaMemcpy(b, d_b, size*sizeof(data_type),cudaMemcpyDeviceToHost));

    

 

    printf("T1 transpose average time:     %0.3f ms\n", time1 / iters);

    printf("T2 transpose average time:     %0.3f ms\n", time2 / iters);

    printf("T3 transpose average time:     %0.3f ms\n", time3 / iters);

    

//    int runs = blocks / 65535;

    

//    <<<blocks

 

 

    for (i = 0; i < 10; i++) {

        printf("%lf ", b[i]);

    }

 

    printf("\n");

 

}

 

 

 

int

main( int argc, char** argv) 

{

    int deviceCount;

 

    CUDA_SAFE_CALL(cudaGetDeviceCount(&deviceCount));

    if (deviceCount == 0)

        printf("There is no device supporting CUDA\n");

    int dev;

    for (dev = 0; dev < deviceCount; ++dev) {

        cudaDeviceProp deviceProp;

        CUDA_SAFE_CALL(cudaGetDeviceProperties(&deviceProp, dev));

        if (dev == 0) {

            if (deviceProp.major == 9999 && deviceProp.minor == 9999)

                printf("There is no device supporting CUDA.\n");

            else if (deviceCount == 1)

                printf("There is 1 device supporting CUDA\n");

            else

                printf("There are %d devices supporting CUDA\n", deviceCount);

        }

        printf("\nDevice %d: \"%s\"\n", dev, deviceProp.name);

        printf("  Major revision number:                         %d\n",

               deviceProp.major);

        printf("  Minor revision number:                         %d\n",

               deviceProp.minor);

        printf("  Total amount of global memory:                 %u bytes\n",

               deviceProp.totalGlobalMem);

    #if CUDART_VERSION >= 2000

        printf("  Number of multiprocessors:                     %d\n",

               deviceProp.multiProcessorCount);

        printf("  Number of cores:                               %d\n",

               8 * deviceProp.multiProcessorCount);

    #endif

        printf("  Total amount of constant memory:               %u bytes\n",

               deviceProp.totalConstMem); 

        printf("  Total amount of shared memory per block:       %u bytes\n",

               deviceProp.sharedMemPerBlock);

        printf("  Total number of registers available per block: %d\n",

               deviceProp.regsPerBlock);

        printf("  Warp size:                                     %d\n",

               deviceProp.warpSize);

        printf("  Maximum number of threads per block:           %d\n",

               deviceProp.maxThreadsPerBlock);

        printf("  Maximum sizes of each dimension of a block:    %d x %d x %d\n",

               deviceProp.maxThreadsDim[0],

               deviceProp.maxThreadsDim[1],

               deviceProp.maxThreadsDim[2]);

        printf("  Maximum sizes of each dimension of a grid:     %d x %d x %d\n",

               deviceProp.maxGridSize[0],

               deviceProp.maxGridSize[1],

               deviceProp.maxGridSize[2]);

        printf("  Maximum memory pitch:                          %u bytes\n",

               deviceProp.memPitch);

        printf("  Texture alignment:                             %u bytes\n",

               deviceProp.textureAlignment);

        printf("  Clock rate:                                    %.2f GHz\n",

               deviceProp.clockRate * 1e-6f);

    #if CUDART_VERSION >= 2000

        printf("  Concurrent copy and execution:                 %s\n",

               deviceProp.deviceOverlap ? "Yes" : "No");

    #endif

    }

 

    test();

 

//    CUT_EXIT(argc, argv);

}