CUDA 如何处理 OpenCL 中的线程偏移？

Question

例如，我可以在 OpenCL 中启动带有线程 id 偏移量的内核，因此当使用多个 GPU 时，第二个 GPU 线程可以直接从任意整数值开始：

GPU1: thread idx 0 to k
GPU2: thread idx k+1 to N

所有 GPU 的内核看起来都完全相同：

__kernel(..)
{
   int threadId=get_global_id(0); // this starts from k+1 for GPU2
   ..
}

在驱动 API 中调用内核时，在 CUDA 中什么是等效的？

OpenCL 有

cl_int clEnqueueNDRangeKernel ( cl_command_queue command_queue,
    cl_kernel kernel,
    cl_uint work_dim,
    const size_t *global_work_offset, // this can be given k+1
    const size_t *global_work_size,
    const size_t *local_work_size,
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list,
    cl_event *event)

我是否必须自己在内核中嵌入偏移量，例如下面的示例才能在多 GPU 上使用外观最相似的内核？

__global__ void vecAdd(.., int * gpuParams)
{
    int offset=gpuParams[0];
    int threadId = offset + blockIdx.x * blockDim.x + threadIdx.x;
}

（这是用于扁平缓冲区访问和内核）

这个偏移量会在每次内核启动之前动态改变，所以我猜我不能使用#define。

Cuda 驱动 api 有这个用于启动内核：

CUresult cuLaunchKernel ( CUfunction f, unsigned int  gridDimX, 
                unsigned int  gridDimY, unsigned int  gridDimZ, 
                unsigned int  blockDimX, unsigned int  blockDimY, 
                unsigned int  blockDimZ, unsigned int  sharedMemBytes, 
                CUstream hStream, void** kernelParams, void** extra )

Answer 1

CUDA 中没有与之等效的功能，但您也不需要它。

在标准 CUDA 模型中，您无法通过单个 API 调用在多个设备上启动内核（从技术上讲，您也无法通过单个 API 调用在多个设备上分配内存）。每个 GPU 都是有效独立的。

NVIDIA 已引入多设备启动和同步方法作为 CUDA 9 中新的 cooperative groups 功能的一部分，但这些与您所询问的完全不同。