《GPU高性能编程CUDA实战》学习笔记（十一）

最新推荐文章于 2024-04-18 10:04:17 发布
最新推荐文章于 2024-04-18 10:04:17 发布
阅读量3.5k 收藏 4
点赞数 2
分类专栏： CUDA 文章标签： nvidia 高性能 cuda gpu c语言
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/w09103419/article/details/52574998
版权

第11章 多GPU系统上的CUDA C

“比在一个GPU上计算要更好的，只有在两个GPU上计算。”

11.1 本章目标

  • 了解如何分配和使用零拷贝内存（Zero-Copy Memory）；
  • 了解如何在同一个应用程序中使用多个GPU；
  • 了解如何分配和使用可移动的固定内存（Portable Pinned Memory）。

11.2 零拷贝主机内存

cudaHostAlloc() 传递参数 cudaHostAllocDefault 分配固定内存；cudaHostAllocMapped分配主机固定内存。
这种内存除了可以用于主机与GPU之间的内存复制外，还可以子CUDA C核函数中直接访问这种类型的主机内存。由于这种内存不需要复制到GPU，因此也称为零拷贝内存。

11.2.1 通过零拷贝内存实现点积运算

GPU访问主机内存：实现点积运算。GPU上执行归约运算，并使用零拷贝作为输入缓冲区和输出缓冲区。
点积运算的主机内存版本：malloc_test
零拷贝版本： cuda_host_alloc_test，cudaHostAllocMapped零拷贝内存标志；cudaHostAllocWriteCombined 运行时将内存分配为“合并式写入（Write-Combined）内存”
cudaHostGetDevicePointer 获取GPU指针 
#include "../common/book.h"

#define imin(a,b) (a<b?a:b)

const int N = 33 * 1024 * 1024;
const int threadsPerBlock = 256;
const int blocksPerGrid =
            imin( 32, (N+threadsPerBlock-1) / threadsPerBlock );


__global__ void dot( int size, float *a, float *b, float *c ) {
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;

    float   temp = 0;
    while (tid < size) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }
    
    // set the cache values
    cache[cacheIndex] = temp;
    
    // synchronize threads in this block
    __syncthreads();

    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0) {
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}


float malloc_test( int size ) {
    cudaEvent_t     start, stop;
    float           *a, *b, c, *partial_c;
    float           *dev_a, *dev_b, *dev_partial_c;
    float           elapsedTime;

    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );

    // allocate memory on the CPU side
    a = (float*)malloc( size*sizeof(float) );
    b = (float*)malloc( size*sizeof(float) );
    partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
                              size*sizeof(float) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b,
                              size*sizeof(float) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
                              blocksPerGrid*sizeof(float) ) );

    // fill in the host memory with data
    for (int i=0; i<size; i++) {
        a[i] = i;
        b[i] = i*2;
    }

    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
    // copy the arrays 'a' and 'b' to the GPU
    HANDLE_ERROR( cudaMemcpy( dev_a, a, size*sizeof(float),
                              cudaMemcpyHostToDevice ) );
    HANDLE_ERROR( cudaMemcpy( dev_b, b, size*sizeof(float),
                              cudaMemcpyHostToDevice ) ); 

    dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
                                            dev_partial_c );
    // copy the array 'c' back from the GPU to the CPU
    HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
                              blocksPerGrid*sizeof(float),
                              cudaMemcpyDeviceToHost ) );

    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
                                        start, stop ) );

    // finish up on the CPU side
    c = 0;
    for (int i=0; i<blocksPerGrid; i++) {
        c += partial_c[i];
    }

    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaFree( dev_b ) );
    HANDLE_ERROR( cudaFree( dev_partial_c ) );

    // free memory on the CPU side
    free( a );
    free( b );
    free( partial_c );

    // free events
    HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );

    printf( "Value calculated:  %f\n", c );

    return elapsedTime;
}


