参加【Nvidia CUDA线上训练营2023春】笔记——如何优化CUDA程序

最新推荐文章于 2024-04-17 20:49:38 发布
最新推荐文章于 2024-04-17 20:49:38 发布
阅读量246 收藏
点赞数
分类专栏： 机器学习 文章标签： c++ c语言 大数据 深度学习 Powered by 金山文档
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/qq_43106579/article/details/128927436
版权

一、知识点梳理

GPU线程调度是一个黑盒，我们无法得知其中到底发生了什么，所以CUDA为每个系统函数提供了一个返回值，我们只需要定义专有变量接收函数返回值并对照文档对各个返回值的定义进行输出，就能快速获知到问题的发生的现场在哪。另外，CUDA提供"事件（event）"机制，通过简单的函数调用，就可以监测CUDA程序的时间性能。并且，CUDA可以利用独特的共享储存单元优化应用。

二、实验例程

1.错误检测

//错误码变量类型
cudaError_t error_code;
示例代码
//检测函数示例
#pragma once
#include <stdio.h>

#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)

在使用中，只需像如下这样就可以判断了。

CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));

就像这样

2.监测时间性能

CPU程序使用CPU自身的定时器来监测程序时间性能，而这种方式并不适用于GPU，但是GPU可以利用自身的时间戳来监测程序时间性能。

/*******Event相关的函数********/
//声明
cudaEvent_t event;
//创建
cudaError_t cudaEventCreate(cudaEvent_t* event);
//销毁
cudaError_t cudaEventDestroy(cudaEvent_t* event);
//添加事件到当前执行流
cudaError_t cudaEventRecord(cudaEvent_t event,cudaStream_t stream = 0);
//等待事件完成，设立flag
cudaError_t cudaEventSynchronize(cudaEvent_t event); //阻塞
cudaError_t cudaEventQuery (cudaEvent_t event); //非阻塞
//记录执行的事件
cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start,cudaEvent_t stop);
示例代码
#include <stdio.h>
#include <math.h>
#include "error.cuh"

#define BLOCK_SIZE 16

__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{ 
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int sum = 0;
    if( col < k && row < m) 
    {
        for(int i = 0; i < n; i++) 
        {
            sum += a[row * n + i] * b[i * k + col];
        }
        c[row * k + col] = sum;
    }
} 

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0.0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

int main(int argc, char const *argv[])
{
    int m=100;
    int n=100;
    int k=100;

    int *h_a, *h_b, *h_c, *h_cc;
    cudaMallocHost((void **) &h_a, sizeof(int)*m*n);
    cudaMallocHost((void **) &h_b, sizeof(int)*n*k);
    cudaMallocHost((void **) &h_c, sizeof(int)*m*k);
    cudaMallocHost((void **) &h_cc, sizeof(int)*m*k);
    
    cudaEvent_t start,stop; //声明event
    cudaEventCreate(&start);  //创建event
    cudaEventCreate(&stop);

    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = rand() % 1024;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = rand() % 1024;
        }
    }

    int *d_a, *d_b, *d_c;
    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
    cudaMalloc((void **) &d_b, sizeof(int)*n*k);
    cudaMalloc((void **) &d_c, sizeof(int)*m*k);

    // copy matrix A and B from host to device memory
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice);

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);

   //添加事件到当前执行流
   cudaEventRecord(start);
    gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); 
    cudaEventRecord(stop);
    //等待事件完成，设立flag
    cudaEventSynchronize(stop);
    //记录时间
    float timeout;
    cudaEventElapsedTime(&timeout,start,stop);
    printf("Timeout is %g ms,\n",timeout);
    //销毁event
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    cudaMemcpy(h_c, d_c, sizeof(int)*m*k, cudaMemcpyDeviceToHost);
    //cudaThreadSynchronize(&timeout,start,stop);

    cpu_matrix_mult(h_a, h_b, h_cc, m, n, k);

    int ok = 1;
    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cc[i*k + j] - h_c[i*k + j])>(1.0e-10))
            {
                
                ok = 0;
            }
        }
    }

    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }

    // free memory
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    cudaFreeHost(h_a);
    cudaFreeHost(h_b);
    cudaFreeHost(h_c);
    cudaFreeHost(h_cc);
    return 0;
}

  1. 利用共享储存单元优化应用

CUDA硬件划分了多种储存单元，按照读取速度由高到低分别为，Register file,Shared Memory,Constant Memory,Texture Memory,Local Memory and Global Memory。在这里我们用到Shared Memory，所以只着重说一下这个。

