CUDA编程练习

最新推荐文章于 2023-07-24 14:09:41 发布
最新推荐文章于 2023-07-24 14:09:41 发布
阅读量490 收藏 5
点赞数
分类专栏： CUDA 文章标签： cuda
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/weixin_42764932/article/details/115641500
版权

文章目录

1. 设定GPU

#define CHECK(call) \
{                   \
    const cudaError_t error = call; \
    if(error!=cudaSuccess)          \
    {               \
    printf("error: %s: %d, ", __FILE__, __LINE__); \
    printf("code:%d, reason: %S\n", error, cudaGetErrorString(error)); \
    exit(1);\
    }\
}
#include <stdio.h>
#include <cuda_runtime.h>

int main() {
    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("Using device %d: %s\n",dev, deviceProp.name);
    CHECK(cudaSetDevice(dev));

    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

2. 简单的向量相加

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

void initdata(float *p, int size)
{
    for(int i=0; i<size; i++)
    {
        p[i] = (float)(rand()&0xFF) / 10.0f;
    }
}
__global__ void sumarray(float *a, float *b, float *c, const int n)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if(idx<n)
        c[idx] = a[idx] + b[idx];
}
void sumOnCpu(float *a, float *b, float *c, const int n) {
    for (int i = 0; i < n; i++) {

        c[i] = a[i] + b[i];
    }
}
void sumprintf(float *a, float *b, float *c, const int n) {
    for (int i = 0; i < n; i++) {
        printf("%f + %f = %f\n",a[i], b[i], c[i]);
    }
}
void check(float *a, float *b, const int n)
{
    double e = 1.0E-8;
    for(int i = 0; i < n; i++)
    {
        if(abs(a[i] - b[i]) > e)
        {
            printf("\ndo not match");
            printf("host %5.2f gpu %5.2f at current %d\n", a[i],
                   b[i], i);
            return;
        }
    }
    printf("\nsuccess\n");
    return;
}

int main() {
    int dev = 0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    cudaSetDevice(dev);
    printf("device %d: %s",dev, deviceProp.name);

    int nElem = 1 << 4;
    int nbytes = sizeof(float) * nElem;

    float *ha, *hb, *res_g, *hc;
    ha = (float *)malloc(nbytes);
    hb = (float *)malloc(nbytes);
    res_g = (float *)malloc(nbytes);
    hc = (float *)malloc(nbytes);


    initdata(ha, nElem);
    initdata(hb, nElem);
    memset(res_g, 0, nbytes);
    memset(hc, 0, nbytes);

    float *da, *db, *dc;
    cudaMalloc((float**)&da, nbytes);
    cudaMalloc((float**)&db, nbytes);
    cudaMalloc((float**)&dc, nbytes);

    cudaMemcpy(da, ha, nbytes, cudaMemcpyHostToDevice);
    cudaMemcpy(db, hb, nbytes, cudaMemcpyHostToDevice);


    dim3 block(32);
    dim3 grid((nElem + block.x - 1)/block.x);

    sumarray<<<grid, block>>>(da, db, dc, nElem);
    cudaMemcpy(res_g, dc, nbytes, cudaMemcpyDeviceToHost);

    sumOnCpu(ha, hb, hc, nElem);
    check(res_g, hc, nElem);

//    sumprintf(ha, hb, hc, nElem);
//    sumprintf(ha, hb, res_g, nElem);

    cudaFree(da);
    cudaFree(db);
    cudaFree(dc);
    free(ha);
    free(hb);
    free(res_g);

    return EXIT_SUCCESS;
}

3. 简单的矩阵相加

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

void initdata(float *p, int size)
{
    for(int i=0; i<size; i++)
    {
        p[i] = (float)(rand()&0xFF) / 10.0f;
    }
}

__global__ void summetrix(float *a, float *b, float *c, int nx, int ny)
{
    unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
    unsigned int idx = ix + iy * nx;
    if(ix<nx&&iy<ny)
        c[idx] = a[idx] + b[idx];
}
void sumMatrixOnHost(float *A, float *B, float *C, const int nx,
                     const int ny)
{
    float *ia = A;
    float *ib = B;
    float *ic = C;

    for (int iy = 0; iy < ny; iy++)
    {
        for (int ix = 0; ix < nx; ix++)
        {
            ic[ix] = ia[ix] + ib[ix];

        }

        ia += nx;
        ib += nx;
        ic += nx;
    }

    return;
}


void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
            match = 0;
            printf("host %f gpu %f\n", hostRef[i], gpuRef[i]);
            break;
        }
    }

    if (match)
        printf("Arrays match.\n\n");
    else
        printf("Arrays do not match.\n\n");
}