float cuda_host_alloc_test( int size ) {
    cudaEvent_t     start, stop;
    float           *a, *b, c, *partial_c;
    float           *dev_a, *dev_b, *dev_partial_c;
    float           elapsedTime;

    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );

    // allocate the memory on the CPU
    HANDLE_ERROR( cudaHostAlloc( (void**)&a,
                              size*sizeof(float),
                              cudaHostAllocWriteCombined |
                                     cudaHostAllocMapped ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&b,
                              size*sizeof(float),
                              cudaHostAllocWriteCombined |
                                     cudaHostAllocMapped ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&partial_c,
                              blocksPerGrid*sizeof(float),
                              cudaHostAllocMapped ) );

    // find out the GPU pointers
    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_a, a, 0 ) );
    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_b, b, 0 ) );
    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_partial_c,
                                            partial_c, 0 ) );

    // fill in the host memory with data
    for (int i=0; i<size; i++) {
        a[i] = i;
        b[i] = i*2;
    }

    HANDLE_ERROR( cudaEventRecord( start, 0 ) );

    dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
                                            dev_partial_c );

    HANDLE_ERROR( cudaThreadSynchronize() );
    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
                                        start, stop ) );

    // finish up on the CPU side
    c = 0;
    for (int i=0; i<blocksPerGrid; i++) {
        c += partial_c[i];
    }

    HANDLE_ERROR( cudaFreeHost( a ) );
    HANDLE_ERROR( cudaFreeHost( b ) );
    HANDLE_ERROR( cudaFreeHost( partial_c ) );

    // free events
    HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );

    printf( "Value calculated:  %f\n", c );

    return elapsedTime;
}


int main( void ) {
    cudaDeviceProp  prop;
    int whichDevice;
    HANDLE_ERROR( cudaGetDevice( &whichDevice ) );
    HANDLE_ERROR( cudaGetDeviceProperties( &prop, whichDevice ) );
    if (prop.canMapHostMemory != 1) {
        printf( "Device can not map memory.\n" );
        return 0;
    }

    float           elapsedTime;

    HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );

    // try it with malloc
    elapsedTime = malloc_test( N );
    printf( "Time using cudaMalloc:  %3.1f ms\n",
            elapsedTime );

    // now try it with cudaHostAlloc
    elapsedTime = cuda_host_alloc_test( N );
    printf( "Time using cudaHostAlloc:  %3.1f ms\n",
            elapsedTime );
}

11.2.2 零拷贝内存的性能

独立GPU自己拥有专门的DRAM，通常位于CPU之外的电路板上。位于PCIE槽中的扩展卡。
集成GPU是系统芯片组中内置的图形处理器，通常与CPU共享系统内存。

对于集成GPU，使用零拷贝通常有性能提升，因为内存在物理上与主机是共享的，可以避免数据复制， 但它占用系统内存，所以系统性能会降低。
当输入内存和输出内存都只能使用一次时，那么独立GPU上使用零拷贝内存性能会提升，因为GPU在设计时考虑了隐藏内存访问带来的延迟，因此它将减轻PCIE总线上的读写等操作的延迟，从而提升性能， 但不会缓存零拷贝内存的内容，如果多次读取内存，将得不偿失，还不如一开始就将数据复制到GPU。
查看GPU是集成还是独立：方法1，打开机箱来观察；方法2，通过代码cudaGetDeviceProperties() 返回的结构来判断属性integrated，如果是集成GPU，则这个值是true，否则为false。

11.3 使用多个GPU

NVIDIA的SLI技术将它们桥接。
我们将改点积运算为多个GPU， 
#include "../common/book.h"


#define imin(a,b) (a<b?a:b)

#define     N    (33*1024*1024)
const int threadsPerBlock = 256;
const int blocksPerGrid =
            imin( 32, (N/2+threadsPerBlock-1) / threadsPerBlock );


__global__ void dot( int size, float *a, float *b, float *c ) {
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;

    float   temp = 0;
    while (tid < size) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }
    
    // set the cache values
    cache[cacheIndex] = temp;
    
    // synchronize threads in this block
    __syncthreads();

    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0) {
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}


