7. 编程接口¶

7.1. 内存管理（Memory Management）¶

内存管理系列函数用于分配和释放设备内存，访问在全局内存空间中声明的变量的内存，并在主机和设备内存之间传输数据。

线性内存的分配和释放使用topsMalloc()进行分配，并使用topsFree()进行释放。

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

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

topsGetSymbolAddress()用于获取指向在全局内存空间中声明的变量分配的内存地址。通过topsGetSymbolSize()可以获取分配内存的大小。

下表中列出了所有使用topsMalloc()函数分配的线性内存和数组之间进行内存拷贝的各种函数。

以下示例代码将主机内存中的数组数据拷贝到设备内存：

float data[256];
int size = sizeof(data);
void *devPtr;
topsMalloc(&devPtr, size);
topsMemcpy(devPtr, data, size, topsMemcpyHostToDevice);

以下示例代码用主机内存中的数组数据初始化常量内存：

__constant__ float constData[256];
float data[256];
topsMemcpyToSymbol(constData, data, sizeof(data));

topsHostMalloc()和topsHostFree()用于分配和释放页锁定（page-locked）的主机内存。相比使用malloc()分配的常规页内存（pageable memory），页锁定的主机内存在与设备内存进行数据传输时具有更高的带宽。需要说明的是，页锁定的主机内存是一种相对稀缺的资源，相比常规页内存分配更容易产生分配失败错误。此外，通过减少操作系统用于页面调度的可用物理内存量，分配过多的页锁定主机内存会降低整体系统性能。

topsHostRegister()用于将malloc()分配的常规页内存进行页锁定操作，从而成为页锁定的主机内存。topsHostUnregister()用于反向解除这种页锁定操作。

7.2. 执行流（Stream）¶

执行流管理函数用于创建和销毁执行流，以及确定执行流中的所有操作是否已完成。

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

topsStream_t stream[2];
for (int i = 0; i < 2; ++i)
    topsStreamCreate(&stream[i]);

每个执行流可以通过以下代码示例定义为以下执行序列：从主机到设备的内存拷贝、内核函数启动、以及从设备到主机的内存拷贝：

for (int i = 0; i < 2; ++i)
    topsMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, topsMemcpyHostToDevice, 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)
    topsMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, topsMemcpyDeviceToHost, stream[i]);
topsDeviceSynchronize();

每一个执行流从输入数组hostPtr拷贝对应部分数据到设备内存中的数组inputDevPtr，通过调用myKernel()函数在设备上处理intputDevPtr，然后将结果outputDevPtr再次复制回hostPtr的对应部分。

示例中使用了两个执行流来处理hostPtr，为了使一个执行流的内存拷贝与另一个执行流的内核执行重叠，hostPtr必须是指向页锁定（page-locked）的主机内存：

float* hostPtr;
topsHostMalloc((void**)&hostPtr, 2 * size);

主机端调用topsDeviceSynchronize()函数可以确保所有执行流中的所有操作已完成。

7.3. 事件（Event）¶

事件管理函数用于创建、记录和销毁时间，并查询两个事件之间的经过时间。

以下示例代码创建两个事件：

topsEvent_t start, stop;
topsEventCreate(&start);
topsEventCreate(&stop);

这些创建的事件可以按照下述代码示例来度量执行流上执行操作的执行时间：

topsEventRecord(start, 0);
for (int i = 0; i < 2; ++i)
    topsMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, topsMemcpyHostToDevice, 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)
    topsMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, topsMemcpyDeviceToHost, stream[i]);
topsEventRecord(stop, 0);
topsEventSynchronize(stop);
float elapsedTime;
topsEventElapsedTime(&elapsedTime, start, stop);

7.4. 默认执行流（Default Stream）¶

在执行核函数启动和数据拷贝操作时，如果没有指定执行流参数或者执行流参数为空，该操作会下发到默认执行流中。运行时库为每个设备维护一个隐式的默认执行流，进程内的多个线程共享该默认执行流，不同设备的默认执行流在执行序上无任何依赖关系，可以并发执行。在同一设备上，默认执行流同常规执行流有额外同步语义，即默认执行流和常规执行流会构建隐式依赖关系。特别地，如果使用topsStreamNonBlocking标志创建常规执行流，该常规执行流与默认执行流在执行上不会构建隐式依赖关系。当前运行时库实现还未完全支持这些同步规则。建议开发者不要使用默认执行流，容易产生故障。