Shared Memory访问速度仅次于Register file，并且Shared Memory是最快的可以让多个线程沟通的地方，所以线程对Shared Memory的访问就很频繁，为了提升效率，CUDA将Shared Memory分为32个（逻辑块）banks。

同常量内存一样，当一个warp中的所有线程访问同一地址的共享内存是，会触发广播（broadcast）机制到warp中所有线程，这是最高效的，例如：

如果同一个half-warp/warp中的线程访问同一个bank中的不同地址是将发生bank conflict，例如：

我们要在代码的编写中避免bank conflict的产生。

//Shared Memory申请关键字
__shared__ int sm[SIZE][SIZE];

大家应该注意到，Shared Memory是有大小的，他的宽度一般为32bit，所以，可能不能覆盖到全部数据，所以，在实际计算中，Shared Memory，是会“滑动”的。例如：

在两个16x16的矩阵Ｍ，Ｎ中，分别设置一块8x8的Shared Memory，这两块共享内存在调用的时候会将16x16的矩阵分为4块，以“滑动”的方式分别计算。

使用共享内存会对程序的时间性能有一定的提升，下面举例说明：

//未使用共享内存的代码
#include <stdio.h>
#include <math.h>
#include "error.cuh"

#define BLOCK_SIZE 16

__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{ 
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int sum = 0;
    if( col < k && row < m) 
    {
        for(int i = 0; i < n; i++) 
        {
            sum += a[row * n + i] * b[i * k + col];
        }
        c[row * k + col] = sum;
    }
} 

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0.0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

int main(int argc, char const *argv[])
{
    int m=100;
    int n=100;
    int k=100;

    int *h_a, *h_b, *h_c, *h_cc;
    cudaMallocHost((void **) &h_a, sizeof(int)*m*n);
    cudaMallocHost((void **) &h_b, sizeof(int)*n*k);
    cudaMallocHost((void **) &h_c, sizeof(int)*m*k);
    cudaMallocHost((void **) &h_cc, sizeof(int)*m*k);
    
    cudaEvent_t start,stop; //声明event
    cudaEventCreate(&start);  //创建event
    cudaEventCreate(&stop);

    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = rand() % 1024;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = rand() % 1024;
        }
    }

    int *d_a, *d_b, *d_c;
    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
    cudaMalloc((void **) &d_b, sizeof(int)*n*k);
    cudaMalloc((void **) &d_c, sizeof(int)*m*k);

    // copy matrix A and B from host to device memory
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice);

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);

   //添加事件到当前执行流
   cudaEventRecord(start);
    gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); 
    cudaEventRecord(stop);
    //等待事件完成，设立flag
    cudaEventSynchronize(stop);
    //记录时间
    float timeout;
    cudaEventElapsedTime(&timeout,start,stop);
    printf("Timeout is %g ms,\n",timeout);
    //销毁event
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    cudaMemcpy(h_c, d_c, sizeof(int)*m*k, cudaMemcpyDeviceToHost);
    //cudaThreadSynchronize(&timeout,start,stop);

    cpu_matrix_mult(h_a, h_b, h_cc, m, n, k);

    int ok = 1;
    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cc[i*k + j] - h_c[i*k + j])>(1.0e-10))
            {
                
                ok = 0;
            }
        }
    }

    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }

    // free memory
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    cudaFreeHost(h_a);
    cudaFreeHost(h_b);
    cudaFreeHost(h_c);
    cudaFreeHost(h_cc);
    return 0;
}

未使用共享内存时，程序所用时间

//使用共享内存的代码
#include <stdio.h>
#include <math.h>
#include "error.cuh"

#define BLOCK_SIZE 16

__global__ void gpu_matrix_mult_shared(int *d_a, int *d_b, int *d_result, int m, int n, int k) 
{
    __shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE];

    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    int tmp = 0;
    int idx;

    for (int sub = 0; sub < gridDim.x; ++sub) 
    {
        idx = row * n + sub * BLOCK_SIZE + threadIdx.x;
        tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0;
        idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;
        tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0;
        
        __syncthreads();
        for (int k = 0; k < BLOCK_SIZE; ++k) 
        {
            tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x];
        }
        __syncthreads();
    }
    if(row < n && col < n)
    {
        d_result[row * n + col] = tmp;
    }
}

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0.0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