int main() {
    int dev;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    cudaSetDevice(dev);

    int nx = 1<<8;
    int ny = 1<<8;
    int nElem = nx * ny;
    int nbytes = sizeof(float) * nElem;

    float *ha, *hb, *hc, *res_g;
    ha = (float*)malloc(nbytes);
    hb = (float*)malloc(nbytes);
    hc = (float*)malloc(nbytes);
    res_g = (float*)malloc(nbytes);

    initdata(ha, nElem);
    initdata(hb, nElem);
    memset(res_g, 0, nbytes);
    memset(hc, 0, nbytes);

    float *da, *db, *dc;
    cudaMalloc((float**)&da, nbytes);
    cudaMalloc((float**)&db, nbytes);
    cudaMalloc((float**)&dc, nbytes);
    cudaMemcpy(da, ha, nbytes, cudaMemcpyHostToDevice);
    cudaMemcpy(db, hb, nbytes, cudaMemcpyHostToDevice);

    dim3 block(32, 32);
    dim3 grid((nx + block.x -1)/block.x, (ny+block.y-1)/block.y);
    summetrix<<<grid, block>>>(da, db, dc, nx, ny);

    cudaMemcpy(res_g, dc, nbytes, cudaMemcpyDeviceToHost);

    sumMatrixOnHost(ha, hb, hc, nx, ny);
    checkResult(res_g, hc, nElem);


    cudaFree(da);
    cudaFree(db);
    cudaFree(dc);
    free(ha);
    free(hb);
    free(res_g);

    cudaDeviceReset();
    return 0;
}

4. 相邻归约，向量求和

#include <stdio.h>

__global__ void sumgpu(float *a, float *b, int n)
{
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;
    unsigned int tid = threadIdx.x;


    float *idata = a + blockIdx.x * blockDim.x;
    if(idx > n)return;
    for(int stride=1; stride<blockDim.x; stride*=2)
    {
        if((tid%(stride*2))==0)
            idata[tid] += idata[tid+stride];
        __syncthreads();
    }
    if(tid==0)
        b[blockIdx.x] = idata[0];
}


void initdata(float *a, int n)
{
    for(int i=0; i<n; i++)
    {
        a[i] = (float)(rand()&0xFF) / 100.0f;
    }
}
int main(){
    int dev = 0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("gpu: %d, name: %s\n", dev, deviceProp.name);
    cudaSetDevice(dev);

    int nElem = 1<<12;
    dim3 block(32);
    dim3 grid((nElem + block.x - 1)/block.x);
    int nbytes = sizeof(float) * nElem;
    float *h_a, *h_g;
    h_a = (float *)malloc(nbytes);
    h_g = (float *)malloc(sizeof(float) * grid.x);
    initdata(h_a, nElem);
    memset(h_g, 0, sizeof(float) * grid.x);

    float *d_a, *d_b;
    cudaMalloc((float **)&d_a, nbytes);
    cudaMalloc((float **)&d_b, sizeof(float) * grid.x);
    cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice);
    cudaMemset(d_b, 0, sizeof(float) * grid.x);

    float res = 0.0;
    for(int i=0; i<nElem; i++)
        res += h_a[i];
    printf("\nsum on cpu:%f\n", res);

    sumgpu<<<grid, block>>>(d_a, d_b, nElem);
    cudaMemcpy(h_g, d_b, sizeof(float) * grid.x, cudaMemcpyDeviceToHost);
    float h_res = 0.0;
    for(int i=0; i<grid.x; i++)
        h_res += h_g[i];
    printf("\nsum on gpu:%f\n", h_res);


    cudaFree(d_a);
    cudaFree(d_b);
    free(h_a);
    free(h_g);
    cudaDeviceReset();
}

5. 相邻归约，改善分化

#include <stdio.h>

__global__ void sumgpu2(float *a, float *b, int n)
{
    unsigned int tid = threadIdx.x;
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;
    if(idx > n)return;

    float *idata = a + blockIdx.x * blockDim.x;
    for(int stride=1; stride<blockDim.x; stride*=2)
    {
        unsigned int index = 2 * stride * tid;
        if(index < blockDim.x)
            idata[index] += idata[index + stride];
        __syncthreads();
    }
    if(tid == 0)b[blockIdx.x] = idata[0];
}

