《CUDA By Example》【Chapter 10】流？

最新推荐文章于 2023-02-10 17:17:19 发布
最新推荐文章于 2023-02-10 17:17:19 发布
阅读量286 收藏 1
点赞数
CC 4.0 BY-SA版权
分类专栏： AI GPU
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.youkuaiyun.com/wondervictor/article/details/80612978
本文介绍CUDA流和页锁定内存的概念，展示了如何利用CUDA流实现GPU的任务并行，提高应用程序的执行效率。通过具体代码示例，详细解释了单个及多个CUDA流的应用方法。

摘要生成于 C知道 ，由 DeepSeek-R1 满血版支持， 前往体验 >

10.1 概述

并行性分数据并行性和任务并行性。GPU是SIMD天生的数据并行，同时也能支持类似于CPU多线程应用程序中的任务并行性。
GPU任务并行性不像CPU上那样灵活，但仍然可以进一步提高GPU上的运行速度。本章，介绍CUDA流，以及如何通过流在GPU上同时执行多个任务。

  • 了解如何分配也锁定（Page-Locked）类型的主机内存
  • 了解CUDA流的概念
  • 了解如何使用CUDA流来加速应用程序

10.2 页锁定主机内存

分配内存，有cudaMalloc（）在GPU上分配内存；用malloc（）在主机上分配内存。cudaHostAlloc（）是CDA独有的机制来分配主机内存。

cudaHostAlloc()和malloc()都是分配主机上内存，有什么区别？

  • malloc() ：分配标准的，可分页的（Pagable）主机内存

  • cudaHostAlloc() ：分配也锁定的主机内存。也锁定内存也成为固定内存（Pinned Memory）或者不可分页内存，有个重要属性，即操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终贮存在物理内存中。因此，操作系统能够安全的使某个应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位。

    因为GPU知道内存的物理地址，因此可以通过DMA来实现GPU和主机之间复制数据。在DMA复制过程中使用固定内存是非常重要的。固定内存虽可以提高DMA效率，但是也会失去虚拟内存的所有功能，此时会更快地耗尽内存。
    建议，进队cudaMemcpy（）调用中的源内存和目标内存，才是用页锁定内存，并且在不需要使用时立即释放，而非等到应用程序关闭时才释放。

    // allocate host locked memory, used to stream
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_a,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );

    // copy the locked memory to the device, async
    HANDLE_ERROR( cudaMemcpyAsync( dev_a, host_a+i,
                               N * sizeof(int),
                               cudaMemcpyHostToDevice,
                               stream ) );

   // copy the data from device to locked memory
   HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c,
                               N * sizeof(int),
                               cudaMemcpyDeviceToHost,
                               stream ) );

    // cleanup the streams and memory
    HANDLE_ERROR( cudaFreeHost( host_a ) );

10.3 CUDA流

HANDLE_ERROR( cudaEventRecord( start, 0 ) );
第二个参数，用于指定插入时间的流（Stream）

CUDA流在加速应用程序方面有重要作用。

10.4 使用单个CUDA流

basic_single_stream.cu

#include "../common/book.h"

#define N   (1024*1024)
#define FULL_DATA_SIZE   (N*20)


__global__ void kernel( int *a, int *b, int *c ) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        int idx1 = (idx + 1) % 256;
        int idx2 = (idx + 2) % 256;
        float   as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
        float   bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
        c[idx] = (as + bs) / 2;
    }
}