int main(int argc, char const *argv[])
{
    int m=100;
    int n=100;
    int k=100;

    int *h_a, *h_b, *h_c, *h_cc;
    cudaMallocHost((void **) &h_a, sizeof(int)*m*n);
    cudaMallocHost((void **) &h_b, sizeof(int)*n*k);
    cudaMallocHost((void **) &h_c, sizeof(int)*m*k);
    cudaMallocHost((void **) &h_cc, sizeof(int)*m*k);
    
    cudaEvent_t start,stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = rand() % 1024;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = rand() % 1024;
        }
    }

    int *d_a, *d_b, *d_c;
    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
    cudaMalloc((void **) &d_b, sizeof(int)*n*k);
    cudaMalloc((void **) &d_c, sizeof(int)*m*k);

    // copy matrix A and B from host to device memory
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice);

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
   
   cudaEventRecord(start);
    gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); 
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float timeout;
    cudaEventElapsedTime(&timeout,start,stop);
    printf("Timeout is %g ms,\n",timeout);
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    cudaMemcpy(h_c, d_c, sizeof(int)*m*k, cudaMemcpyDeviceToHost);
    //cudaThreadSynchronize(&timeout,start,stop);

    cpu_matrix_mult(h_a, h_b, h_cc, m, n, k);

    int ok = 1;
    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cc[i*k + j] - h_c[i*k + j])>(1.0e-10))
            {
                
                ok = 0;
            }
        }
    }

    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }

    // free memory
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    cudaFreeHost(h_a);
    cudaFreeHost(h_b);
    cudaFreeHost(h_c);
    cudaFreeHost(h_cc);
    return 0;
}

使用共享内存时，程序所用时间

可以看到，在这里使用共享内存使程序获得了大约 30%的提升，这是从GPU储存单元的角度对程序优化的手段。

三.总结

CUDA已经提供了很成熟的框架，在这个框架内，我们可以使用的工具很多，这也大大降低了代码编写的难度，时间就可以放在思路理解上。

