简介
CUDAC不是C语言，而是对C语言进行扩展。CUDA对C的扩展主要包括以下四个方面：
- 函数类型限定符，用来确定函数是在CPU还是在GPU上执行，以及这个函数是从CPU调用还是从GPU调用。
- _device_，表示从GPU上调用，在GPU上执行；也就是说其可以被__global__或者__device__修饰的函数调用，此限定符修饰的函数使用有限制，比如在G80/GT200架构上不能使用递归，不能使用函数指针等。
- _global_，表示在CPU上调用，在GPU上执行，也就是所谓的内核(kernel)函数；内核只能够被主机调用，内核并不是一个完整的程序，它只是一个数据并行步骤，其指令流由多个线程执行。
- _host_，表示在CPU上调用，在CPU上执行，这是默认时的情况，也就是传统的C函数。
- 变量类型限定符，用来规定变量存储什么位置上。在传统的CPU程序上，这个任务由编译器承担。在CUDA中，不仅要使用主机端的内存，还要使用设备端的显存和GPU片上的寄存器、共享存储器和缓存。在CUDA存储器模型中，一共抽象出来了8种不同的存储器。复杂的存储器模型使得必须要使用限定符要说明变量的存储位置。
- _device_，__device__表明声明的数据存放在显存中，所有的线程都可以访问，而且主机也可以通过运行时库访问；
- _shared_，__shared__表示数据存放在共享存储器在，只有在所在的块内的线程可以访问，其它块内的线程不能访问；
- _constant_，__constant__表明数据存放在常量存储器中，可以被所有的线程访问，也可以被主机通过运行时库访问；
- Texture，texture表明其绑定的数据可以被纹理缓存加速存取，其实数据本身的存放位置并没有改变，纹理是来源于图形学的一介概念，CUDA使用它的原因一部分在于支持图形处理，另一方面也可以利用它的一些特殊功能；
- 如果在GPU上执行的函数内部的变量没有限定符，那表示它存放在寄存器或者本地存储器中，在寄存器中的数据只归线程所有，其它线程不可见；
- 如果SM的寄存器用完，那么编译器就会将本应放到寄存器中的变量放到本地存储器中。
- 执行配置运算符<<<>>>，用来传递内核函数的执行参数。执行配置有四个参数，第一个参数声明网格的大小，第二个参数声明块的大小，第三个参数声明动态分配的共享存储器大小，默认为0，最后一个参数声明执行的流，默认为0。
- 五个内建变量，用于在运行时获得网格和块的尺寸及线程索引等信息
- gridDim, gridDim是一个包含三个元素x,y,z的结构体，分别表示网格在x,y,z三个方向上的尺寸，虽然其有三维，但是目前只能使用二维；
- blockDim,blockDim也是一个包含三个元素x,y,z的结构体，分别表示块在x,y,z三个方向上的尺寸，对应于执行配置中的第一个参数，对应于执行配置的第二个参数；
- blockIdx,blockIdx也是一个包含三个元素x,y,z的结构体，分别表示当前线程所在块在网格中x,y,z三个方向上的索引；
- threadIdx,threadIdx也是一个包含三个元素x,y,z的结构体，分别表示当前线程在其所在块中x,y,z三个方向上的索引；
- warpSize，warpSize表明warp的尺寸，在计算能力为1.0的设备中，这个值是24，在1.0以上的设备中，这个值是32。
exp1
#include <cuda.h>
#include <iostream>
struct HostObject {
__host__
int value() const { return 42; }
};
struct DeviceObject {
__device__
int value() const { return 3; }
};
template <typename T>
__host__ __device__
int foo(const T &obj) {
return obj.value();
}
/*
template <typename T>
__host__
int foo_host(const T &obj) {
return obj.value();
}
template <typename T>
__device__
int foo_device(const T &obj) {
return obj.value();
}
*/
__global__ void kernel(int *data) {
data[threadIdx.x] = foo(DeviceObject());
}
int main() {
foo(HostObject());
int *data;
cudaMalloc((void**)&data, sizeof(int) * 64);
kernel<<<1, 64>>>(data);
cudaThreadSynchronize();
cudaFree(data);
}
#include<cuda.h>
#include<iostream>
struct HostObject{
__host__
int value() const {return 42;}
};
struct DeviceObject {
__device__
int value() const{return 3;}
};
template< typename T>
__host__ __device__
int foo(const T& obj){
return obj.value();
}
/*
template< typename T>
__host__
int foo_host(const T& obj){
return obj.value（）;
}
template< typename T>
__device__
int foo_device(const T& obj){
return obj.value();
}
*/
__global__ void kernel(int * data){
data [threadIdx.x] = foo(DeviceObject());
}
int main(){
foo(HostObject());
int *data;
cudaMalloc((void **)&data,sizeof(int)* 64);
kernel<<<1,664>>>(data);
cudaThreadSynchronize();
cudaFree(data);
}
#include<iostream>
#include<cstdio>
#pragma hd_warning_disable
template<class Function>
__host__ __device__
void invoke(Function f)
{
f();
}
struct host_only
{
__host__
void operator()(){
std::cout<< "host_only()"<< std::endl;
}
};
struct device_only
{
__device__
void operator()(){
printf("device_only()：thread%d \n",threadIdx.x);
}
};
__global__
void kernel()
{
//从带设备函数的设备使用
invoke(device_only());
// XXX错误
invoke(host_only());
}
int main()
{
//从带有主机函数的主机使用
invoke(host_only());
kernel<<<1,1>>>();
cudaDeviceSynchronize();
// XXX错误
// invoke（device_only（））;
return 0;
}
exp2
#include <iostream>
using namespace std;
template<typename Operation> __global__ void kernel_foreach(Operation o)
{
size_t i = blockIdx.x * blockDim.x + threadIdx.x;
o(i);
}
template<bool onDevice, typename Operation> void foreach(size_t size, Operation o)
{
if(onDevice)
{
size_t blocksize = 32;
size_t gridsize = size/32;
kernel_foreach<<<gridsize,blocksize>>>( o );
}
else
{
for( size_t i = 0; i < size; ++i )
{
o(i);
}
}
}
__global__ void printFirstElementOnDevice(double* vector)
{
printf("dVector[0] = %f\n", vector[0]);
}
template<bool onDevice> void assignScalar( size_t size, double* vector, double a )
{
auto assign = [=] __host__ __device__ (size_t i) { vector[i] = a; };
foreach<onDevice>(size,assign);
}
int main()
{
size_t SIZE = 32;
double* hVector = new double[SIZE];
double* dVector;
cudaMalloc( &dVector, SIZE*sizeof(double) );
// clear memory
for( size_t i = 0; i < SIZE; ++i )
{
hVector[i] = 0;
}
cudaMemcpy(dVector, hVector, SIZE*sizeof(double), cudaMemcpyHostToDevice);
assignScalar<false>(SIZE, hVector,3.0);
cout << "hVector[0] = " << hVector[0] << endl;
assignScalar<true>(SIZE, dVector, 4.0);
printFirstElementOnDevice<<<1,1>>>( dVector );
cudaDeviceSynchronize();
cudaError_t error = cudaGetLastError();
if(error!=cudaSuccess)
{
cout << "ERROR: " << cudaGetErrorString(error);
}
}
nvcc example1.cu --expt-extended-lambda -o example
./example
hVector[0] = 3
dVector[0] = 4.000000
exp3
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
__global__ void myfunc(void){
//printf("myfunc\n");
}
int main(void){
myfunc<<<2,2>>>();
printf("main\n");
return 0;
}
Makefile
CUFLAG = -g -Xcompiler -v \
-O2
CUDA_DIR = /usr/local/cuda
IFLAGS = -I$(CUDA_DIR)/include
LFLAGS = -L$(CUDA_DIR)/lib64
PRG = cuda_test
$(PRG) : example1.cu
nvcc example1.cu -o $(PRG) $(CUFLAG) $(IFLAGS) $(LFLAGS)
上述代码编译后运行生成可执行文件cuda_test，运行cuda_test后将输出：
Hello World!
注意：
调用kernel时需要三个尖括号
包含必要的头文件
CUDA C/C++中引入的新关键字__global__所修饰的函数有以下两方面含义：
此函数代码由设备执行
此函数由主机代码调用
nvcc将源代码分为设备函数和主机函数两大类：
设备函数由NVIDA编译器编译
主机函数由主机上配置的编译器编译
三个尖括号标志着一个从主机代码调用设备代码的函数，称为“启动内核”(kernel launch)
exp3
int main(void) {
int a,b,c;
int * d_a, * d_b, * d_c;
int size = sizeof(int);
cudaMalloc((void**)&d_a,size);
cudaMalloc((void**)&d_b,size);
cudaMalloc((void**)&d_c,size);
printf("Enter two integers with a space to separate them:\n");
scanf("%d %d",&a,&b);
cudaMemcpy(d_a,&a,size,cudaMemcpyHostToDevice);
cudaMemcpy(d_b,&b,size,cudaMemcpyHostToDevice);
integer_add<<<1,1>>>(d_a,d_b,d_c);
cudaMemcpy(&c,d_c,size,cudaMemcpyDeviceToHost);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
printf("Sum is %d\n",c);
return 0;
}
__global__修饰的integer_add函数说明：
integer_add函数将在设备上执行
integer_add函数将被主机调用
由于integer_add函数在设备上执行，所以指针a，b，c应该指向设备内存。这说明需要在设备内存中为变量开辟内存。
设备内存和主机内存在物理上是完全分开的不同电子部件：
设备指针指向GPU内存的某个位置。设备指针可以从主机端传给设备端或者从设备端传给主机端，但是设备指针不能在主机端解引用。
主机指针指向CPU内存的某个位置。主机指针可以从设备端传给主机端或者从主机端传给设备端，但是主机指针不能在设备端解引用。
CUDA API提供的用于处理设备内存的函数有cudaMalloc, cudaFree, cudaMemcpy。语义上分别对应于C语言的malloc, free, memcpy函数。这几个函数的具体使用方法如例2所示。
exp4
#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<time.h>
#define N 512
__global__ void vec_block_add(int * a, int * b, int * c) {
c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}
void rand_ints(int * arr, int count) {
srand(time(NULL));
for(int i=0;i<count;i++) {
arr[i] = rand() % 100;
}
}
int main(void) {
int * a,* b,* c;
int * d_a, * d_b, * d_c;
int size = N * sizeof(int);
cudaMalloc((void**)&d_a,size);
cudaMalloc((void**)&d_b,size);
cudaMalloc((void**)&d_c,size);
a = (int *) malloc(size);
rand_ints(a,N);
b = (int *) malloc(size);
rand_ints(b,N);
c = (int *) malloc(size);
cudaMemcpy(d_a,a,size,cudaMemcpyHostToDevice);
cudaMemcpy(d_b,b,size,cudaMemcpyHostToDevice);
vec_block_add<<<N,1>>>(d_a,d_b,d_c);
cudaMemcpy(c,d_c,size,cudaMemcpyDeviceToHost);
#if 1
for(int i=0;i<N;i++) {
printf("%-5d: a:%-5d b:%-5d c:%-5d\n",i,a[i],b[i],c[i]);
}
#endif
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(a);
free(b);
free(c);
return 0;
}
如果我们要实现两个向量相加：
add<<<1,1>>>() —> add<<<N,1>>>
N表示同时调用N次add函数，这样就可以实现并行的向量相加了。
每个被并行调用的add函数称之为一个 块(block)。
块的集合称之为网格(grid).
每个块可以使用索引值blockIdx.x
通过使用blockIdx.x作为索引，每个块可以处理数组元素中的一部分。
有了这些基础后，就可以实现并行版本的向量相加了。
由于函数是并行执行的，和传统的串行程序在integer_add函数中使用 循环来完成加法相比，相当于由GPU这个加速器使用硬件的方式进行了 循环展开，展开后便可以并行执行了。所以在编写这段代码时，需要使用blockIdx.x来定位当前执行的是 循环的哪个部分。
从硬件的角度看，相当于同时有多个块在并行执行：
块0: c[0]=a[0]+b[0]
块1: c[1]=a[1]+b[1]
块2: c[2]=a[2]+b[2]
块3: c[3]=a[3]+b[3]
exp5 线程(Threads)
简单的说就是一个 块(Block)可以分割成多个 线程(Threads).
所以可以将上述的多块单线程版本改成单块多线程版本。
#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<time.h>
#define N 512
__global__ void vec_thread_add(int * a, int * b, int * c) {
c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
}
void rand_ints(int * arr, int count) {
srand(time(NULL));
for(int i=0;i<count;i++) {
arr[i] = rand() % 100;
}
}
int main(void) {
int * a,* b,* c;
int * d_a, * d_b, * d_c;
int size = N * sizeof(int);
cudaMalloc((void**)&d_a,size);
cudaMalloc((void**)&d_b,size);
cudaMalloc((void**)&d_c,size);
a = (int *) malloc(size);
rand_ints(a,N);
b = (int *) malloc(size);
rand_ints(b,N);
c = (int *) malloc(size);
cudaMemcpy(d_a,a,size,cudaMemcpyHostToDevice);
cudaMemcpy(d_b,b,size,cudaMemcpyHostToDevice);
vec_thread_add<<<1,N>>>(d_a,d_b,d_c);
cudaMemcpy(c,d_c,size,cudaMemcpyDeviceToHost);
#if 1
for(int i=0;i<N;i++) {
printf("%-5d: a:%-5d b:%-5d c:%-5d\n",i,a[i],b[i],c[i]);
}
#endif
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(a);
free(b);
free(c);
return 0;
}
上述代码vec_thread_add<<<1,N>>>指定向量相加操作仅有1个块完成，但是这个块可以分割成512个线程来完成这个计算。和块类似的，由于函数vec_thread_add是被多个线程并行展开循环的，所以需要根据线程编号来确定当前循环应该完成的计算部分。每个线程使用threadIdx.x来标识当前线程。
exp6 索引
实现多块多线程并行程序的核心是利用进行准确的索引计算，即准确的利用索引将某个线程和该线>程要完成计算的数据对应起来。对于多块多线程的情况，我们假设要处理的向量包含128个元素，可以设计16个块，每个块中再包>含8个线程来完成计算（因为16*8=128）。对于任何一个元素，我们都应该能准确的计算出来由哪>个线程来完成这个相加计算。比如对于编号为71的元素：
块号为：71/8=8
线程号为：71%8=7
所以71号元素应该由：
71 = threadIdx.x + blockIdx.x * 8 = 7 + 8 * 8
在CUDA中可以使用blockDim.x来表示每个块中的线程数量，所以以上的计算可以修改为
元素编号 = threadIdx.x + blockIdx.x * blockDim.x
多块多线程的向量相加程序如下：
#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<time.h>
#define N (16*16)
#define THREAD_PER_BLOCK 32
__global__ void vec_block_thread_add(int * a, int * b, int * c) {
int index = threadIdx.x + blockIdx.x * blockDim.x;
c[index] = a[index] + b[index];
}
void rand_ints(int * arr, int count) {
srand(time(NULL));
for(int i=0;i<count;i++) {
arr[i] = rand() % 100;
}
}
int main(void) {
int * a,* b,* c;
int * d_a, * d_b, * d_c;
int size = N * sizeof(int);
cudaMalloc((void**)&d_a,size);
cudaMalloc((void**)&d_b,size);
cudaMalloc((void**)&d_c,size);
a = (int *) malloc(size);
rand_ints(a,N);
b = (int *) malloc(size);
rand_ints(b,N);
c = (int *) malloc(size);
cudaMemcpy(d_a,a,size,cudaMemcpyHostToDevice);
cudaMemcpy(d_b,b,size,cudaMemcpyHostToDevice);
vec_block_thread_add<<<N/THREAD_PER_BLOCK,THREAD_PER_BLOCK>>>(d_a,d_b,d_c);
cudaMemcpy(c,d_c,size,cudaMemcpyDeviceToHost);
#if 1
for(int i=0;i<N;i++) {
printf("%-5d: a:%-5d b:%-5d c:%-5d\n",i,a[i],b[i],c[i]);
}
#endif
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(a);
free(b);
free(c);
return 0;
}
当需要计算元素的数量不能够整除每块中线程数量
#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<time.h>
#define N 100
#define M 32
__global__ void vec_block_thread_add(int * a, int * b, int * c, int n )
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
if(index < n) {
c[index] = a[index] + b[index];
}
}
void rand_ints(int * arr, int count) {
srand(time(NULL));
for(int i=0;i<count;i++) {
arr[i] = rand() % 100;
}
}
int main(void) {
int * a,* b,* c;
int * d_a, * d_b, * d_c;
int size = N * sizeof(int);
cudaMalloc((void**)&d_a,size);
cudaMalloc((void**)&d_b,size);
cudaMalloc((void**)&d_c,size);
a = (int *) malloc(size);
rand_ints(a,N);
b = (int *) malloc(size);
rand_ints(b,N);
c = (int *) malloc(size);
cudaMemcpy(d_a,a,size,cudaMemcpyHostToDevice);
cudaMemcpy(d_b,b,size,cudaMemcpyHostToDevice);
vec_block_thread_add<<<(N+M-1)/M,M>>>(d_a,d_b,d_c,N);
cudaMemcpy(c,d_c,size,cudaMemcpyDeviceToHost);
#if 1
for(int i=0;i<N;i++) {
printf("%-5d: a:%-5d b:%-5d c:%-5d\n",i,a[i],b[i],c[i]);
}
#endif
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(a);
free(b);
free(c);
return 0;
}
exp7 共享内存(Shared Memory)
提供两个数组，输入数组in和输出数组out. 一共有n个元素，将in中的值映射到out，映射的规则是：对于任意的i，out[i] = out[ i - radius ] + out[ i - radius + 1] + out[i-radius+2] + … + out[i+radius].
这里我们会发现同一块内的线程比如，2号线程和3号线程访问的数据将会有大量的数据重复。如果能有cache机制，将有效的降低从GPU的DRAM中加载数据所消耗的时间。共享内存(shared memory)就是用来实现这个功能的：
- 在同一个块内线程通过共享内存实现数据共享或者说在同一个块内，线程通过共享内存实现线程间的通信
- 硬件上看，共享内存是速度极高的片上内存。而设备内存，又称之为全局内存(global memory)的速度较慢
- 可以将共享内存看做是用户管理的cache
- 使用关键字__shared__定义的存储空间将在共享内存上为每个块开辟空间
- 一个块内的共享内存数据对另外一个块的线程来说是不可见的
对每一个块来说:
需要读入2 * radius + blockDim.x 个元素的数据。
输出blockDim.x个元素数据到全局内存中.
#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<time.h>
#define N 256
#define RADIUS 2
#define BLOCK_SIZE 32
__global__ void stencil_1d(int * in, int *out)
{
__shared__ int temp[BLOCK_SIZE + 2 * RADIUS];
int g_index = threadIdx.x + blockIdx.x * blockDim.x + RADIUS;
int s_index = threadIdx.x + RADIUS;
temp[s_index] = in[g_index];
if(threadIdx.x < RADIUS) {
temp[s_index - RADIUS] = in[g_index - RADIUS];
temp[s_index + BLOCK_SIZE] = in[g_index + BLOCK_SIZE];
}
int result = 0;
for(int offset = -RADIUS; offset <= RADIUS; offset++) {
result = result + temp[s_index+offset];
}
out[g_index-RADIUS] = result;
}
void rand_ints(int * arr, int count) {
srand(time(NULL));
for(int i=0;i<count;i++) {
arr[i] = rand() % 100;
}
}
int main(void) {
int * in, * out;
int size_in = sizeof(int) * ( N + 2 * RADIUS );
in = (int *) malloc(size_in);
rand_ints(in+RADIUS,N);
int size_out = sizeof(int) * N;
out = (int *) malloc(size_out);
int * d_in, * d_out;
cudaMalloc((void**)&d_in,size_in);
cudaMalloc((void**)&d_out,size_out);
cudaMemcpy(d_in,in,size_in,cudaMemcpyHostToDevice);
stencil_1d<<<(N+BLOCK_SIZE-1)/BLOCK_SIZE,BLOCK_SIZE>>>(d_in,d_out);
cudaMemcpy(out,d_out,size_out,cudaMemcpyDeviceToHost);
#if 1
for(int i=0;i<N+2*RADIUS;i++) {
printf("%-5d ",in[i]);
}
printf("\n");
for(int i=0;i<N;i++) {
printf("%-5d ",out[i]);
}
printf("\n");
#endif
cudaFree(d_in);
cudaFree(d_out);
free(in);
free(out);
return 0;
}
exp8 线程同步函数(_syncthreads())
上述代码因为没有考虑数据竞争可能产生错误的结果。比如第2个块中的第三个线程进入for循环开始进行相加计算了，但是它需要读的数据还没有被4号线程写入到存储空间就会产生数据竞争问题。
CUDA提供了__syncthreads函数来进行数据同步，这样就可以保证所有的数据都就绪后再开始进行计算,修改后的代码如下：
#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<cutil.h>
#include<time.h>
#define N 256
#define RADIUS 2
#define BLOCK_SIZE 32
__global__ void stencil_1d(int * in, int *out)
{
__shared__ int temp[BLOCK_SIZE + 2 * RADIUS];
int g_index = threadIdx.x + blockIdx.x * blockDim.x + RADIUS;
int s_index = threadIdx.x + RADIUS;
temp[s_index] = in[g_index];
if(threadIdx.x < RADIUS) {
temp[s_index - RADIUS] = in[g_index - RADIUS];
temp[s_index + BLOCK_SIZE] = in[g_index + BLOCK_SIZE];
}
__syncthreads();
int result = 0;
for(int offset = -RADIUS; offset <= RADIUS; offset++) {
result = result + temp[s_index+offset];
}
out[g_index-RADIUS] = result;
}
void rand_ints(int * arr, int count) {
srand(time(NULL));
for(int i=0;i<count;i++) {
arr[i] = rand() % 100;
}
}
int main(void) {
int * in, * out;
int size_in = sizeof(int) * ( N + 2 * RADIUS );
in = (int *) malloc(size_in);
rand_ints(in+RADIUS,N);
int size_out = sizeof(int) * N;
out = (int *) malloc(size_out);
int * d_in, * d_out;
cudaMalloc((void**)&d_in,size_in);
cudaMalloc((void**)&d_out,size_out);
cudaMemcpy(d_in,in,size_in,cudaMemcpyHostToDevice);
stencil_1d<<<(N+BLOCK_SIZE-1)/BLOCK_SIZE,BLOCK_SIZE>>>(d_in,d_out);
cudaMemcpy(out,d_out,size_out,cudaMemcpyDeviceToHost);
#if 1
for(int i=0;i<N+2*RADIUS;i++) {
printf("%-5d ",in[i]);
}
printf("\n");
for(int i=0;i<N;i++) {
printf("%-5d ",out[i]);
}
printf("\n");
#endif
cudaFree(d_in);
cudaFree(d_out);
free(in);
free(out);
return 0;
}
exp9 异步操作(asynchronous operations)
- CUDA kernel函数可以异步启动，普通的kernel函数需要等到执行完毕后才能将控制流返回到CPU端。异步启动方式指的是，kernel启动之后控制流立即返回到CPU端。
- CUDA 的cudaMemcpy函数是等待数据拷贝完毕后才返回到CPU端。CUDA提供cudaMemcpyAsync来支持异步的数据拷贝。
- CUDA提供的接口函数cudaDeviceSynchronize函数将阻塞CPU直到所有的CUDA函数调用都执行完毕。
exp10 故障处理(handle errors)
- 所有的CUDA函数都会返回一个错误码，其类型为cudaError_t。错误可能是当前函数调用产生的也可能是之前发起的异步函数执行后产生的。
- 得到最后一个错误的函数是： cudaError_t cudaGetLastError(void);
- 将错误码转换为字符串描述信息的函数是： char * cudaGetErrorString(cudaError_t);
exp11设备管理(managing devices)
- 应用程序可以查询GPU的数量或者是选择GPU.
- 查询设备数量的接口： cudaGetDeviceCount(int * count);
- 选择设备的接口： cudaSetDevice(int device);
- 获取设备的接口：cudaGetDevice(int * device);
- 获取设备属性的接口： cudaGetDeviceProperties(cudaDeviceProp * prop, int device);
- 多个主机线程可以共享一个设备。
- 单个主机线程可以管理多个GPU设备。
1096
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