void initdata(float *a, int n)
{
    for(int i=0; i<n; i++)
    {
        a[i] = (float)(rand()&0xFF) / 100.0f;
    }
}
int main(){
    int dev = 0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("gpu: %d, name: %s\n", dev, deviceProp.name);
    cudaSetDevice(dev);

    int nElem = 1<<12;
    dim3 block(32);
    dim3 grid((nElem + block.x - 1)/block.x);
    int nbytes = sizeof(float) * nElem;
    float *h_a, *h_g;
    h_a = (float *)malloc(nbytes);
    h_g = (float *)malloc(sizeof(float) * grid.x);
    initdata(h_a, nElem);
    memset(h_g, 0, sizeof(float) * grid.x);

    float *d_a, *d_b;
    cudaMalloc((float **)&d_a, nbytes);
    cudaMalloc((float **)&d_b, sizeof(float) * grid.x);
    cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice);
    cudaMemset(d_b, 0, sizeof(float) * grid.x);

    float res = 0.0;
    for(int i=0; i<nElem; i++)
        res += h_a[i];
    printf("\nsum on cpu:%f\n", res);

    sumgpu2<<<grid, block>>>(d_a, d_b, nElem);
    cudaMemcpy(h_g, d_b, sizeof(float) * grid.x, cudaMemcpyDeviceToHost);
    float h_res = 0.0;
    for(int i=0; i<grid.x; i++)
        h_res += h_g[i];
    printf("\nsum on gpu:%f\n", h_res);


    cudaFree(d_a);
    cudaFree(d_b);
    free(h_a);
    free(h_g);
    cudaDeviceReset();
}

6. 交错归约

#include <stdio.h>

__global__ void sumgpu3(float *a, float *b, int n)
{
    unsigned int tid = threadIdx.x;
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;

    if(idx>n)return;
    float *idata = a + blockDim.x * blockIdx.x;
    for(int stride=blockDim.x/2; stride>0; stride>>=1)
    {
        if(tid<stride)
            idata[tid] += idata[tid+stride];
        __syncthreads();
    }
    if(tid==0)b[blockIdx.x] = idata[0];
}

void initdata(float *a, int n)
{
    for(int i=0; i<n; i++)
    {
        a[i] = (float)(rand()&0xFF) / 100.0f;
    }
}
int main(){
    int dev = 0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("gpu: %d, name: %s\n", dev, deviceProp.name);
    cudaSetDevice(dev);

    int nElem = 1<<12;
    dim3 block(32);
    dim3 grid((nElem + block.x - 1)/block.x);
    int nbytes = sizeof(float) * nElem;
    float *h_a, *h_g;
    h_a = (float *)malloc(nbytes);
    h_g = (float *)malloc(sizeof(float) * grid.x);
    initdata(h_a, nElem);
    memset(h_g, 0, sizeof(float) * grid.x);

    float *d_a, *d_b;
    cudaMalloc((float **)&d_a, nbytes);
    cudaMalloc((float **)&d_b, sizeof(float) * grid.x);
    cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice);
    cudaMemset(d_b, 0, sizeof(float) * grid.x);

    float res = 0.0;
    for(int i=0; i<nElem; i++)
        res += h_a[i];
    printf("\nsum on cpu:%f\n", res);

    sumgpu3<<<grid, block>>>(d_a, d_b, nElem);
    cudaMemcpy(h_g, d_b, sizeof(float) * grid.x, cudaMemcpyDeviceToHost);
    float h_res = 0.0;
    for(int i=0; i<grid.x; i++)
        h_res += h_g[i];
    printf("\nsum on gpu:%f\n", h_res);


    cudaFree(d_a);
    cudaFree(d_b);
    free(h_a);
    free(h_g);
    cudaDeviceReset();
}

7. 展开归约

先将块叠加，即展开归约，然后再块内交错归约。提高内存吞吐量。内存延迟可以更好地被隐藏起来。

#include <stdio.h>
void initdata(float *a, int n)
{
    for(int i=0; i<n; i++)
    {
        a[i] = (float)(rand()&0xFF) / 100.f;
    }
}

__global__ void sumInterleaveUnroll(float *d, float *tmp, int n)
{
    unsigned int tid = threadIdx.x;
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x * 8;


    if(idx + blockDim.x * 7 < n)
    {
        float a1 = d[idx];
        float a2 = d[idx + blockDim.x];
        float a3 = d[idx + blockDim.x * 2];
        float a4 = d[idx + blockDim.x * 3];
        float a5 = d[idx + blockDim.x * 4];
        float a6 = d[idx + blockDim.x * 5];
        float a7 = d[idx + blockDim.x * 6];
        float a8 = d[idx + blockDim.x * 7];
        d[idx] = a1 + a2 + a3 + a4 + a5 + a6 + a7 + a8;
    }
    __syncthreads();
    float *idata = d + blockDim.x * blockIdx.x * 8;
    for(int stride=blockDim.x/2; stride>0; stride>>=1)
    {
        if(tid < stride)
            idata[tid] += idata[tid + stride];
        __syncthreads();
    }
    if(tid==0)tmp[blockIdx.x] = idata[0];
}
int main()
{
    int dev=0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("gpu: %d, name: %s \n", dev, deviceProp.name);
    cudaSetDevice(dev);

    int nElem = 1<<12;
    int nbytes = sizeof(float) * nElem;
    dim3 block(32);
    dim3 grid((nElem + block.x - 1) / block.x);

    float *ha, *htmp;
    ha = (float*)malloc(nbytes);
    htmp = (float*)malloc(sizeof(float)*grid.x / 8);
    initdata(ha, nElem);
    memset(htmp, 0, sizeof(float)*grid.x / 8);

    float *da, *dtmp;
    cudaMalloc((float**)&da, nbytes);
    cudaMalloc((float**)&dtmp, sizeof(float)*grid.x / 8);
    cudaMemcpy(da, ha, nbytes, cudaMemcpyHostToDevice);
    cudaMemset(dtmp, 0, sizeof(float)*grid.x / 8);

    float res = 0.0;
    for(int i=0; i<nElem; i++)
        res += ha[i];
    printf("correct sum: %f \n", res);

    sumInterleaveUnroll<<<grid.x / 8, block>>>(da, dtmp, nElem);

    cudaMemcpy(htmp, dtmp, sizeof(float)*grid.x / 8, cudaMemcpyDeviceToHost);
    float res_d = 0.0;
    for(int i=0; i<grid.x / 8; i++)
        res_d += htmp[i];
    printf("gpu sum: %f \n", res_d);


    cudaFree(da);
    cudaFree(dtmp);
    free(ha);
    free(htmp);
}

