▶ 简单的将纯 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();
const unsigned int size = ;
float hVector[size], hResultVector[size], *dVector;
for (unsigned int i = ; 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 << <, >>>(dVector, hFunctionPtr, size);
cudaMemcpyFromSymbol(&hFunctionPtr, dDivideByTwoPtr, sizeof(deviceFunc));
transformVector << <, >> > (dVector, hFunctionPtr, size); cudaMemcpy(hResultVector, dVector, sizeof(float) * size, cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
if (dVector)
cudaFree(dVector);
for (int i = ; i < size; ++i)
{
if (fabs(hVector[i] - hResultVector[i]) > EPS)
{
printf("\nError at i == %d, hVector[i] == %f, hResultVector[i] == %f", i, hVector[i], hResultVector[i]);
return ;
}
}
return ;
} int main()
{
printf("\n\tStart.\n");
printf("\n\tFinish: %s\n", test() ? "Pass" : "Fail");
getchar();
return ;
}
● 输出结果：
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(), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost));
// 从指定符号 symbol 处偏移 offset 字节处，拷贝 count 字节到 dst，默认模式为设备拷到主机