blob: d96cc7c8881916b65942613df9a3acb7aae26242 (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
|
#include <iostream>
#include "file1.h"
#include "file2.h"
#ifdef _WIN32
#define EXPORT __declspec(dllexport)
#define IMPORT __declspec(dllimport)
#else
#define EXPORT
#define IMPORT
#endif
result_type __device__ file1_func(int x);
result_type_dynamic __device__ file2_func(int x);
IMPORT void __host__ cuda_dynamic_lib_func();
static
__global__
void mixed_kernel(result_type* r, int x)
{
*r = file1_func(x);
result_type_dynamic rd = file2_func(x);
}
EXPORT int mixed_launch_kernel(int x)
{
cuda_dynamic_lib_func();
result_type* r;
cudaError_t err = cudaMallocManaged(&r, sizeof(result_type));
if(err != cudaSuccess)
{
std::cerr << "mixed_launch_kernel: cudaMallocManaged failed: "
<< cudaGetErrorString(err) << std::endl;
return x;
}
mixed_kernel <<<1,1>>> (r,x);
err = cudaGetLastError();
if(err != cudaSuccess)
{
std::cerr << "mixed_kernel [SYNC] failed: "
<< cudaGetErrorString(err) << std::endl;
return x;
}
err = cudaDeviceSynchronize();
if(err != cudaSuccess)
{
std::cerr << "mixed_kernel [ASYNC] failed: "
<< cudaGetErrorString(cudaGetLastError()) << std::endl;
return x;
}
int result = r->sum;
err = cudaFree(r);
if(err != cudaSuccess)
{
std::cerr << "mixed_launch_kernel: cudaFree failed: "
<< cudaGetErrorString(err) << std::endl;
return x;
}
return result;
}
|