8. 展开线程的归约

用线程束内隐式同步代替部分__syncthreads()

#include <stdio.h>
void initdata(float *a, int n)
{
    for(int i=0; i<n; i++)
    {
        a[i] = (float)(rand()&0xFF) / 100.f;
    }
}


__global__ void sumInterleaveUnrollWarp(float *d, float *tmp, int n)
{
    unsigned int tid = threadIdx.x;
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x * 8;
    if(idx + 7 * blockDim.x<n)
    {
        float a1 = d[idx];
        float a2 = d[idx + blockDim.x];
        float a3 = d[idx + blockDim.x * 2];
        float a4 = d[idx + blockDim.x * 3];
        float a5 = d[idx + blockDim.x * 4];
        float a6 = d[idx + blockDim.x * 5];
        float a7 = d[idx + blockDim.x * 6];
        float a8 = d[idx + blockDim.x * 7];
        d[idx] = a1 + a2 + a3 + a4 + a5 + a6 + a7 + a8;
    }
    __syncthreads();
    float *idata = d + 8 * blockDim.x * blockIdx.x;
    for(int stride=blockDim.x/2; stride>32; stride>>=1)
    {
        if(tid < stride)
            idata[tid] += idata[tid + stride];
        __syncthreads();
    }
    if(tid<32)
    {
        volatile float *vmem = idata;
        vmem[tid] += vmem[tid + 32];
        vmem[tid] += vmem[tid + 16];
        vmem[tid] += vmem[tid + 8];
        vmem[tid] += vmem[tid + 4];
        vmem[tid] += vmem[tid + 2];
        vmem[tid] += vmem[tid + 1];
    }
    if(tid==0)tmp[blockIdx.x] = idata[0];
}
int main()
{
    int dev=0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("gpu: %d, name: %s \n", dev, deviceProp.name);
    cudaSetDevice(dev);

    int nElem = 1<<14;
    int nbytes = sizeof(float) * nElem;
    dim3 block(64);
    dim3 grid((nElem + block.x - 1) / block.x);

    float *ha, *htmp;
    ha = (float*)malloc(nbytes);
    htmp = (float*)malloc(sizeof(float)*grid.x / 8);
    initdata(ha, nElem);
    memset(htmp, 0, sizeof(float)*grid.x / 8);

    float *da, *dtmp;
    cudaMalloc((float**)&da, nbytes);
    cudaMalloc((float**)&dtmp, sizeof(float)*grid.x / 8);
    cudaMemcpy(da, ha, nbytes, cudaMemcpyHostToDevice);
    cudaMemset(dtmp, 0, sizeof(float)*grid.x / 8);

    float res = 0.0;
    for(int i=0; i<nElem; i++)
        res += ha[i];
    printf("correct sum: %f \n", res);

    sumInterleaveUnrollWarp<<<grid.x / 8, block>>>(da, dtmp, nElem);

    cudaMemcpy(htmp, dtmp, sizeof(float)*grid.x / 8, cudaMemcpyDeviceToHost);
    float res_d = 0.0;
    for(int i=0; i<grid.x / 8; i++)
        res_d += htmp[i];
    printf("gpu sum: %f \n", res_d);


    cudaFree(da);
    cudaFree(dtmp);
    free(ha);
    free(htmp);
}

9. 完全展开的归约

#include <stdio.h>
void initdata(float *a, int n)
{
    for(int i=0; i<n; i++)
    {
        a[i] = (float)(rand()&0xFF) / 100.f;
    }
}


