/********************************************************************
    created:    2009/09/19
    created:    19:9:2009   12:00
    filename:     vector_matrix_multiplication.cu
    file base:    vector_matrix_multiplication
    file ext:    CUDA
    author:        zhao.kaiyong(at)gmail.com
    purpose:    vector matrix multiplication
    copyright:    everyone can use this code, please specify source.
                任意使用，请注明出处;
http://www.hpctech.com
http://openhero.net
*********************************************************************/

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cutil.h>

/************************************************************************/
/* Init CUDA                                                            */
/************************************************************************/
#if __DEVICE_EMULATION__

bool InitCUDA(void){return true;}

#else
bool InitCUDA(void)
{
    int count = 0;
    int i = 0;

    cudaGetDeviceCount(&count);
    if(count == 0) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    for(i = 0; i < count; i++) {
        cudaDeviceProp prop;
        if(cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if(prop.major >= 1) {
                break;
            }
        }
    }
    if(i == count) {
        fprintf(stderr, "There is no device supporting CUDA.\n");
        return false;
    }
    cudaSetDevice(i);

    printf("CUDA initialized.\n");
    return true;
}

#endif
/************************************************************************/
/* Example                                                              */
/************************************************************************/
__global__ static void vector_matrix_mult_kernel(float* A, long wA, float* B, long wB, float* C)
{
    __shared__ float subA[64];

    A = A + threadIdx.x;
    B = B + blockIdx.x * 64 + threadIdx.x;
    C = C + blockIdx.x * 64 + threadIdx.x;

    float subC = 0.0;

    for (int i = 0; i < wA; i+=64)
    {
        subA[threadIdx.x] = A[i];
        __syncthreads();

#pragma unroll
        for (int j = 0; j < 64; j++, B += wB)
        {
            subC += subA[j] * B[0];
        }
        __syncthreads();
    }

    C[0] = subC;
}

//__global__ static void vector_matrix_mult_kernel_32T(float* A, long wA, float* B, long wB, float* C)
//{

//}

#define  RUN_TEST

#define  WA 32
#define  WB 64
/************************************************************************/
/* HelloCUDA                                                            */
/************************************************************************/
int main(int argc, char* argv[])
{

    if(!InitCUDA()) {
        return 0;
    }

    srand(2009);

    long wA = 64* WA;
    long wB = 64* WB;

    long size_A = wA;
    long size_B = wA*wB;
    long size_C = wB;

    float    *hA = (float*)malloc(sizeof(float) * size_A);
    float    *hB = (float*)malloc(sizeof(float) * size_B);
    float    *hC = (float*)malloc(sizeof(float) * size_C);
    float    *testhC = (float*)malloc(sizeof(float) * size_C);

    for (int i = 0; i < size_A; i++)
    {
        hA[i] = (float)rand()/(float)RAND_MAX;
    }

    for (int i = 0; i < size_B; i++)
    {
        hB[i] = (float)rand()/(float)RAND_MAX;
    }

    float    *dA    = 0;
    float    *dB = 0;
    float    *dC = 0;

    CUDA_SAFE_CALL( cudaMalloc((void**) &dA, sizeof(float) * size_A));
    CUDA_SAFE_CALL( cudaMalloc((void**) &dB, sizeof(float) * size_B));
    CUDA_SAFE_CALL( cudaMalloc((void**) &dC, sizeof(float) * size_C));

    CUDA_SAFE_CALL( cudaMemcpy(dA, hA, sizeof(float)*size_A, cudaMemcpyHostToDevice));
    CUDA_SAFE_CALL( cudaMemcpy(dB, hB, sizeof(float)*size_B, cudaMemcpyHostToDevice));

    unsigned int timer = 0;
    CUT_SAFE_CALL( cutCreateTimer( &timer));
    CUT_SAFE_CALL( cutStartTimer( timer));

    dim3 threads = 64;
    dim3 blocks = wB/64;
    vector_matrix_mult_kernel<<<blocks, threads>>>(dA, wA, dB, wB, dC);

    CUT_CHECK_ERROR("Kernel execution failed\n");

    CUDA_SAFE_CALL( cudaMemcpy(hC, dC, sizeof(float) * size_C, cudaMemcpyDeviceToHost));
    CUT_SAFE_CALL( cutStopTimer( timer));
    printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutResetTimer( timer));

    for (int i = 0; i < wB; i++)
    {
        float subC = 0.0;
        for (int j = 0; j <wA; j++)
        {
            subC += hA[j] * hB[j*wB + i];
        }
        testhC[i] = subC;
    }

    CUT_SAFE_CALL( cutStopTimer( timer));
    printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));

    CUT_SAFE_CALL( cutDeleteTimer( timer));