int main( void ) {
    cudaDeviceProp  prop;
    int whichDevice;
    HANDLE_ERROR( cudaGetDevice( &whichDevice ) );
    HANDLE_ERROR( cudaGetDeviceProperties( &prop, whichDevice ) );
    if (!prop.deviceOverlap) {
        printf( "Device will not handle overlaps, so no speed up from streams\n" );
        return 0;
    }

    cudaEvent_t     start, stop;
    float           elapsedTime;

    cudaStream_t    stream;
    int *host_a, *host_b, *host_c;
    int *dev_a, *dev_b, *dev_c;

    // start the timers
    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );

    // initialize the stream
    HANDLE_ERROR( cudaStreamCreate( &stream ) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_c,
                              N * sizeof(int) ) );

    // allocate host locked memory, used to stream
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_a,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_b,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_c,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );

    for (int i=0; i<FULL_DATA_SIZE; i++) {
        host_a[i] = rand();
        host_b[i] = rand();
    }

    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
    // now loop over full data, in bite-sized chunks
    for (int i=0; i<FULL_DATA_SIZE; i+= N) {
        // copy the locked memory to the device, async
        HANDLE_ERROR( cudaMemcpyAsync( dev_a, host_a+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_b, host_b+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream ) );

        kernel<<<N/256,256,0,stream>>>( dev_a, dev_b, dev_c );

        // copy the data from device to locked memory
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c,
                                       N * sizeof(int),
                                       cudaMemcpyDeviceToHost,
                                       stream ) );

    }
    // copy result chunk from locked to full buffer
    HANDLE_ERROR( cudaStreamSynchronize( stream ) );

    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );

    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
                                        start, stop ) );
    printf( "Time taken:  %3.1f ms\n", elapsedTime );

    // cleanup the streams and memory
    HANDLE_ERROR( cudaFreeHost( host_a ) );
    HANDLE_ERROR( cudaFreeHost( host_b ) );
    HANDLE_ERROR( cudaFreeHost( host_c ) );
    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaFree( dev_b ) );
    HANDLE_ERROR( cudaFree( dev_c ) );
    HANDLE_ERROR( cudaStreamDestroy( stream ) );

    return 0;
}

10.5 使用多个CUDA流

对个流，可以使GPU同时执行核函数和内存的复制。虽然可以加速，但是不需要核函数有任何改变。

basic_double_stream.cu

#include "../common/book.h"

#define N   (1024*1024)
#define FULL_DATA_SIZE   (N*20)


__global__ void kernel( int *a, int *b, int *c ) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        int idx1 = (idx + 1) % 256;
        int idx2 = (idx + 2) % 256;
        float   as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
        float   bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
        c[idx] = (as + bs) / 2;
    }
}