✿柚汁儿
关注
  • 0
    点赞
  • 踩
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
CUDA程序优化
01-21
CUDA程序优参考资料，从指令基础讲解，比较实用，供大家参考使用！
CUDA学习笔记之程序优化
yujiflying的专栏
12-16 965
CUDA程序优化 CUDA程序优化应该考虑的点： 精度：只在关键步骤使用双精度，其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡；            延迟：需要首先缓冲一部分数据，缓冲的大小应该可以保证每个内核程序处理的一批数据能够让GPU慢负荷工作；            计算量：计算量太小的程序使用CUDA很不合算；当需要计算的问题的计算密集度很低的时候，执行计算的时间远远比
CUDA优化入门
最新发布
小裘的博客
04-17 666
通过优化矩阵乘法来学习CUDA
CUDA: CUDA程序优化步骤
923723914
10-12 247
在《runtime API创建CUDA程序》中，我们做了一个计算一大堆数字的平方和的程序。不过，我们也提到这个程序的执行效率并不理想。当然，实际上来说，如果只是要做计算平方和的动作，用CPU做会比用 GPU 快得多。这是因为平方和的计算并不需要太多运算能力，所以几乎都是被内存带宽所限制。因此，光是把数据复制到显卡内存上的这个动作，所需要的时间，可能已经和直接在CPU上进行计算差不多了。 　　不过...
CUDA程序的优化
wufenxia的专栏
10-21 393
CUDA程序的优化 CUDA程序员优化的目的，是以最短的时间，在允许的误差范围内完成给定的计算任务。在这里，“最短的时间”是指整个程序的运行时间，更侧重于计算的吞吐量，而不是单个数据的延迟。在开始考虑GUP和CPU协同计算之前，应该先粗略地评估使用CUDA是否能达到预想效果，包括以下几个方面： 1.精度 目前，GPU的单精度计算性能要远远超过双精度计算性能，整数乘法、除...
CUDA——性能优化（一）
01-20
CUDA全局内存的合并访问（个人理解） 每个warp去访问全局内存，会有400-600个时钟周期的内存延迟，这个代价很昂贵，所以为了减少访问全局内存的指令次数，我们将满足字节大小和对齐要求的warp合并起来访问全局内存，...
NVIDIA GPU CUDA代码性能优化基础
06-06
Fundamental Optimizations in CUDA Optimization Overview GPU architecture Kernel optimization — Memory optimization — Latency optimization — Instruction optimization CPU-GPU interaction ...
NVIDIA CUDA官方工具11.4.0版本win10系统在线安装包
07-23
CUDA 11.4.0是NVIDIA针对Windows 10系统提供的一个更新版本，旨在提升GPU计算效率，增强软件兼容性和性能优化。在这个版本中，用户可以期待一些关键的改进、新特性和修复。 首先，CUDA 11.4.0包含了一些重要的性能...
CUDA-Fortran-Book_nvidia_CUDA-Fortran_
10-04
《CUDA-Fortran-Book_nvidia_CUDA-Fortran_》是关于NVIDIA CUDA Fortran的一本详细教程，旨在帮助开发者理解和应用CUDA Fortran进行高性能计算。CUDA Fortran是NVIDIA公司为Fortran语言开发的一种扩展，它允许程序员...
ubuntu16.04 cuda10.0 pytorch安装（csdn）————程序.pdf
12-04
CUDA是NVIDIA开发的一个并行计算平台，用于加速高性能计算任务，而PyTorch则是一个开源的深度学习框架，它支持GPU计算，能够极大地提升训练神经网络的速度。 首先，确保你的系统已经安装了NVIDIA驱动和CUDA 10.0。...
cuda 优化步骤
08-04
cuda概述，cuda架构，cuda内核优化详细步骤说明
CUDA程序优化[转]
03-23 65
本文章转载自CSDN，原博客地址http://blog.csdn.net/ouczoe/article/details/5137063 CUDA程序优化 CUDA程序优化应该考虑的点： 精度：只在关键步骤使用双精度，其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡； 延迟：需要首先缓冲一部分数据，缓冲的大小应该可以保证每个内核程序处理的一批数据能够让G...
CUDA学习之CUDA程序优化
weixin_30425949的博客
06-25 89
一 测量程序运行时间 1主机端测时 由于CUDA API函数都是异步的，为了准确测量CUDA调用运行的时间，首先要使用cudaThreadSynchronize(),同步cpu与gpu之后，才能结束测时。 2设备端测时 使用clock()函数，这个函数测的结果是一个block在gpu中上下文保持的时间，而不是实际执行的时间。（一般大于执行时间） 二任务划分 1原则 划分为cp...
CUDA程序性能调优
yu132563的专栏
12-28 815
介绍了GPU的结构以及资源的控制要素（GPU硬件结构和程序具体参数设置_yu132563的专栏-CSDN博客），下面就可以对CUDA进行程序的调优，从而在不同的GPU上面运行同一个CUDA程序的参数设置方法。 对于一个CUDA kernel function而言，其通常由如下几个部分组成： kernel function paras local variables shared memory with __syncthreads__ call device function call loop/
CUDA程序优化方法
changyi9995的博客
04-14 871
使用共享内存减少全局内存读取次数 减少全局内存的重复数据的重复访问，此处大有学问，需要设计我们的线程组织模式，最大可能利用共享内存，可参考矩阵乘法优化问题； 把全局内存绑定为纹理； 纹理的存取速度要远高于全局内存 减少bank conflict, 让不同线程读取连续内存。 Tesla 的每个 SM 拥有 16KB 共享存储器，用于同一个线程块内的线程间通信。为了使一个 half-warp 内的线程能够在一个内核周期中并行访问，共享存储器被组织成 16 个 bank，每个 bank 拥有 32bi..
CUDA 程序的优化(1) 概述
不算太晚
09-24 901
CUDA程序优化的最终目的，是以最短的时间，在允许的误差范围内完成给定的计算任务。在这里，“最短的时间”是指整个程序的运行时间，更侧重于计算的吞吐量，而不是单个数据的延迟。在开始考虑使用GPU和CPU协同计算之前，应该先粗略地评估使用CUDA是否能达到预想的效果，包括以下几个方面: 1.精度 目前，GPU的单精度计算性能要远远超过双精度计算性能，整数乘法、除法、求模等运算的指令吞吐量也较为有限。在科学计算中，由于需要处理的数据量巨大，往往只有在采用双精度或者四精度时才能获得可靠的结果。目前，采用Tesla架
CUDA程序基本优化
sinat_38368658的博客
04-08 343
有效的数据并行算法 + 针对GPU架构特性的优化 = 最优性能 1.Parallel Reduction 并行规约 这个过程类似篮球锦标赛的淘汰过程：n个元素进行log(n)个回合，如何在CUDA上实现？ //累加存在shared memory内的元素 __shared__ float partialSum[element_num]; unsigned int t = thread...
CUDA程序优化小记（七）
开发游戏，辉煌全中国
01-26 3736
CUDA全称Computer Unified Device Architecture（计算机同一设备架构），它的引入为计算机计算速度质的提升提供了可能，从此微型计算机也能有与大型机相当计算的能力。可是不恰当地使用CUDA技术，不仅不会让应用程序获得提升，反而会比普通CPU的计算还要慢。最近我通过学习《GPGPU编程技术》这本书，深刻地体会到了这一点，并且用CUDA Runtime应用改写书上的例子程序；来体会CUDA技术给我们计算能力带来的提升。
CUDA详细参数（算力+CUDA支持版本）
08-19
CUDA详细参数（算力+CUDA支持版本）

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

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

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值