struct DataStruct {
    int     deviceID;
    int     size;
    float   *a;
    float   *b;
    float   returnValue;
};


void* routine( void *pvoidData ) {
    DataStruct  *data = (DataStruct*)pvoidData;
    HANDLE_ERROR( cudaSetDevice( data->deviceID ) );

    int     size = data->size;
    float   *a, *b, c, *partial_c;
    float   *dev_a, *dev_b, *dev_partial_c;

    // allocate memory on the CPU side
    a = data->a;
    b = data->b;
    partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
                              size*sizeof(float) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b,
                              size*sizeof(float) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
                              blocksPerGrid*sizeof(float) ) );

    // copy the arrays 'a' and 'b' to the GPU
    HANDLE_ERROR( cudaMemcpy( dev_a, a, size*sizeof(float),
                              cudaMemcpyHostToDevice ) );
    HANDLE_ERROR( cudaMemcpy( dev_b, b, size*sizeof(float),
                              cudaMemcpyHostToDevice ) ); 

    dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
                                            dev_partial_c );
    // copy the array 'c' back from the GPU to the CPU
    HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
                              blocksPerGrid*sizeof(float),
                              cudaMemcpyDeviceToHost ) );

    // finish up on the CPU side
    c = 0;
    for (int i=0; i<blocksPerGrid; i++) {
        c += partial_c[i];
    }

    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaFree( dev_b ) );
    HANDLE_ERROR( cudaFree( dev_partial_c ) );

    // free memory on the CPU side
    free( partial_c );

    data->returnValue = c;
    return 0;
}


int main( void ) {
    int deviceCount;
    HANDLE_ERROR( cudaGetDeviceCount( &deviceCount ) );
    if (deviceCount < 2) {
        printf( "We need at least two compute 1.0 or greater "
                "devices, but only found %d\n", deviceCount );
        return 0;
    }

    float   *a = (float*)malloc( sizeof(float) * N );
    HANDLE_NULL( a );
    float   *b = (float*)malloc( sizeof(float) * N );
    HANDLE_NULL( b );

    // fill in the host memory with data
    for (int i=0; i<N; i++) {
        a[i] = i;
        b[i] = i*2;
    }

    // prepare for multithread
    DataStruct  data[2];
    data[0].deviceID = 0;
    data[0].size = N/2;
    data[0].a = a;
    data[0].b = b;

    data[1].deviceID = 1;
    data[1].size = N/2;
    data[1].a = a + N/2;
    data[1].b = b + N/2;

    CUTThread   thread = start_thread( routine, &(data[0]) );
    routine( &(data[1]) );
    end_thread( thread );


    // free memory on the CPU side
    free( a );
    free( b );

    printf( "Value calculated:  %f\n",
            data[0].returnValue + data[1].returnValue );

    return 0;
}
每个GPU都需要由一个不同的CPU线程来控制，start_thread创建了新的线程【这个函数是自己写的哟】。

11.4 可移动的固定内存

固定内存只是相对于分配他的线程是固定内存，其他线程让然当它是正常内存可以拷贝、各种操作。
一种补救措施：我们将固定内存分配为可移动的，就是可以在主机线程之间移动这块内存，并且每个线程都视其为固定内存。
需要使用cudaHostAlloc() 来分配内存，并且调用新的标志 cudaHostAllocPortable，这个标志可以与其他标志一起使用，如cudaHostAllocWriteCombined和cudaHostAllocMapped。这意味着分配主机内存时，可将其视为可移动、零拷贝以及合并式写入等的任意组合。

例子，点积运算 
#include "../common/book.h"


#define imin(a,b) (a<b?a:b)

#define     N    (33*1024*1024)
const int threadsPerBlock = 256;
const int blocksPerGrid =
            imin( 32, (N/2+threadsPerBlock-1) / threadsPerBlock );


__global__ void dot( int size, float *a, float *b, float *c ) {
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;

    float   temp = 0;
    while (tid < size) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }
    
    // set the cache values
    cache[cacheIndex] = temp;
    
    // synchronize threads in this block
    __syncthreads();

    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0) {
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}