__global__ void sumCompleteUnroll(float *d, float *tmp, int n)
{
    unsigned int tid = threadIdx.x;
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x * 8;
    if(idx + 7 * blockDim.x<n)
    {
        float a1 = d[idx];
        float a2 = d[idx + blockDim.x];
        float a3 = d[idx + blockDim.x * 2];
        float a4 = d[idx + blockDim.x * 3];
        float a5 = d[idx + blockDim.x * 4];
        float a6 = d[idx + blockDim.x * 5];
        float a7 = d[idx + blockDim.x * 6];
        float a8 = d[idx + blockDim.x * 7];
        d[idx] = a1 + a2 + a3 + a4 + a5 + a6 + a7 + a8;
    }
    __syncthreads();
    float *idata = d + 8 * blockDim.x * blockIdx.x;

    if(blockDim.x>=1024 && tid <512) idata[tid] += idata[tid + 512];
    __syncthreads();
    if(blockDim.x>=512 && tid <256) idata[tid] += idata[tid + 256];
    __syncthreads();
    if(blockDim.x>=256 && tid <128) idata[tid] += idata[tid + 128];
    __syncthreads();
    if(blockDim.x>=128 && tid <64) idata[tid] += idata[tid + 64];
    __syncthreads();
    if(tid<32)
    {
        volatile float *vmem = idata;
        vmem[tid] += vmem[tid + 32];
        vmem[tid] += vmem[tid + 16];
        vmem[tid] += vmem[tid + 8];
        vmem[tid] += vmem[tid + 4];
        vmem[tid] += vmem[tid + 2];
        vmem[tid] += vmem[tid + 1];
    }
    if(tid==0)tmp[blockIdx.x] = idata[0];
}
int main()
{
    int dev=0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("gpu: %d, name: %s \n", dev, deviceProp.name);
    cudaSetDevice(dev);

    int nElem = 1<<14;
    int nbytes = sizeof(float) * nElem;
    dim3 block(64);
    dim3 grid((nElem + block.x - 1) / block.x);

    float *ha, *htmp;
    ha = (float*)malloc(nbytes);
    htmp = (float*)malloc(sizeof(float)*grid.x / 8);
    initdata(ha, nElem);
    memset(htmp, 0, sizeof(float)*grid.x / 8);

    float *da, *dtmp;
    cudaMalloc((float**)&da, nbytes);
    cudaMalloc((float**)&dtmp, sizeof(float)*grid.x / 8);
    cudaMemcpy(da, ha, nbytes, cudaMemcpyHostToDevice);
    cudaMemset(dtmp, 0, sizeof(float)*grid.x / 8);

    float res = 0.0;
    for(int i=0; i<nElem; i++)
        res += ha[i];
    printf("correct sum: %f \n", res);


    sumCompleteUnroll<<<grid.x / 8, block>>>(da, dtmp, nElem);

    cudaMemcpy(htmp, dtmp, sizeof(float)*grid.x / 8, cudaMemcpyDeviceToHost);
    float res_d = 0.0;
    for(int i=0; i<grid.x / 8; i++)
        res_d += htmp[i];
    printf("gpu sum: %f \n", res_d);


    cudaFree(da);
    cudaFree(dtmp);
    free(ha);
    free(htmp);
}

10.模板函数的归约

blockDim.x 自内核启动时一旦确定，就不能更改了，所以模板函数解决了这个问题，当编译时编译器会去检查，blockDim.x 是否固定，如果固定，直接可以删除掉内核中不可能的部分也就是上半部分，下半部分是要执行的。

#include <stdio.h>
void initdata(float *a, int n)
{
    for(int i=0; i<n; i++)
    {
        a[i] = (float)(rand()&0xFF) / 100.f;
    }
}