#ifdef RUN_TEST

    CUTBoolean res = cutCompareL2fe(testhC, hC, size_C, 1e-6f);
    printf("Test %s \n", (1 == res) ? "PASSED":"FAILED");
#endif

    CUDA_SAFE_CALL( cudaFree(dA));
    CUDA_SAFE_CALL( cudaFree(dB));
    CUDA_SAFE_CALL( cudaFree(dC));
    free(hA);
    free(hB);
    free(hC);
    CUT_EXIT(argc, argv);

    return 0;
}

http://cuda.csdn.net/Contest/pro/nvidia_results.html

历时5个月的NVIDIA CUDA 在8月落下帷幕，大赛举办期间。有300多个学校700多名学生报名。

最终结果

特别贡献奖

奖 别 姓 名 学 校 参赛作品

特别贡献奖 彭江锋 施少怀 漆舒汉 杨植群 华南理工大学/软件学院 基于GPU的多模式网页精确匹配系统

自选赛奖项

奖 别 姓 名 学 校 参赛作品

一等奖 陈实富 中国科学院/深圳先进技术研究院 Glirt:基于CUDA的多模态三维医学图像配准

二等奖 陈伯君 北京大学/工学院 基于CUDA的DSMC高性能并行计算

二等奖 喻勤 电子科技大学 /电子工程学院 基于CUDA的高性能SAR成像模拟

三等奖 胡晓菡 上海交通大学 /生命科学技术学院 CUDA平台下的复杂疾病全基因组基因相互作用计算

三等奖 陆扬 中国科学技术大学 /工程科学学院 使用CUDA模拟顶点、几何着色器-原型演示

命题赛奖项

奖 别 姓 名 学 校 参赛作品

一等奖 翟艳堂 中国科学院 /计算机网络信息中心 稀疏大矩阵与矢量的乘积运算

二等奖 包南森 上海大学/计算机学院 粒子邻居搜索

二等奖 尚书杰 谢迪 浙江大学 /计算机学院 稀疏大矩阵与矢量的乘积运算

三等奖 褚艳利 邵帅 年华 西安电子科技大学 /计算机学院 基于CSR表示的稀疏矩阵-矢量乘法的CUDA实现

三等奖 张广勇 内蒙古大学 /计算机学院 基于GPU 平台的稀疏矩阵与矢量乘积

优秀作品奖

姓 名 学 校 参赛作品

朱小松 大连理工大学 /土木水利学院 基于CUDA的有限差分法求解简单热传递问题

冯涛 沈乐 清华大学 /工程物理系 CUDA 平台下三维反应堆中子通量分布和有效增

尹康学 长安大学 /信息学院 稀疏大矩阵与矢量相乘基于CUDA的一种可行算法

朱德东 复旦大学 /软件学院 基于CUDA的多分辨率图像融合算法

杨植群 华南理工大学 /计算机科学与工程学院 粒子邻居搜索

石丹 范正娟 赵安元 四川大学 /软件学院 CUDA平台下的超声图像斑点噪声抑制

陈晓熹 厦门大学 /信息科学与技术学院 CuParcone

李斌 段元泽 云南大学 /资源环境与地球科学学院 即时计算的动态分形屏保

优秀作品奖得主将会获得由NVIDIA 中国公司颁发的荣誉证书

恭喜以上获奖选手，请保持联系方式畅通，近期会与您联系相关事宜。