炼数成金 › 门户 ›CUDA › 查看内容

CUDA学习之主机运行时组件

2015-9-8 16:16| 发布者: 炼数成金_小数| 查看: 942| 评论: 0|原作者: yanghangjun|来自: yanghangjun的博客

摘要: 只有主机函数才能使用主机运行时组件。它提供了具有以下功能的函数：n 设备管理；n 上下文管理；n 存储器管理；n 代码模块管理；n 执行控制；n 纹理参考管理；n 与 OpenGL 和 Direct3D ...

管理基础存储内核 CUDA

只有主机函数才能使用主机运行时组件。

它提供了具有以下功能的函数：

n 设备管理；

n 上下文管理；

n 存储器管理；

n 代码模块管理；

n 执行控制；

n 纹理参考管理；

n 与 OpenGL 和 Direct3D 的互操作性。

它包含两个 API：

n 一个称为 CUDA 驱动程序 API 的低级 API；

n 一个称为 CUDA 运行时 API 的高级 API，它是在 CUDA 驱动程序 API 的基础之上实现的。

这些 API 是互斥的：一个应用程序仅能使用其中之一。

CUDA 驱动程序 API 是通过 nvcuda 动态库提供的，其所有入口点都带有 cu 前缀。

CUDA 运行时 API 是通过 cudart 动态库提供的，其所有入口点都带有 cuda 前缀。

运行时API

cudaGetDeviceCount() 和 cudaGetDeviceProperties() 提供了一种方法，用于枚举这些设备并检索其属性：

int deviceCount;

cudaGetDeviceCount(&deviceCount);

int device;

for (device = 0; device < deviceCount; ++device) {

cudaDeviceProp deviceProp;

cudaGetDeviceProperties(&deviceProp, device);

}

cudaSetDevice() 用于选择与主机线程相关的设备：

cudaSetDevice(device);

必须首先选择设备，之后才能调用 _global_ 函数或任何来自运行时 API 的函数。如果未通过显式调用 cudaSetDevice() 完成此任务，将自动选中设备 0，随后对 cudaSetDevice() 的任何显式调用都将无效。

存储器管理：

线性存储器是使用 cudaMalloc() 或 cudaMallocPitch() 分配的，使用 cudaFree() 释放。

以下示例代码将在线性存储器中分配一个包含 256 个浮点元素的数组：

float* devPtr;

cudaMalloc((void**)&devPtr, 256 * sizeof(float));

建议在分配二维数组时使用 cudaMallocPitch()，因为它能确保合理填充已分配的存储器，满足第 5.1.2.1 节介绍的对齐要求，从而确保访问行地址或执行二维数组与设备存储器的其他区域之间的复制（使用 cudaMemcpy2D()）时获得最优性能。所返回的间距（或步幅）必须用于访问数组元素。以下代码示例将分配一个 widthxheight 的二维浮点值数组，并显示如何在设备代码中循环遍历数组元素：

// host code

float* devPtr;

int pitch;

cudaMallocPitch((void**)&devPtr, &pitch,

width * sizeof(float), height);

myKernel<<<100, 512>>>(devPtr, pitch);

// device code

__global__ void myKernel(float* devPtr, int pitch)

{

for (int r = 0; r < height; ++r) {

float* row = (float*)((char*)devPtr + r * pitch);

for (int c = 0; c < width; ++c) {

float element = row[c];

}

CUDA 数组是使用 cudaMallocArray() 分配的，使用 cudaFreeArray() 释放。cudaMallocArray() 需要使用 cudaCreateChannelDesc() 创建的格式描述。

以下代码示例分配了一个 widthxheight 的 CUDA 数组，包含一个 32 位的浮点组件：

cudaChannelFormatDesc channelDesc =

cudaCreateChannelDesc<float>();

cudaArray* cuArray;

cudaMallocArray(&cuArray, &channelDesc, width, height);

cudaGetSymbolAddress() 用于检索指向为全局存储器空间中声明的变量分配的存储器的地址。所分配存储器的大小是通过 cudaGetSymbolSize() 获取的。

参考手册列举了用于在 cudaMalloc() 分配的线性存储器、cudaMallocPitch() 分配的线性存储器、CUDA 数组和为全局或固定存储器空间中声明的变量分配的存储器之间复制存储器的所有函数。

下面的代码示例将二维数组复制到之前代码示例中分配的 CUDA 数组中：

cudaMemcpy2DToArray(cuArray, 0, 0, devPtr, pitch,

width * sizeof(float), height,

cudaMemcpyDeviceToDevice);

下面的代码示例将一些主机存储器数组复制到设备存储器中：

float data[256];

int size = sizeof(data);

float* devPtr;

cudaMalloc((void**)&devPtr, size);

cudaMemcpy(devPtr, data, size, cudaMemcpyHostToDevice);

下面的代码示例将一些主机存储器数组复制到固定存储器中：

__constant__ float constData[256];

float data[256];

cudaMemcpyToSymbol(constData, data, sizeof(data));

流管理：

以下代码示例创建两个流：

cudaStream_t stream[2];

for (int i = 0; i < 2; ++i)

cudaStreamCreate(&stream[i]);

这些流均通过以下代码示例定义为一个序列，包括一次从主机到设备的存储器复制、一次内核启动、一次从设备到主机的存储器复制：

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,

size, cudaMemcpyHostToDevice, stream[i]);

for (int i = 0; i < 2; ++i)

myKernel<<<100, 512, 0, stream[i]>>>

(outputDevPtr + i * size, inputDevPtr + i * size, size);

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,

size, cudaMemcpyDeviceToHost, stream[i]);