int main( void ) {
    cudaDeviceProp  prop;
    int whichDevice;
    HANDLE_ERROR( cudaGetDevice( &whichDevice ) );
    HANDLE_ERROR( cudaGetDeviceProperties( &prop, whichDevice ) );
    if (!prop.deviceOverlap) {
        printf( "Device will not handle overlaps, so no speed up from streams\n" );
        return 0;
    }

    cudaEvent_t     start, stop;
    float           elapsedTime;

    cudaStream_t    stream0, stream1;
    int *host_a, *host_b, *host_c;
    int *dev_a0, *dev_b0, *dev_c0;
    int *dev_a1, *dev_b1, *dev_c1;

    // start the timers
    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );

    // initialize the streams
    HANDLE_ERROR( cudaStreamCreate( &stream0 ) );
    HANDLE_ERROR( cudaStreamCreate( &stream1 ) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a0,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b0,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_c0,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a1,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b1,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_c1,
                              N * sizeof(int) ) );

    // allocate host locked memory, used to stream
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_a,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_b,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_c,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );

    for (int i=0; i<FULL_DATA_SIZE; i++) {
        host_a[i] = rand();
        host_b[i] = rand();
    }

    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
    // now loop over full data, in bite-sized chunks
    for (int i=0; i<FULL_DATA_SIZE; i+= N*2) {
        // copy the locked memory to the device, async
        HANDLE_ERROR( cudaMemcpyAsync( dev_a0, host_a+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream0 ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_b0, host_b+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream0 ) );

        kernel<<<N/256,256,0,stream0>>>( dev_a0, dev_b0, dev_c0 );

        // copy the data from device to locked memory
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c0,
                                       N * sizeof(int),
                                       cudaMemcpyDeviceToHost,
                                       stream0 ) );


        // copy the locked memory to the device, async
        HANDLE_ERROR( cudaMemcpyAsync( dev_a1, host_a+i+N,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream1 ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_b1, host_b+i+N,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream1 ) );

        kernel<<<N/256,256,0,stream1>>>( dev_a1, dev_b1, dev_c1 );

        // copy the data from device to locked memory
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i+N, dev_c1,
                                       N * sizeof(int),
                                       cudaMemcpyDeviceToHost,
                                       stream1 ) );
    }
    HANDLE_ERROR( cudaStreamSynchronize( stream0 ) );
    HANDLE_ERROR( cudaStreamSynchronize( stream1 ) );

    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );

    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
                                        start, stop ) );
    printf( "Time taken:  %3.1f ms\n", elapsedTime );

    // cleanup the streams and memory
    HANDLE_ERROR( cudaFreeHost( host_a ) );
    HANDLE_ERROR( cudaFreeHost( host_b ) );
    HANDLE_ERROR( cudaFreeHost( host_c ) );
    HANDLE_ERROR( cudaFree( dev_a0 ) );
    HANDLE_ERROR( cudaFree( dev_b0 ) );
    HANDLE_ERROR( cudaFree( dev_c0 ) );
    HANDLE_ERROR( cudaFree( dev_a1 ) );
    HANDLE_ERROR( cudaFree( dev_b1 ) );
    HANDLE_ERROR( cudaFree( dev_c1 ) );
    HANDLE_ERROR( cudaStreamDestroy( stream0 ) );
    HANDLE_ERROR( cudaStreamDestroy( stream1 ) );

    return 0;
}

10.6 GPU工作调度机制

硬件在处理内存复制和核函数执行时分别采用了不同的引擎。

10.7 高效地使用多个CUDA流

如果同时调度某个流的所有操作， 则容易在无意中阻塞另一个流的复制操作或者核函数执行。所以，解决办法是将操作放入流的队列时应采用宽度有限方式，而非深度有限方式。
###basic_double_stream_correct.cu

#include "../common/book.h"

#define N   (1024*1024)
#define FULL_DATA_SIZE   (N*20)


__global__ void kernel( int *a, int *b, int *c ) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        int idx1 = (idx + 1) % 256;
        int idx2 = (idx + 2) % 256;
        float   as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
        float   bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
        c[idx] = (as + bs) / 2;
    }
}