template <unsigned int iBlockSize>
__global__ void sumCompleteUnrolltemplate(float *d, float *tmp, int n)
{
    unsigned int tid = threadIdx.x;
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x * 8;
    if(idx + 7 * blockDim.x<n)
    {
        float a1 = d[idx];
        float a2 = d[idx + blockDim.x];
        float a3 = d[idx + blockDim.x * 2];
        float a4 = d[idx + blockDim.x * 3];
        float a5 = d[idx + blockDim.x * 4];
        float a6 = d[idx + blockDim.x * 5];
        float a7 = d[idx + blockDim.x * 6];
        float a8 = d[idx + blockDim.x * 7];
        d[idx] = a1 + a2 + a3 + a4 + a5 + a6 + a7 + a8;
    }
    __syncthreads();
    float *idata = d + 8 * blockDim.x * blockIdx.x;

    if(iBlockSize>=1024 && tid <512) idata[tid] += idata[tid + 512];
    __syncthreads();
    if(iBlockSize>=512 && tid <256) idata[tid] += idata[tid + 256];
    __syncthreads();
    if(iBlockSize>=256 && tid <128) idata[tid] += idata[tid + 128];
    __syncthreads();
    if(iBlockSize>=128 && tid <64) idata[tid] += idata[tid + 64];
    __syncthreads();
    if(tid<32)
    {
        volatile float *vmem = idata;
        vmem[tid] += vmem[tid + 32];
        vmem[tid] += vmem[tid + 16];
        vmem[tid] += vmem[tid + 8];
        vmem[tid] += vmem[tid + 4];
        vmem[tid] += vmem[tid + 2];
        vmem[tid] += vmem[tid + 1];
    }
    if(tid==0)tmp[blockIdx.x] = idata[0];
}
int main()
{
    int dev=0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("gpu: %d, name: %s \n", dev, deviceProp.name);
    cudaSetDevice(dev);

    int nElem = 1<<14;
    int nbytes = sizeof(float) * nElem;
    dim3 block(64);
    dim3 grid((nElem + block.x - 1) / block.x);

    float *ha, *htmp;
    ha = (float*)malloc(nbytes);
    htmp = (float*)malloc(sizeof(float)*grid.x / 8);
    initdata(ha, nElem);
    memset(htmp, 0, sizeof(float)*grid.x / 8);

    float *da, *dtmp;
    cudaMalloc((float**)&da, nbytes);
    cudaMalloc((float**)&dtmp, sizeof(float)*grid.x / 8);
    cudaMemcpy(da, ha, nbytes, cudaMemcpyHostToDevice);
    cudaMemset(dtmp, 0, sizeof(float)*grid.x / 8);

    float res = 0.0;
    for(int i=0; i<nElem; i++)
        res += ha[i];
    printf("correct sum: %f \n", res);


    sumCompleteUnrolltemplate<64><<<grid.x / 8, block>>>(da, dtmp, nElem);

    cudaMemcpy(htmp, dtmp, sizeof(float)*grid.x / 8, cudaMemcpyDeviceToHost);
    float res_d = 0.0;
    for(int i=0; i<grid.x / 8; i++)
        res_d += htmp[i];
    printf("gpu sum: %f \n", res_d);


    cudaFree(da);
    cudaFree(dtmp);
    free(ha);
    free(htmp);
}

11. 使用共享内存

共享内存，减少全局内存访问

#include <stdio.h>
void initdata(float *a, int n)
{
    for(int i=0; i<n; i++)
    {
        a[i] = (float)(rand()&0xFF) / 100.f;
    }
}

template <unsigned int iBlockSize>
__global__ void sumCompleteUnrolltemplateSMEM(float *d, float *tmp, int n)
{
    __shared__ float smem[iBlockSize];
    unsigned int tid = threadIdx.x;
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x * 8;
    float temsum = 0.0;
    if(idx + 7 * blockDim.x<n)
    {
        float a1 = d[idx];
        float a2 = d[idx + blockDim.x];
        float a3 = d[idx + blockDim.x * 2];
        float a4 = d[idx + blockDim.x * 3];
        float a5 = d[idx + blockDim.x * 4];
        float a6 = d[idx + blockDim.x * 5];
        float a7 = d[idx + blockDim.x * 6];
        float a8 = d[idx + blockDim.x * 7];
        temsum = a1 + a2 + a3 + a4 + a5 + a6 + a7 + a8;
    }
    smem[tid] = temsum;
    
    __syncthreads();

    if(iBlockSize>=1024 && tid <512) smem[tid] += smem[tid + 512];
    __syncthreads();
    if(iBlockSize>=512 && tid <256) smem[tid] += smem[tid + 256];
    __syncthreads();
    if(iBlockSize>=256 && tid <128) smem[tid] += smem[tid + 128];
    __syncthreads();
    if(iBlockSize>=128 && tid <64) smem[tid] += smem[tid + 64];
    __syncthreads();
    if(tid<32)
    {
        volatile float *vmem = smem;
        vmem[tid] += vmem[tid + 32];
        vmem[tid] += vmem[tid + 16];
        vmem[tid] += vmem[tid + 8];
        vmem[tid] += vmem[tid + 4];
        vmem[tid] += vmem[tid + 2];
        vmem[tid] += vmem[tid + 1];
    }
    if(tid==0)tmp[blockIdx.x] = smem[0];
}

