0_Simple__simpleSeparateCompilation

把代码文件和主程序文件分开编译，使用头文件的形式进行引用。

▶ 源代码

// simpleDeviceLibrary.cuh
#ifndef SIMPLE_DEVICE_LIBRARY_CUH
#define SIMPLE_DEVICE_LIBRARY_CUH

extern __device__ float multiplyByTwo(float number);

extern __device__ float divideByTwo(float number);

#endif

// simpleDeviceLibrary.cu
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
__device__ float multiplyByTwo(float number)
{
    return number * 2.0f;
}

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

#include <stdio.h>
#include <iostream>
#include <vector>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include "simpleDeviceLibrary.cuh"

using std::cout;
using std::endl;
using std::vector;

#define EPS 1e-5

typedef float(*deviceFunc)(float);

__device__ deviceFunc dMultiplyByTwoPtr = multiplyByTwo;
__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]);
}

bool test()
{
    bool result = true;
    cudaSetDevice(0);

    const unsigned int kVectorSize = 1000;
    vector<float> hVector(kVectorSize);
    for (unsigned int i = 0; i < kVectorSize; ++i)
        hVector[i] = rand() / static_cast<float>(RAND_MAX);
    float *dVector;
    cudaMalloc(&dVector, kVectorSize * sizeof(float));
    cudaMemcpy(dVector, &hVector[0], kVectorSize * sizeof(float), cudaMemcpyHostToDevice);

    dim3 dimGrid(1);
    dim3 dimBlock(1024);

    // 函数指针需要用 cudaMemcpyFromSymbol 放入设备常量内存
    deviceFunc hFunctionPtr;
    cudaMemcpyFromSymbol(&hFunctionPtr, dMultiplyByTwoPtr, sizeof(deviceFunc));
    transformVector << <dimGrid, dimBlock >> > (dVector, hFunctionPtr, kVectorSize);
    cudaGetLastError();

    cudaMemcpyFromSymbol(&hFunctionPtr, dDivideByTwoPtr, sizeof(deviceFunc));
    transformVector << <dimGrid, dimBlock >> > (dVector, hFunctionPtr, kVectorSize);
    cudaGetLastError();

    vector<float> hResultVector(kVectorSize);
    cudaMemcpy(&hResultVector[0], dVector, kVectorSize * sizeof(float), cudaMemcpyDeviceToHost);

    // 检查结果
    for (int i = 0; i < kVectorSize; ++i)
    {
        if (fabs(hVector[i] - hResultVector[i]) > EPS)
        {
            printf("
	Error at %d, gpu[i] = %f, cpu[i] = %f
", i, hResultVector[i], hVector[i]);
            result = false;
            break;
        }
    }
    return result;
}

int main(int argc, char **argv)
{
    printf("
	Start
");
    printf("
	Finish, %s
", test() ? "Passed" : "Failed");

    getchar();
    return 0;
}

▶ 输出结果

　　未测试

▶ 涨姿势

●写在其他 .cpp 文件中的设备函数，需要用函数 cudaMemcpyFromSymbol() 放入设备常量内存才能使用。

typedef float(*deviceFunc)(float);
deviceFunc hFunctionPtr;
cudaMemcpyFromSymbol(&hFunctionPtr, dMultiplyByTwoPtr, sizeof(deviceFunc));