int main( void ) {
    cudaDeviceProp  prop;
    int whichDevice;
    HANDLE_ERROR( cudaGetDevice( &whichDevice ) );
    HANDLE_ERROR( cudaGetDeviceProperties( &prop, whichDevice ) );
    if (!prop.deviceOverlap) {
        printf( "Device will not handle overlaps, so no speed up from streams\n" );
        return 0;
    }

    cudaEvent_t     start, stop;
    float           elapsedTime;

    cudaStream_t    stream0, stream1;
    int *host_a, *host_b, *host_c;
    int *dev_a0, *dev_b0, *dev_c0;
    int *dev_a1, *dev_b1, *dev_c1;

    // start the timers
    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );

    // initialize the streams
    HANDLE_ERROR( cudaStreamCreate( &stream0 ) );
    HANDLE_ERROR( cudaStreamCreate( &stream1 ) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a0,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b0,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_c0,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a1,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b1,
                              N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_c1,
                              N * sizeof(int) ) );

    // allocate host locked memory, used to stream
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_a,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_b,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&host_c,
                              FULL_DATA_SIZE * sizeof(int),
                              cudaHostAllocDefault ) );

    for (int i=0; i<FULL_DATA_SIZE; i++) {
        host_a[i] = rand();
        host_b[i] = rand();
    }

    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
    // now loop over full data, in bite-sized chunks
    for (int i=0; i<FULL_DATA_SIZE; i+= N*2) {
        // enqueue copies of a in stream0 and stream1
        HANDLE_ERROR( cudaMemcpyAsync( dev_a0, host_a+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream0 ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_a1, host_a+i+N,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream1 ) );
        // enqueue copies of b in stream0 and stream1
        HANDLE_ERROR( cudaMemcpyAsync( dev_b0, host_b+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream0 ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_b1, host_b+i+N,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream1 ) );

        // enqueue kernels in stream0 and stream1   
        kernel<<<N/256,256,0,stream0>>>( dev_a0, dev_b0, dev_c0 );
        kernel<<<N/256,256,0,stream1>>>( dev_a1, dev_b1, dev_c1 );

        // enqueue copies of c from device to locked memory
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c0,
                                       N * sizeof(int),
                                       cudaMemcpyDeviceToHost,
                                       stream0 ) );
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i+N, dev_c1,
                                       N * sizeof(int),
                                       cudaMemcpyDeviceToHost,
                                       stream1 ) );
    }
    HANDLE_ERROR( cudaStreamSynchronize( stream0 ) );
    HANDLE_ERROR( cudaStreamSynchronize( stream1 ) );

    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );

    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
                                        start, stop ) );
    printf( "Time taken:  %3.1f ms\n", elapsedTime );

    // cleanup the streams and memory
    HANDLE_ERROR( cudaFreeHost( host_a ) );
    HANDLE_ERROR( cudaFreeHost( host_b ) );
    HANDLE_ERROR( cudaFreeHost( host_c ) );
    HANDLE_ERROR( cudaFree( dev_a0 ) );
    HANDLE_ERROR( cudaFree( dev_b0 ) );
    HANDLE_ERROR( cudaFree( dev_c0 ) );
    HANDLE_ERROR( cudaFree( dev_a1 ) );
    HANDLE_ERROR( cudaFree( dev_b1 ) );
    HANDLE_ERROR( cudaFree( dev_c1 ) );
    HANDLE_ERROR( cudaStreamDestroy( stream0 ) );
    HANDLE_ERROR( cudaStreamDestroy( stream1 ) );

    return 0;
}

疑问

  1. 异步内存复制的原理？
CUDA by Example（英文原书+自带源代码）
01-02
CUDA by Example.An Introduction to General-Purpose GPU Programming（英文原书+自带源代码）源代码是nvidia官网下的。
CUDA by example Chapter10 流
wolfcsharp的博客
09-05 290
流为什么要使用CUDA流页锁存主机内存CUDA流使用单个CUDA流使用多个CUDA流GPU工作调度机制高效地使用多个CUDA流 为什么要使用CUDA流 前面章节实现的GPU并行性都是把同一个任务分成可以并行执行的子任务，利用GPU中数量众多的cuda核来实现单个任务的并行性。而如果要实现任务并行性，就需要使用CUDA流来实现。 页锁存主机内存 1. 页锁存内存的概念及用途 页所存内存也称为固定内...
CUDA By Example(八)——流
hhhcbw的博客，欢迎各位来访
02-10 2588
CUDA 流操作，使用Stream加速程序执行
CUDA By Example（一）
2012到2014
05-27 250
        cudaMalloc是在设备上分配内存，第一个参数是新分配内存的地址，第二个参数是分配内存大小。在主机上不能对这块内存做任何的修改。主机指针只能访问主机代码中的内存，设备指针只能访问设备代码中的内存。使用完设备内存后要用cudaFree来释放掉分配的内存。在主机代码中可以调用cudaMemcpy来访问设备上的内存。 #include "cuda_runtime.h" #in...
