Skip to content
Projects
Groups
Snippets
Help
Loading...
Sign in / Register
Toggle navigation
N
ngraph
Project
Project
Details
Activity
Cycle Analytics
Repository
Repository
Files
Commits
Branches
Tags
Contributors
Graph
Compare
Charts
Issues
0
Issues
0
List
Board
Labels
Milestones
Merge Requests
0
Merge Requests
0
CI / CD
CI / CD
Pipelines
Jobs
Schedules
Charts
Packages
Packages
Wiki
Wiki
Snippets
Snippets
Members
Members
Collapse sidebar
Close sidebar
Activity
Graph
Charts
Create a new issue
Jobs
Commits
Issue Boards
Open sidebar
submodule
ngraph
Commits
b6c34bd3
Commit
b6c34bd3
authored
Feb 09, 2018
by
Fenglei Tian
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
test on cuda kernel
parent
144d790b
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
159 additions
and
41 deletions
+159
-41
gpu_cuda_kernel_emitters.cpp
src/ngraph/runtime/gpu/gpu_cuda_kernel_emitters.cpp
+115
-2
gpu_cuda_kernel_emitters.hpp
src/ngraph/runtime/gpu/gpu_cuda_kernel_emitters.hpp
+44
-39
No files found.
src/ngraph/runtime/gpu/gpu_cuda_kernel_emitters.cpp
View file @
b6c34bd3
...
...
@@ -14,8 +14,55 @@
#include <algorithm>
#include <map>
#include "ngraph/runtime/gpu/gpu_kernel_emitters.hpp"
#include "ngraph/runtime/gpu/gpu_
cuda_
kernel_emitters.hpp"
#include <nvrtc.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>
#include <cudnn_v7.h>
#include "ngraph/node.hpp"
#include "ngraph/ops/broadcast.hpp"
#include "ngraph/ops/concatenate.hpp"
#include "ngraph/ops/constant.hpp"
#include "ngraph/ops/convolution.hpp"
#include "ngraph/ops/dot.hpp"
#include "ngraph/ops/function_call.hpp"
#include "ngraph/ops/get_output_element.hpp"
#include "ngraph/ops/max_pool.hpp"
#include "ngraph/ops/one_hot.hpp"
#include "ngraph/ops/reduce.hpp"
#include "ngraph/ops/replace_slice.hpp"
#include "ngraph/ops/reshape.hpp"
#include "ngraph/ops/reverse.hpp"
#include "ngraph/ops/slice.hpp"
#include "ngraph/ops/sum.hpp"
#include "ngraph/util.hpp"
#define NVRTC_SAFE_CALL(x) \
do
{
\
nvrtcResult
result
=
x
;
\
if
(
result
!=
NVRTC_SUCCESS
)
{
\
std
::
cerr
<<
"
\n
error: "
#
x
" failed with error "
\
<<
nvrtcGetErrorString
(
result
)
<<
'\n'
;
\
exit
(
1
);
\
}
\
}
while
(
0
)
#define CUDA_SAFE_CALL(x) \
do
{
\
CUresult
result
=
x
;
\
if
(
result
!=
CUDA_SUCCESS
)
{
\
const
char
*
msg
;
\
cuGetErrorName
(
result
,
&
msg
);
\
std
::
cerr
<<
"
\n
error: "
#
x
" failed with error "
\
<<
msg
<<
'\n'
;
\
exit
(
1
);
\
}
\
}
while
(
0
)
namespace
ngraph
{
...
...
@@ -28,9 +75,75 @@ namespace ngraph
namespace
kernel
{
void
emit_abs
(
float
*
in
,
float
*
out
,
size_t
count
)
void
emit_abs
(
void
*
in
,
void
*
out
,
size_t
count
)
{
const
char
*
op_abs
=
R"(
extern "C" __global__
void cuda_op_abs(float* in, float* out, size_t n)
{
size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid < n)
{
out[tid] = fabsf(in[tid]);
}
})"
;
size_t
numBlocks
=
4
;
size_t
numThreads
=
4
;
// Create an instance of nvrtcProgram with the code string.
nvrtcProgram
prog
;
NVRTC_SAFE_CALL
(
nvrtcCreateProgram
(
&
prog
,
// prog i
op_abs
,
// buffer
"op_abs.cu"
,
// name
0
,
// numHeaders
NULL
,
// headers
NULL
));
// includeNames
const
char
*
opts
[]
=
{
"--gpu-architecture=compute_35"
,
"--relocatable-device-code=true"
};
nvrtcResult
compileResult
=
nvrtcCompileProgram
(
prog
,
// prog
2
,
// numOptions
opts
);
// options
// Obtain compilation log from the program.
size_t
logSize
;
NVRTC_SAFE_CALL
(
nvrtcGetProgramLogSize
(
prog
,
&
logSize
));
char
*
log
=
new
char
[
logSize
];
NVRTC_SAFE_CALL
(
nvrtcGetProgramLog
(
prog
,
log
));
std
::
cout
<<
log
<<
'\n'
;
delete
[]
log
;
if
(
compileResult
!=
NVRTC_SUCCESS
)
{
exit
(
1
);
}
size_t
ptxSize
;
NVRTC_SAFE_CALL
(
nvrtcGetPTXSize
(
prog
,
&
ptxSize
));
char
*
ptx
=
new
char
[
ptxSize
];
NVRTC_SAFE_CALL
(
nvrtcGetPTX
(
prog
,
ptx
));
// Destroy the program.
NVRTC_SAFE_CALL
(
nvrtcDestroyProgram
(
&
prog
));
// Load the generated PTX and get a handle to the parent kernel.
CUdevice
cuDevice
;
CUcontext
context
;
CUmodule
module
;
CUfunction
cuda_op_abs_kernel
;
CUDA_SAFE_CALL
(
cuInit
(
0
));
CUDA_SAFE_CALL
(
cuDeviceGet
(
&
cuDevice
,
0
));
CUDA_SAFE_CALL
(
cuCtxCreate
(
&
context
,
0
,
cuDevice
));
CUDA_SAFE_CALL
(
cuModuleLoadDataEx
(
&
module
,
ptx
,
0
,
0
,
0
));
CUDA_SAFE_CALL
(
cuModuleGetFunction
(
&
cuda_op_abs_kernel
,
module
,
"cuda_op_abs"
));
void
*
argsList
[]
=
{
In
,
Out
,
&
count
};
CUDA_SAFE_CALL
(
cuLaunchKernel
(
cuda_op_abs_kernel
,
count
,
1
,
1
,
// grid dim
1
,
1
,
1
,
// block dim
0
,
NULL
,
// shared mem and stream
argsList
,
0
));
// arguments
CUDA_SAFE_CALL
(
cuCtxSynchronize
());
// Retrieve and print output.
}
void
emit_broadcast
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
...
...
src/ngraph/runtime/gpu/gpu_cuda_kernel_emitters.hpp
View file @
b6c34bd3
...
...
@@ -23,56 +23,61 @@ namespace ngraph
{
namespace
gpu
{
namespace
kernel
namespace
cuda
{
void
emit_broadcast
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
string
&
arg0
,
// replacement context
const
std
::
string
&
out
,
const
Shape
&
arg0_shape
,
const
Shape
&
out_shape
,
const
AxisSet
&
broadcast_axes
);
void
emit_concat
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
vector
<
std
::
string
>&
args
,
const
std
::
string
&
out
,
const
std
::
vector
<
Shape
>&
in_shapes
,
const
Shape
&
out_shape
,
const
size_t
concatenation_axis
);
namespace
kernel
{
void
emit_abs
(
void
**
in
,
void
**
out
,
size_t
count
);
void
emit_replace_slice
(
codegen
::
CodeWriter
&
writer
,
void
emit_broadcast
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
string
&
arg0
,
// replacement context
const
std
::
string
&
arg1
,
// replacement value
const
std
::
string
&
out
,
const
Shape
&
arg
1
_shape
,
const
Shape
&
arg
0
_shape
,
const
Shape
&
out_shape
,
const
Coordinate
&
lower_bounds
,
const
Coordinate
&
upper_bounds
,
const
Strides
&
strides
);
void
emit_slice
(
codegen
::
CodeWriter
&
writer
,
const
AxisSet
&
broadcast_axes
);
void
emit_concat
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
vector
<
std
::
string
>&
args
,
const
std
::
string
&
out
,
const
std
::
vector
<
Shape
>&
in_shapes
,
const
Shape
&
out_shape
,
const
size_t
concatenation_axis
);
void
emit_replace_slice
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
string
&
arg0
,
// replacement context
const
std
::
string
&
arg1
,
// replacement value
const
std
::
string
&
out
,
const
Shape
&
arg1_shape
,
const
Shape
&
out_shape
,
const
Coordinate
&
lower_bounds
,
const
Coordinate
&
upper_bounds
,
const
Strides
&
strides
);
void
emit_slice
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
string
&
arg0
,
// replacement context
const
std
::
string
&
out
,
const
Shape
&
arg0_shape
,
const
Shape
&
out_shape
,
const
Coordinate
&
lower_bounds
,
const
Coordinate
&
upper_bounds
,
const
Strides
&
strides
);
void
emit_reshape
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
string
&
arg0
,
// replacement context
const
std
::
string
&
out
,
const
Shape
&
arg0_shape
,
const
Shape
&
out_shape
,
const
AxisVector
&
arg0_axis_order
);
void
emit_sum
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
string
&
arg0
,
// replacement context
const
std
::
string
&
out
,
const
Shape
&
arg0_shape
,
const
Shape
&
out_shape
,
const
Coordinate
&
lower_bounds
,
const
Coordinate
&
upper_bounds
,
const
Strides
&
strides
);
void
emit_reshape
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
string
&
arg0
,
// replacement context
const
std
::
string
&
out
,
const
Shape
&
arg0_shape
,
const
Shape
&
out_shape
,
const
AxisVector
&
arg0_axis_order
);
void
emit_sum
(
codegen
::
CodeWriter
&
writer
,
const
std
::
string
&
element_type
,
const
std
::
string
&
arg0
,
// replacement context
const
std
::
string
&
out
,
const
Shape
&
arg0_shape
,
const
Shape
&
out_shape
,
const
AxisSet
&
reduction_axes
);
const
AxisSet
&
reduction_axes
);
}
}
}
}
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment