0_Simple__simpleSeparateCompilation

▶ 简单的将纯 C/C++ 函数放到另一个文件中，利用头文件引用到主体 .cu 中来，编译时共同编译。

▶ 源代码，把 C++ 的部分去掉了

// simpleDeviceLibrary.cuh
#ifndef SIMPLE_DEVICE_LIBRARY_CUH
#define SIMPLE_DEVICE_LIBRARY_CUH

extern "C" __device__ float multiplyByTwo(float number);

extern "C" __device__ float divideByTwo(float number);

#endif

// simpleDeviceLibrary.cu
#include <cuda_runtime.h>

extern "C" __device__ float multiplyByTwo(float number) 
{ 
    return number * 2.0f;
}

extern "C" __device__ float divideByTwo(float number) 
{ 
    return number * 0.5f; 
}

// simpleSeparateCompilation.cu
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include "simpleDeviceLibrary.cuh"

#define EPS 1e-5

typedef float(*deviceFunc)(float);
__device__ deviceFunc dMultiplyByTwoPtr = multiplyByTwo;    // 本地声明，直接在代码中调用 multiplyByTwo / divideByTwo 会导致运行时错误
__device__ deviceFunc dDivideByTwoPtr = divideByTwo;

__global__ void transformVector(float *v, deviceFunc f, unsigned int size)
{
    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < size)    
        v[tid] = (*f)(v[tid]);
}

int test()
{
    cudaSetDevice(0);
    const unsigned int size = 1000;
    float hVector[size], hResultVector[size], *dVector;
    for (unsigned int i = 0; i < size; ++i)
    {
        hVector[i] = rand() / (float)RAND_MAX;
        hResultVector[i] = 0.0f;
    }
    cudaMalloc((void **)&dVector, size * sizeof(float));
    cudaMemcpy(dVector, hVector, sizeof(float) * size, cudaMemcpyHostToDevice);

    deviceFunc hFunctionPtr;                                                    // 作为调用参数的函数指针
    cudaMemcpyFromSymbol(&hFunctionPtr, dMultiplyByTwoPtr, sizeof(deviceFunc)); // 给 hFunctionPtr 一个地址，方便调用
    transformVector << <1, 1024 >>>(dVector, hFunctionPtr, size);
    cudaMemcpyFromSymbol(&hFunctionPtr, dDivideByTwoPtr, sizeof(deviceFunc));
    transformVector << <1, 1024 >> > (dVector, hFunctionPtr, size); 
    
    cudaMemcpy(hResultVector, dVector, sizeof(float) * size, cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    if (dVector)
        cudaFree(dVector);
    for (int i = 0; i < size; ++i)
    {
        if (fabs(hVector[i] - hResultVector[i]) > EPS)
        {
            printf("
Error at i == %d, hVector[i] == %f, hResultVector[i] == %f", i, hVector[i], hResultVector[i]);
            return 0;
        }
    }    
    return 1;
}

int main()
{
    printf("
	Start.
");
    printf("
	Finish: %s
", test() ? "Pass" : "Fail");
    getchar();
    return 0;
}

● 输出结果：

        Start.

        Finish: Pass

▶ 涨姿势

// cuda_runtime_api.h
#define __dv(v) 
        = v

extern __host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const void *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost));
    // 从指定符号 symbol 处偏移 offset 字节处，拷贝 count 字节到 dst，默认模式为设备拷到主机