struct DataStruct {
    int     deviceID;
    int     size;
    int     offset;
    float   *a;
    float   *b;
    float   returnValue;
};


void* routine( void *pvoidData ) {
    DataStruct  *data = (DataStruct*)pvoidData;
    if (data->deviceID != 0) {
        HANDLE_ERROR( cudaSetDevice( data->deviceID ) );
        HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );
    }

    int     size = data->size;
    float   *a, *b, c, *partial_c;
    float   *dev_a, *dev_b, *dev_partial_c;

    // allocate memory on the CPU side
    a = data->a;
    b = data->b;
    partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_a, a, 0 ) );
    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_b, b, 0 ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
                              blocksPerGrid*sizeof(float) ) );

    // offset 'a' and 'b' to where this GPU is gets it data
    dev_a += data->offset;
    dev_b += data->offset;

    dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
                                            dev_partial_c );
    // copy the array 'c' back from the GPU to the CPU
    HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
                              blocksPerGrid*sizeof(float),
                              cudaMemcpyDeviceToHost ) );

    // finish up on the CPU side
    c = 0;
    for (int i=0; i<blocksPerGrid; i++) {
        c += partial_c[i];
    }

    HANDLE_ERROR( cudaFree( dev_partial_c ) );

    // free memory on the CPU side
    free( partial_c );

    data->returnValue = c;
    return 0;
}


int main( void ) {
    int deviceCount;
    HANDLE_ERROR( cudaGetDeviceCount( &deviceCount ) );
    if (deviceCount < 2) {
        printf( "We need at least two compute 1.0 or greater "
                "devices, but only found %d\n", deviceCount );
        return 0;
    }

    cudaDeviceProp  prop;
    for (int i=0; i<2; i++) {
        HANDLE_ERROR( cudaGetDeviceProperties( &prop, i ) );
        if (prop.canMapHostMemory != 1) {
            printf( "Device %d can not map memory.\n", i );
            return 0;
        }
    }

    float *a, *b;
    HANDLE_ERROR( cudaSetDevice( 0 ) );
    HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&a, N*sizeof(float),
                              cudaHostAllocWriteCombined |
                              cudaHostAllocPortable |
                              cudaHostAllocMapped ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&b, N*sizeof(float),
                              cudaHostAllocWriteCombined |
                              cudaHostAllocPortable      |
                              cudaHostAllocMapped ) );

    // fill in the host memory with data
    for (int i=0; i<N; i++) {
        a[i] = i;
        b[i] = i*2;
    }

    // prepare for multithread
    DataStruct  data[2];
    data[0].deviceID = 0;
    data[0].offset = 0;
    data[0].size = N/2;
    data[0].a = a;
    data[0].b = b;

    data[1].deviceID = 1;
    data[1].offset = N/2;
    data[1].size = N/2;
    data[1].a = a;
    data[1].b = b;

    CUTThread   thread = start_thread( routine, &(data[1]) );
    routine( &(data[0]) );
    end_thread( thread );


    // free memory on the CPU side
    HANDLE_ERROR( cudaFreeHost( a ) );
    HANDLE_ERROR( cudaFreeHost( b ) );

    printf( "Value calculated:  %f\n",
            data[0].returnValue + data[1].returnValue );

    return 0;
}
注意routine() 中的if很微妙， 
    if (data->deviceID != 0) {
        HANDLE_ERROR( cudaSetDevice( data->deviceID ) );
        HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );
    }
yongwuzhijing800
关注
  • 2
    点赞
  • 踩
  • 4
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 0
    评论