CUDA by Example 源代码
gongchensz的专栏
11-25 1508
CUDA by Example ,源代码
cuda by example .h file and sample code of every chapter
04-05
"CUDA by Example .h file and sample code of every chapter" 提供了学习CUDA编程的一个宝贵资源，包含了书中每一章的头文件（.h）和示例代码。这些文件可以帮助读者深入理解CUDA编程的基本概念、API用法以及并行...
CUDA编程实践：《CUDA By Example》书中代码解析
《CUDA by Example》这本书是由Jason Sanders和Edward Kandrot合著，是一本面向希望利用NVIDIA CUDA技术进行GPU编程的开发者的入门书籍。 这本书通过大量的实例代码，帮助读者了解CUDA编程的核心概念和方法。它详细...
cuda_by_example实例代码
02-19
《CUDA by Example》是一本由Jason Katzer和Edward Kandrot合著的关于CUDA编程的书籍，这本书通过实例深入浅出地介绍了CUDA编程技术。CUDA是NVIDIA公司推出的计算平台，它允许程序员利用GPU（图形处理器）的强大计算...
CUDA By Example实践代码解析与应用
- "chapter03"、"chapter05"、"chapter09"、"chapter11" 和 "chapter04" 这些文件夹名暗示了包含在相应文件夹中的代码对应着书籍《CUDA By Example》中各章节的示例代码。每章可能专注于不同的CUDA编程技术或概念，...
CUDA by Example源码分析与实践指南
### CUDA编程入门与实践：cuda_by_example源码解析 CUDA（Compute Unified Device Architecture）是由NVIDIA推出的一种通用并行计算架构。它能够让开发者使用NVIDIA的GPU进行计算密集型任务的开发。《cuda_by_...
cuda by example中文版
02-06
cuda 经典的书 很好的入门教程 建议学习gpu的可以看看
《GPU高性能编程 CUDA实战》/《CUDA By Example》中的book.h
11-12
《GPU高性能编程 CUDA实战》/《CUDA By Example》中的book.h
CUDA By Examples 10 - Streams
Full_Speed_Turbo
07-12 706
CUDA Stream 学习总结.
《cuda by example》 book.h 头文件解析
fulva的专栏
11-26 6102
cpu_bitmap.h 头函数解析                blog.youkuaiyun.com/ghost129/article/details/4409565 都是一些函数，等用到的时候再过来解析，哈哈！ #ifndef __BOOK_H__ #define __BOOK_H__ #include /* cuda检错程序 */ static
CUDA之Stream介绍
silencez_w的博客
12-10 6567
一、Stream的概念 用到CUDA的程序一般需要处理海量的数据，内存带宽经常会成为主要的瓶颈。在Stream的帮助下，CUDA程序可以有效地将内存读取和数值运算并行，从而提升数据的吞吐量。 Cuda stream是指一堆异步的cuda操作，他们按照host代码调用的顺序执行在device上。 典型的cuda编程模式我们已经熟知了： · 将输入数据从host转移到device · 在device上执行kernel · 将结果从device上转移回host Cuda Streams 所有的cuda操作（包括k
CUDA By Example(一)——为什么需要CUDA
hhhcbw的博客，欢迎各位来访
11-07 1544
主要介绍了并行计算，为什么需要CUDA，以及CUDA架构是什么
Cuda10.1 +VS2019 samples 成功运行，但是新项目总是报错
weixin_40969653的博客
11-03 1565
我的samples 运行都没有问题。 但是新建的项目总是跑不起来，跑起来就说找不到 project.exe system can not find the file specified. 即使我把samples的代码复制到新项目中都还是同样的错误， 而且每个都会报错MSB３７２１. 好苦恼，有人能帮帮我吗 ...
CUDA by Example 源代码分析与应用实例
标题与描述提到的《CUDA by Example》是一本介绍CUDA编程的书籍，作者是Jason Sanders和Edward Kandrot。本书通过实例教学的方式，向读者展示了如何利用NVIDIA的CUDA编程模型来实现GPU加速计算。该书不仅适合对CUDA...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值