int main()
{
    int dev=0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("gpu: %d, name: %s \n", dev, deviceProp.name);
    cudaSetDevice(dev);

    int nElem = 1<<14;
    int nbytes = sizeof(float) * nElem;
    dim3 block(64);
    dim3 grid((nElem + block.x - 1) / block.x);

    float *ha, *htmp;
    ha = (float*)malloc(nbytes);
    htmp = (float*)malloc(sizeof(float)*grid.x / 8);
    initdata(ha, nElem);
    memset(htmp, 0, sizeof(float)*grid.x / 8);

    float *da, *dtmp;
    cudaMalloc((float**)&da, nbytes);
    cudaMalloc((float**)&dtmp, sizeof(float)*grid.x / 8);
    cudaMemcpy(da, ha, nbytes, cudaMemcpyHostToDevice);
    cudaMemset(dtmp, 0, sizeof(float)*grid.x / 8);

    float res = 0.0;
    for(int i=0; i<nElem; i++)
        res += ha[i];
    printf("correct sum: %f \n", res);

    sumCompleteUnrolltemplateSMEM<64><<<grid.x / 8, block>>>(da, dtmp, nElem);

    cudaMemcpy(htmp, dtmp, sizeof(float)*grid.x / 8, cudaMemcpyDeviceToHost);
    float res_d = 0.0;
    for(int i=0; i<grid.x / 8; i++)
        res_d += htmp[i];
    printf("gpu sum: %f \n", res_d);


    cudaFree(da);
    cudaFree(dtmp);
    free(ha);
    free(htmp);
}
我想静静，
关注
  • 0
    点赞
  • 踩
  • 5
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
cuda:英伟达 GPU 计算
06-29
CUDA编程 本文档的目的只是为了让新开发人员更容易理解整体 CUDA 范式和 NVIDIA GPU 功能。 本文档只是对 CUDA 功能和 CUDA 编程模型的高级概述。 在安装 CUDA Toolkit 之前阅读它可能更简单，阅读 CUDA C 编程指南、CUDA 运行时 API 和代码示例。 有关 CUDA 工具包和 CUDA 运行时 API 和 CUDA 库的完整而深入的概述，请访问 NVIDIA 网站： CUDA区： 加入 CUDA 注册开发者计划： 关于 CUDA 工具包的信息： 下载最新的 CUDA 工具包（Mac OSX / Linux x86 / Linux ARM / Windows）： 目录 记忆 主机内存 设备内存 统一内存 异步并发执行 [并发数据访问](#并发数据访问) 线程同步 原子函数 主机和设备之间的并发执行 并发数据传输 并发内核执行 数
巧用可变参数宏、编译器内置宏和printf输出调试信息
热门推荐
大道至简，持之以恒
06-15 1万+
前言：我们在写程序的时候，总是或多或少会加入一些printf之类的语句用于输出调试信息，但是printf语句有个很不方便的地方就是当我们需要发布程序的时候要一条一条的把这些语句删除，而一旦需要再次调试的时候，这些语句又不得不一条条的加上，这给我们带来了很大的不便，浪费了我们很多的时间，也造成了调试的效率低下。所以，很多人会选择使用宏定义的方式来输出调试语句。比如，定义一个宏开关：#define _
C语言琐碎记录之#define宏的用法【顺便提提 __FILE__,__LINE__的用法】
weixin_34138377的博客
03-05 191
C语言的#define 简而言之 就是 替换你的代码为目标代码的。 最近我有在用，比如 #define PRINT_LINE(FP,MSG) fprintf(FP,"%s:%d %s\n",__FILE__,__LINE__,MSG)#define free_safe(EXP) if((EXP)!=NULL &amp;&amp; \ (...
Cuda-Marketplace-Example:如何使用Nvidia CUDA解决简单的“市场问题”的示例
05-26
CUDA市场 项目资料 客观的： 开始使用CUDA进行编程 项目介绍： 基于具有多个购买选项的一系列产品，为这些产品中的每一个确定最佳购买选项。 “最佳购买选择权”标准仅基于选择权的价格。 资料库： 作者： 主题： 最后期限： 2015年6月23日 课程： 2014-2015 学院： 大学： 介绍 领域 在此练习中，我们假设我们是在线购物门户网站的管理者，最近该门户网站收到的服务器请求量显着增加。 由于此增加，平台的性能受到了影响，我们必须优化应用程序的某些部分，以便继续在我们的网站上提供顺畅的导航。 我们的销售门户是“市场”类型的。 通过这些，我们整合了从各个关联商店收到的目录。 值得强调的事实是，我们在不同商店的目录之间有很大的重叠部分，因此，同一商品具有多个购买选项，每个选项都有其相关的价格。 最终用户是选择要与哪个会员购买每种产品的人。 这是您的购物车配置为能够下
德国耶拿大学Cuda编程教材（英文版）
02-09
德国耶拿大学Cuda编程教材（英文版） 包含13课的讲课投影 6套练习题 指导教材 由浅入深介绍了cuda编程和优化以及并行算法
GPU编程 CUDA C++ 运行时错误检测Debug专用 宏函数 CHECK
行子的博客
03-31 424
可以捕捉核函数 add() 调用时（或配置时）的错误！
【CUDA基础练习】向量内积计算的若干种方法
12-18 1694
【显卡】 GeForce GTX 1650 显卡 | NVIDIAGeForce GTX 1650 显卡拥有突破性的图形性能，能为热门游戏带来加速的游戏体验。https://www.nvidia.cn/geforce/graphics-cards/gtx-1650/ Device0:"GeForce GTX 1650" Total amount of global memory 4101898240 bytes Number of mltiprocessors ..
NVIDIA CUDA 高度并行处理器编程（三）：CUDA存储器习题
ikun的博仔的博客
05-21 687
NVIDIA CUDA 高度并行处理器编程（三）：CUDA存储器习题 习题一中第一题，矩阵加法。可以使用共享存储器减少全局存储器对宽带的消耗吗？ 答：不可以，在每个线程计算一个的 kernel 函数中，每个线程只访问两个所需元素，且每个元素只被访问并加载一次。因为没有元素的重复访问，所以不能使用共享存储器减少全局存储器对宽带的消耗。 对于分块矩阵乘法，证明全局存储器宽带的减少正比于块的维度大小。 证明：不用分块的每个元素要加载 n 次，假设块维度 i，则每个块需要加载 n/i + 1 次。得证。
CUDA面试题库（含答案）一
tarzanlee666的博客
07-24 723
涵盖：CUDA历史概述、GPU和CPU的区别以及各自的优点和应用、并行计算模型：SIMD、MIMD、SPMD等、NVIDIA的GPU硬件架构，包括流多处理器（SM）、线程处理核心（CUDA cores）、全局内存、共享内存等、线程、线程块、网格以及它们在GPU上的执行方式。3、请详细描述NVIDIA的GPU硬件架构，包括流多处理器（SM）、线程处理核心（CUDA cores）、全局内存、共享内存等，并解释其在并行计算中的角色和作用。10、在NVIDIA的GPU架构中，流处理器（SM）的主要功能是什么？
《CUDA C编程权威指南》——1.6节习题
weixin_33851429的博客
05-02 342
本节书摘来自华章社区《CUDA C编程权威指南》一书中的第1章，第1.6节习题，作者[美]　马克斯·格罗斯曼（Max Grossman）　，更多章节内容可以访问云栖社区“华章社区”公众号查看 1.6　习题1.参考图1-5，分析以下几种数据划分形式：（1）对于二维数据，沿x轴进行块划分（2）对于二维数据，沿y轴进行周期划分（3）对于三维数据，沿z轴进行周期...
CUDA C/C++ 新手入门代码练习
01-11
CUDA c 编程学习入门，新手编代码必学
cuda编程讲义
09-22
cuda 2011最新讲义 共7个讲义， 有讲解及练习等
cuda 权威指南习题答案及coda
07-20
cuda 权威指南习题答案pdf，以及在Linux code！！！！
#define _CLNEW new(__FILE__, __LINE__)记录使用了new的文件名及行号,调试用
sealbird的专栏
12-06 1128
//Macro for creating new objects #if defined(LUCENE_ENABLE_MEMLEAKTRACKING) #define _CLNEW new(__FILE__, __LINE__) #define LUCENE_BASE public CL_NS(debug)::LuceneBase #elif defined(LUCENE_EN...
CUDA编程实战——第一个程序
Cross
06-29 676
#include &lt;iostream&gt; #include "book.h" using namespace std; //__global__修饰符告诉编译器这个函数应编译为在设备上运行（而不是主机） __global__ void add(int a, int b, int *c) { *c = a + b; } int main(){ int c; int *de...
CUDA By Example(三)——CUDA C并行编程
hhhcbw的博客，欢迎各位来访
11-09 1155
CUDA C并行编程，主要通过矢量求和与Julia集的绘制进行说明。以及对编译运行代码过程中遇到的问题及解决方案进行说明
NVIDIA CUDA 高度并行处理器编程（四）：性能优化习题
ikun的博仔的博客
07-18 412
优化后的归约与矩阵乘法算法，以及分块矩阵的内存连续访问证明
读书笔记 - CUDA编程指南5.0 + 练习编译运行 01
参谋谋的博客
06-19 947
前言：会根据CUDA编程指南一点点更新，欢迎讨论。图形处理的需求，推动可编程图形处理器（GPU）向着高并行度和多线程演化。从图形处理走向通用并行计算。GPU和CPU浮点计算能力的差异: GPU并行度高，晶体管更多用于数据处理而非缓存和流控。数据并行处理：将数据映射到并行处理的线程上，并进行加速。CUDA：一种通用并行计算架构CUDA核心的三个重点抽象（能提供犀利的的数据、任务并行）：线程组层次、共享存储器和栅栏同步threadIdx是一个有三个分量的向量(1d,2d,3d) 以上为两段代码,分别是一维向量和
解决 nvprof 报错ERR_NVGPUCTRPERM - The user does not have permission to profile on the target device
weixin_42764932的博客
03-18 2385
做 CUDA 开发，用 nvprof 命令行分析性能时报错： ==14973== Warning: ERR_NVGPUCTRPERM - The user does not have permission to profile on the target device. See the following link for instructions to enable permissions and get more information: https://developer.nvidia.com/ERR
windows cuda编程
最新发布
09-11
在Windows系统下，可以使用Visual Studio(VS)或者QTCreator来实现CUDA编程。对于QTCreator实现CUDA编程，可以按照以下步骤进行操作： 1. 在QTCreator中创建一个项目，生成一个模板kernel.cu文件。 2. 在左侧的解决...

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值