专栏目录
CUDA_主机内存
weixin_48151686的博客
11-09 821
主机内存 系统中被CPU访问的内存，分为两种类型：可分页内存（pageable memory,一般应用中默认使用）和页锁定内存（page-locked或者pinned）。 可分页内存即为通过操作系统api（malloc(), new()）分配的存储器空间；而页锁定内存始终不会被分配到低速的虚拟内存中，能够保证存在于物理内存中，并且能够通过DMA加速与设备端的通信。 为了让硬件使用DMA，操作系统允许主机内存进行页锁定，并且因为性能原因，CUDA包含了开发者使用这些操作系统工具的API。页锁定后的且映射为cu
GPU高性能编程CUDA实战.pdf.zip
07-25
《GPU高性能编程CUDA实战》是一本深度探讨GPU计算能力与CUDA编程技术的专著，旨在帮助读者掌握如何利用CUDA架构实现高效能计算。CUDA（Compute Unified Device Architecture）是由NVIDIA公司推出的并行计算平台和...
cudaMallocHost函数详解
Bruce_0712的博客
07-02 2万+
在CUDA2.2以下，仅提供cudaMallocHost函数用于分配页锁定内存，与C语言函数malloc分配分页内存相对应。 而从CUDA2.2开始，页锁定内存增加三种新的类型用于主机多线程的portable，用于高效写回write-combined以及零拷贝的mapping，用cudaHostAlloc进行分配，其中采用四个可选参数标志用以指定使用何种特性： cudaHostAllocDef
C语言与硬件加速编程：OpenCL、CUDA C在异构计算中的应用（三）
最新发布
weixin_56154577的博客
04-18 1060
C语言作为硬件加速编程的基础语言，与OpenCL、CUDA C等框架一道，在异构计算领域发挥着核心作用。面对新兴技术和新型硬件带来的机遇与挑战，持续的技术创新与生态建设将是保持其竞争力的关键。
CUDA Samples: dot product(使用零拷贝内存)
网络资源是无限的
08-02 4600
CUDA Samples: dot product(使用零拷贝内存)
OpenCV源码解读：Mat与GpuMat的内存拷贝（copyTo和clone）
TracelessLe的专栏
02-04 3486
前言 OpenCV基础类型Mat以及GPU上的GpuMat是整个OpenCV工程体系中两种非常基础常用的类型。通常情况下Mat拥有的方法，GpuMat都会有对应的方法。对于内存拷贝，有两个常用的方法copyTo()和clone()，在使用的时候需要特别注意。 Mat clone Mat Mat::clone() const { Mat m; copyTo(m); return m; } copyTo void Mat::copyTo( OutputArray _dst ) co
GPU基础与CUDA编程入门
qq_43398345的博客
05-04 2181
预取是一种内存管理策略，旨在减少内存访问延迟，从而提高计算性能。预取机制通过预先加载数据到高速缓存（例如，从全局内存到共享内存或纹理内存）来实现这一目标，以便在执行计算任务时减少等待时间。GPU的预取机制有两种形式：硬件预取：这是由GPU硬件自动实现的预取机制，不需要程序员进行显式操作。GPU内部的内存控制器会预测内存访问模式，提前将可能需要的数据加载到高速缓存中。这种预取机制在许多现代GPU架构（如NVIDIA的Pascal、Volta和Ampere架构）中都有实现。
GPU高性能编程CUDA实战
11-01
UDA是一种专门为提高并行程序开发效率而设计的计算架构。..., 《GPU高性能编程CUDA实战》适合具备C或者C++知识的应用程序开发人员、数值计算库开发人员等，也可以作为学习并行计算的学生和教师的教辅。
GPU高性能编程CUDA实战—示例代码
10-28
本压缩包中的“GPU高性能编程CUDA实战—示例代码”提供了详细的学习和实践CUDA编程的实例，旨在帮助读者深入理解和掌握CUDA编程技术。 首先，CUDA编程的核心概念是线程和线程块。在CUDA中，程序员可以创建大量的...
GPU高性能编程CUDA实战代码
01-12
本书“GPU高性能编程CUDA实战”旨在帮助读者掌握CUDA编程技术，实现GPU的高效计算潜力。CUDA不仅为开发者提供了与硬件交互的底层接口，还提供了一套丰富的库，如cuBLAS、cuFFT等，以简化常见计算任务。 在CUDA编程...
第9章 多GPU编程
weixin_47955824的博客
03-04 7526
多GPU管理； 跨多GPU执行核函数； GPU间的叠加计算和通信； GPU间的同步； 使用CUDA-aware MPI交换数据； 使用GPU Direct RDAM的CUDA-aware MPI交换数据； 跨GPU加速集群扩展应用程序； 理解CPU和GPU的亲和性；
CUDA平台下多核GPU高性能并行编程
01-03
CUDA编程是一种异构系统编程,即CPu+GPU 混合编程,一个完整的CuDA程序分为cPu串行部分 与GPU并行部分D’。
CUDA by example （中文：GPU高性能编程CUDA实战）代码实例
08-28
CUDA by example代码实例
浅析GPU计算——cuda编程
热门推荐
方亮的专栏
02-13 4万+
在《浅析GPU计算——CPU和GPU的选择》一文中，我们分析了在遇到什么瓶颈时需要考虑使用GPU去进行计算。本文将结合cuda编程来讲解实际应用例子。（转载请指明出于breaksoftware的csdn博客） 之前我们讲解过，CPU是整个计算机的核心，它的主要工作是负责调度各种资源，包括其自身的计算资源以及GPU的计算计算资源。比如一个浮点数相乘逻辑，理论上我们可...
第8章 多CPU与GPU解决方案 摘录
weixin_47955824的博客
12-05 954
8.1 简介 在现代计算机系统中通常由多个设备：CPU和GPU。在CPU方面将讨论卡槽和核。卡槽是指主板用来放置CPU的物理插槽。一个CPU可能包含一个或多个核。每个核实际上式一个独立的实体。许多CPU和GPU卡槽都位于单个节点或计算机系统。 了解核、卡槽已经节点的物理布局，会使得我们更有效地调度和分配任务。 8.2 局部性 局部性原理在GPU和CPU上是相当不错的。接近设备的内存(GPU上共享内存或CPU上的高速缓存)会更快地访问。一个卡槽内(比如核之间)的通信远远快于和不同卡槽上的另一核的通信.
CUDA学习笔记05：Mem申请的所有方式
JackZhangNJU的专栏
10-01 903
昨天小组讨论，有人一直坚持让我申请制定位置的内存，我真的是一脸懵逼，作为编程又5年的我从来没想过这个问题，就好像C++中强制你申请L1 Cache缓存那样，我内心中有万只草泥马路过：“臣妾做不到啊！！！”，为了给给所有人讲清楚CUDA编程是什么，接下来的几篇博客主要就是讲解CUDA的硬件和软件结构，首先应该写硬件部分，但是我这里想先写CUDA的内存申请部分。 cudaMallocHost函数
gpu高性能编程cuda实战 源码
12-08
GPU高性能编程CUDA实战是一本介绍GPU编程和CUDA编程技术的实战指南。该书主要包括了大量的实例和源码示例，可供读者学习和实践。 该书首先介绍了GPU编程的基础知识，包括GPU架构和CUDA编程模型等。接着，书中详细介绍了CUDA编程的各个方面，包括线程管理、内存管理、流管理、并行算法等。每一章都有相应的实例和源码示例，能够帮助读者快速上手并理解各种GPU编程技术。 除了基础知识和技术介绍，该书还着重讲解了GPU编程的实战应用。它涵盖了各个领域的应用案例，包括图像处理、机器学习、物理模拟等。每一章都有相应的源码示例，读者可以通过实际案例学习如何将GPU编程技术应用到实际项目中。 此外，该书还介绍了一些优化和调试技术，帮助读者提高CUDA程序的性能。它涵盖了内存访问模式、并行算法优化等内容，并提供了一些常见问题的解决方案。读者可以通过学习这些技术，提高自己编写高性能CUDA程序的能力。 总之，GPU高性能编程CUDA实战是一本非常实用的书。它通过大量的实例和源码示例，帮助读者深入理解GPU编程和CUDA技术，提供了丰富的实战经验。无论是初学者还是有一定经验的读者，都可以从中获得很多实用的知识和技巧。

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

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

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

yongwuzhijing800

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

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

余额充值