[已注销]对《GPU高性能运算之CUDA》的笔记(5)

GPU高性能运算之CUDA
  • 书名: GPU高性能运算之CUDA
  • 作者: 张舒/褚艳利
  • 页数: 276
  • 出版社: 中国水利水电出版社
  • 出版年: 2009-10
  • 第35页
    带状划分
    
    dim3 grid(4, 1, 1);
    dim3 threads(4, 1, 1);
    
    grid:
    | block | block | block | block |
    block:
    | thread | thread | thread | thread |
    棋盘划分
    dim3 grid(2, 2, 1);
    dim3 threads(2, 2, 1);
    
    grid:
    | block | block |
    | block | block |
    block:
    | thread | thread |
    | thread | thread |
    2012-02-13 20:47:19 回应
  • 第45页

    == local memory
    c 语言中常说的 auto variable 在 cuda c 中是不存在的,因为 gpu 执行线程时使用的不是传统的调用栈模型,所有的局部变量都是分配在片上的 register 中,相当于 c 中的 register variable。
    而当 kernel 函数中声明了较大的数组或结构体时,变量会分配到显存中去,即 local memory。
    == shared memory
    和 local memory 相对的是 shared memory,其中的数据可以被同一 block 中的 thread 共享。使用 shared memory 是实现线程间通信延迟最小的方法。用关键字 __shared__ 声明。
    == global memory
    显存中还有一块 global memory,cpu 和 gpu 都可以访问。它的特点是带宽大,延迟高。
    global memory(又称线性内存)的分配和释放:
    //主机端代码
    int main() {
        size_t size = N * sizeof(float);
        float* d_A;
        cudaMalloc((void**)&d_A, size);
        //处理
        cudaFree(d_A);
    }
    
    
    
    == constant memory
    用于只读的数据分配,节约带宽。必须定义在所有函数之外。
    __constant__ int num = 11;
    
    __global__ static void XX() {...}
    
    
    == texture memory
    2012-03-03 16:00:37 回应
  • 第72页
    volatile 是 c 里面就有的一个关键字,用来防止编译器的优化导致程序出错。
    int ref1 = array[tid];
    int ref2 = array[tid]; 
    
    正常情况下,编译器会把第一次的赋值操作编译为一个内存读操作,而第二次编译为
    int ref2 = ref1;
    
    而 ref1 可能在寄存器中,这样可以避免访问内存,节省了时间。因为编译器认为 array[tid] 的值在期间没有变化,但是如果有多个线程时就不能保证这点。
    可以将 array 声明为 volatile 变量,这样的话强制让每次的赋值操作都会编译为访寸操作。
    2012-02-16 21:34:29 回应
  • 第157页
    warp 是调度和执行的基本单位,half-warp 是存储器操作的基本单位标为代码。
    shared memory(位于GPU片内)bank conflict
    每个 bank 大小为 32 Bit,在分配数组时让每个元素占一个 bank,可以避免 bank conflict。
    __shared__ char shared[32];
    char data = shared[BaseIndex + 4 * tid]; 
    
    访问 local memory 或 global memory 时需要 400~600 周期的访问延时,因此在访问时尽量要想办法把这段延时隐藏起来,比如异步执行。
    1.cpu 和 gpu 之间的异步执行
    2. gpu不同流之间的异步执行
    ==== ==== ==== ==== 数据拷贝
    ==== ==== ==== ==== 执行
    就相当于本来是先盛完了饭再吃,现在是一边盛一边吃。
    2012-02-21 12:28:55 回应
  • 第190页
    第5章暂时还用不到,先列个提纲
    -应用
    ----基本
    ---------双调排序网络
    ---------求all-prefix-sums
    ---------数值计算库-CUBLAS
    ---------FFT库-CUFFT
    ----高级
    ---------共轭梯度法
    ---------AC多模式匹配算法
    2012-02-21 12:53:12 回应

[已注销]的其他笔记  · · · · · ·  ( 全部295条 )

The Design of Everyday Things
1
About Face 3
6
Engineering a Compiler
1
人有人的用处
8
Understanding the Linux Kernel
4
计算机程序设计艺术
1
公正
1
The Art of Doing Science and Engineering: Learning to Learn
1
科学革命的结构
7
罗素论教育
3
三十六大
1
娱乐至死
3
Real World Haskell
2
Writing Analytically
1
Is Parallel Programming Hard, And, If So, What Can You Do About It?
1
计算机与人脑
1
组合数学
2
菊与刀
1
Rework
5
翻译新究
4
计算机程序的构造和解释(原书第2版)
5
The Laws of Simplicity
4
计算机组成与设计硬件/软件接口
6
写给无神论者
2
放任自流的时光
3
哥德尔、艾舍尔、巴赫
2
树上的男爵
2
C++语言的设计与演化
1
Land of LISP
7
C陷阱与缺陷
2
CUDA by Example
3
C++沉思录
1
世界尽头与冷酷仙境
4
Head First C
2
刀锋
1
并行编程模式
2
The Ph.D. Grind
2
计算机系统结构
2
禅与摩托车维修艺术
14
流浪的面包树
2
翻译研究
18
An Introduction to Programming in Emacs Lisp
1
GNU Emacs Lisp 编程入门
1
计算机系统概论
1
编码
3
拖延心理学
1
古今数学思想(一)
1
挪威的森林
9
奇特的一生
7
那些年,我们一起追的女孩
8
十八岁给我一个姑娘
2
C++编程思想(第1卷)
9
多核计算与程序设计
8
少有人走的路
5
忧伤的情欲
3
Hackers & Painters
7
哲学的慰藉
9
男人来自火星 女人来自金星
8
旅行的艺术
14
活着活着就老了
1
如何阅读一本书
7
Data-intensive Text Processing With Mapreduce
1
学习GNU Emacs
1
给研究生的学术建议
3
C专家编程
2
Spring揭秘
2
Head First Java(第二版·中文版)
1
自私的基因
1
C程序设计语言
2
计算机网络
3
自由在高处
2
大话设计模式
16
计算机网络
12