cudaThreadSynchronize();

两个流均会将其输入数组 hostPtr 的一部分复制到设备存储器的 inputDevPtr 数组中，通过调用 myKernel() 处理设备上的 inputDevPtr，并将结果 outputDevPtr 复制回 hostPtr 的相同部分。使用两个流处理 hostPtr 允许一个流的存储器复制与另外一个流的内核执行相互重叠。hostPtr 必须指向分页锁定的主机存储器，这样才能出现重叠：

float* hostPtr;

cudaMallocHost((void**)&hostPtr, 2 * size);

最后调用了 cudaThreadSynchronize()，目的是在进一步处理之前确定所有流均已完成。cudaStreamsynchronize() 可用于同步主机与特定流，允许其他流继续在该设备上执行。通过调用 cudaStreamDestroy() 可释放流。

时间管理：

下面的代码示例创建了两个事件：

cudaEvent_t start, stop;

cudaEventCreate(&start);

cudaEventCreate(&stop);

这些事件可用于为上一节的代码示例计时，方法如下：

cudaEventRecord(start, 0);

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(inputDev + i * size, inputHost + i * size,

size, cudaMemcpyHostToDevice, stream[i]);

for (int i = 0; i < 2; ++i)

myKernel<<<100, 512, 0, stream[i]>>>

(outputDev + i * size, inputDev + i * size, size);

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(outputHost + i * size, outputDev + i * size,

size, cudaMemcpyDeviceToHost, stream[i]);

cudaEventRecord(stop, 0);

cudaEventSynchronize(stop);

float elapsedTime;

cudaEventElapsedTime(&elapsedTime, start, stop);

cudaEventDestroy(start);

cudaEventDestroy(stop);

纹理参考管理：

在内核使用纹理参考从纹理存储器中读取之前，必须使用 cuTexRefSetAddress() 或 cuTexRefSetArray()将纹理参考绑定到纹理。

如果模块 cuModule 包含定义如下的纹理参考 texRef：

texture<float, 2, cudaReadModeElementType> texRef;

则下面的代码示例将检索 texRef 的句柄：

CUtexref cuTexRef;

cuModuleGetTexRef(&cuTexRef, cuModule, “texRef”);

下面的代码示例将 texRef 绑定到 devPtr 指向的线性存储器：

cuTexRefSetAddress(NULL, cuTexRef, devPtr, size);

下面的代码示例将 texRef 绑定到 CUDA 数组 cuArray：

cuTexRefSetArray(cuTexRef, cuArray, CU_TRSA_OVERRIDE_FORMAT);

参考手册列举了用于设置寻址模式、过滤模式和其他针对纹理参考的标记的各种函数。在将纹理绑定到纹理参考时所指定的格式必须与声明纹理参考时指定的参数相匹配；否则纹理获取的结果将无法确定。

Open GL互操作性

必须使用 cuGLInit() 初始化与 OpenGL 的互操作性。

首先必须将一个缓冲对象注册到 CUDA，之后才能进行映射。可通过 cuGLRegisterBufferObject() 完成：

GLuint bufferObj;

cuGLRegisterBufferObject(bufferObj);

注册完成后，内核即可使用 cuGLMapBufferObject() 返回的设备存储器地址读取或写入缓冲对象：

GLuint bufferObj;

CUdeviceptr devPtr;

int size;

cuGLMapBufferObject(&devPtr, &size, bufferObj);

解除映射是通过 cuGLUnmapBufferObject() 完成的，可使用 cuGLUnregisterBufferObject() 取消注册。

Direct3D 互操作性要求在创建 CUDA 上下文时指定 Direct3D 设备。通过使用 cuD3D9CtxCreate() 而非cuCtxCreate() 创建 CUDA 上下文即可实现此目标。。

随后即可使用 cuD3D9RegisterResource() 将 Direct3D 资源注册到 CUDA：

LPDIRECT3DVERTEXBUFFER9 buffer;

cuD3D9RegisterResource(buffer, CU_D3D9_REGISTER_FLAGS_NONE);

LPDIRECT3DSURFACE9 surface;

cuD3D9RegisterResource(surface, CU_D3D9_REGISTER_FLAGS_NONE);

cuD3D9RegisterResource() 可能具有较高的开销，通常仅为每个资源调用一次。使用 cuD3D9UnregisterVertexBuffer() 可取消注册。

将资源注册到 CUDA 之后，即可在需要时分别使用 cuD3D9MapResources() 和 cuD3D9UnmapResources()任意多次地映射和解除映射。内核可使用 cuD3D9ResourceGetMappedPointer() 返回的设备存储器地址和 cuD3D9ResourceGetMappedSize()、cuD3D9ResourceGetMappedPitch() 及 cuD3D9ResourceGetMappedPitchSlice() 返回的大小和间距信息来读取和写入已映射的资源。通过 Direct3D 访问已映射的资源将导致不确定的结果。

下面的代码示例使用 0 填充了一个缓冲区：

CUdeviceptr devPtr;

cuD3D9ResourceGetMappedPointer(&devPtr, buffer);

size_t size;

cuD3D9ResourceGetMappedSize(&size, buffer);

cuMemset(devPtr, 0, size);

在下面的代码示例中，每个线程都访问大小为 (width, height) 的二维表面的一个像素，像素格式为 float4：

// host code

CUdeviceptr devPtr;

cuD3D9ResourceGetMappedPointer(&devPtr, surface);

size_t pitch;

cuD3D9ResourceGetMappedPitch(&pitch, surface);