[CUDA] 同一个函数根据运行设备选择不同的运行方式

一、简介

有时我们需要在host和device端实现同样的一个函数。
一个简单的方法是分别在host和device端上实现两个同名函数，例如fun_host()和fun_device()，在host和device上分别使用不同的函数。
另一种方法只使用一个函数fun()，在编译时根据运行目标的不同选择使用不同的实现方式。该方法可以使用宏 __CUDA__ARCH__实现，本文给出了该方法的实现示例。

二、代码示例

main.cu程序代码：

#include <math.h>
#include <stdio.h>
#include <cuda_runtime.h>
__host__ __device__ inline void fun()
{
#ifdef __CUDA_ARCH__
    float a = fdividef(1.0f, 3.0f);
    printf("In device, result is %f\n", a);
#else
    float a = 1.0f / 3.0f;
    printf("In host, result is %f\n", a);
#endif
};

__global__ void f()
{
    fun();
}

int main()
{
    fun();         // in host
    f<<<1, 1>>>(); // in device
    cudaDeviceSynchronize();
    return 0;
}

CMakeLists.txt文件内容：

cmake_minimum_required(VERSION 3.0.0)
project(HelloWorld VERSION 0.1.0 LANGUAGES C CXX)
find_package(CUDA REQUIRED)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --std=c++11")
cuda_add_executable(HelloWorld main.cu)

运行结果：

In host, result is 0.333333
In device, result is 0.333333