Skip to content

Commit e68c78c

Browse files
authored
Merge pull request #22 from mwootton/miopen_integration
Add support for MIOpen
2 parents 633c416 + 7e3319f commit e68c78c

32 files changed

+1839
-11
lines changed

aten/src/ATen/CMakeLists.txt

Lines changed: 11 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,8 @@ FILE(GLOB cuda_cpp "cuda/*.cpp" "cuda/detail/*.cpp")
5050
FILE(GLOB cuda_cu "cuda/*.cu" "cuda/detail/*.cu")
5151
FILE(GLOB cudnn_h "cudnn/*.h" "cudnn/*.cuh")
5252
FILE(GLOB cudnn_cpp "cudnn/*.cpp")
53+
FILE(GLOB miopen_h "miopen/*.h")
54+
FILE(GLOB miopen_cpp "miopen/*.cpp")
5355
FILE(GLOB mkl_cpp "mkl/*.cpp")
5456
FILE(GLOB mkldnn_cpp "mkldnn/*.cpp")
5557

@@ -58,6 +60,7 @@ FILE(GLOB native_sparse_cpp "native/sparse/*.cpp")
5860
FILE(GLOB native_sparse_cuda_cu "native/sparse/cuda/*.cu")
5961
FILE(GLOB native_sparse_cuda_cpp "native/sparse/cuda/*.cpp")
6062
FILE(GLOB native_cudnn_cpp "native/cudnn/*.cpp")
63+
FILE(GLOB native_miopen_cpp "native/miopen/*.cpp")
6164
FILE(GLOB native_cuda_cu "native/cuda/*.cu")
6265
FILE(GLOB native_cuda_cpp "native/cuda/*.cpp")
6366
FILE(GLOB native_mkl_cpp "native/mkl/*.cpp")
@@ -74,9 +77,14 @@ endif()
7477
IF(USE_CUDA OR USE_ROCM)
7578
list(APPEND ATen_CUDA_INCLUDE ${CMAKE_CURRENT_SOURCE_DIR}/cuda)
7679
set(ATen_CUDA_SRCS ${ATen_CUDA_SRCS} ${cuda_cu} ${native_cuda_cu} ${native_sparse_cuda_cu})
77-
set(all_cuda_cpp ${native_cudnn_cpp} ${native_sparse_cuda_cpp} ${cuda_cpp} ${native_cuda_cpp} ${cuda_generated_cpp} ${ATen_CUDA_SRCS})
78-
IF(CUDNN_FOUND)
79-
SET(all_cuda_cpp ${all_cuda_cpp} ${cudnn_cpp})
80+
set(all_cuda_cpp ${native_sparse_cuda_cpp} ${cuda_cpp} ${native_cuda_cpp} ${cuda_generated_cpp} ${ATen_CUDA_SRCS})
81+
IF(USE_CUDA)
82+
SET(all_cuda_cpp ${native_cudnn_cpp} ${native_miopen_cpp} ${all_cuda_cpp})
83+
IF(CUDNN_FOUND)
84+
SET(all_cuda_cpp ${all_cuda_cpp} ${cudnn_cpp})
85+
ENDIF()
86+
ELSEIF(USE_ROCM)
87+
SET(all_cuda_cpp ${native_cudnn_cpp} ${native_miopen_cpp} ${miopen_cpp} ${all_cuda_cpp})
8088
ENDIF()
8189
endif()
8290

aten/src/ATen/cuda/CUDAConfig.h.in

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,3 +5,4 @@
55
// c.f. https://stackoverflow.com/questions/33759787/generating-an-error-if-checked-boolean-macro-is-not-defined
66

77
#define AT_CUDNN_ENABLED() @AT_CUDNN_ENABLED@
8+
#define AT_MIOPEN_ENABLED() @AT_MIOPEN_ENABLED@

aten/src/ATen/cuda/detail/CUDAHooks.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -119,6 +119,10 @@ bool CUDAHooks::compiledWithCuDNN() const {
119119
return AT_CUDNN_ENABLED();
120120
}
121121

122+
bool CUDAHooks::compiledWithMIOpen() const {
123+
return AT_MIOPEN_ENABLED();
124+
}
125+
122126
bool CUDAHooks::supportsDilatedConvolutionWithCuDNN() const {
123127
#if AT_CUDNN_ENABLED()
124128
cudaDeviceProp* prop =

aten/src/ATen/cuda/detail/CUDAHooks.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,7 @@ struct CUDAHooks : public at::CUDAHooksInterface {
1818
Allocator* getPinnedMemoryAllocator() const override;
1919
void registerCUDATypes(Context*) const override;
2020
bool compiledWithCuDNN() const override;
21+
bool compiledWithMIOpen() const override;
2122
bool supportsDilatedConvolutionWithCuDNN() const override;
2223
long versionCuDNN() const override;
2324
double batchnormMinEpsilonCuDNN() const override;

aten/src/ATen/cudnn/Descriptors.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -339,7 +339,7 @@ union Constant
339339
double d;
340340
Constant(cudnnDataType_t dataType, double value) {
341341
if (dataType == CUDNN_DATA_HALF || dataType == CUDNN_DATA_FLOAT) {
342-
f = (float) value;
342+
f = static_cast<float>(value);
343343
} else {
344344
d = value;
345345
}

aten/src/ATen/cudnn/Handles.cpp renamed to aten/src/ATen/cudnn/Handle.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
#include "Handles.h"
1+
#include "Handle.h"
22

33
#include "ATen/cuda/Exceptions.h"
44

File renamed without changes.

aten/src/ATen/cudnn/Utils.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@
44
#include "ATen/cuda/Exceptions.h"
55
#include "THC/THC.h"
66
#include "cudnn-wrapper.h"
7-
#include "Handles.h"
7+
#include "Handle.h"
88

99
namespace at { namespace native {
1010

aten/src/ATen/detail/CUDAHooksInterface.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -85,6 +85,10 @@ struct AT_API CUDAHooksInterface {
8585
return false;
8686
}
8787

88+
virtual bool compiledWithMIOpen() const {
89+
return false;
90+
}
91+
8892
virtual bool supportsDilatedConvolutionWithCuDNN() const {
8993
return false;
9094
}
Lines changed: 116 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,116 @@
1+
#include "Descriptors.h"
2+
#include <ATen/ATen.h>
3+
4+
namespace at { namespace native {
5+
6+
namespace {
7+
8+
inline miopenDataType_t getDataType(const at::Type& t) {
9+
auto scalar_type = t.scalarType();
10+
if (scalar_type == at::kFloat) {
11+
return miopenFloat;
12+
} else if (scalar_type == at::kHalf) {
13+
return miopenHalf;
14+
}
15+
throw std::runtime_error("TensorDescriptor only supports float and half tensors");
16+
}
17+
18+
inline miopenDataType_t getDataType(const at::Tensor& t) {
19+
return getDataType(t.type());
20+
}
21+
22+
} // anonymous namespace
23+
24+
25+
void TensorDescriptor::set(const at::Tensor &t, size_t pad) {
26+
set(getDataType(t), t.sizes(), t.strides(), pad);
27+
}
28+
29+
static int MIOPEN_DIM_MAX = 4;
30+
31+
void TensorDescriptor::set(miopenDataType_t datatype, IntList t_sizes, IntList t_strides, size_t pad) {
32+
size_t dim = t_sizes.size();
33+
if (dim > MIOPEN_DIM_MAX || pad > MIOPEN_DIM_MAX)
34+
#define _STR(X) #X
35+
#define STR(X) _STR(X)
36+
throw std::runtime_error("MIOpen supports only up to " STR(MIOPEN_DIM_MAX) " dimensions");
37+
#undef _STR
38+
#undef STR
39+
int size[MIOPEN_DIM_MAX];
40+
int stride[MIOPEN_DIM_MAX];
41+
for (size_t i = 0; i < dim; ++i) {
42+
size[i] = static_cast<int>(t_sizes[i]);
43+
stride[i] = static_cast<int>(t_strides[i]);
44+
}
45+
for (size_t i = dim; i < pad; ++i) {
46+
size[i] = 1;
47+
stride[i] = 1;
48+
}
49+
set(datatype, static_cast<int>(std::max(dim, pad)), size, stride);
50+
}
51+
52+
std::string miopenTypeToString(miopenDataType_t dtype) {
53+
switch (dtype) {
54+
case miopenFloat:
55+
return "miopenFloat";
56+
case miopenHalf:
57+
return "miopenHalf";
58+
default:
59+
std::ostringstream oss;
60+
oss << "(unknown data-type " << static_cast<int>(dtype) << ")";
61+
return oss.str();
62+
}
63+
}
64+
65+
std::ostream& operator<<(std::ostream & out, const TensorDescriptor& d) {
66+
out << "TensorDescriptor " << static_cast<void*>(d.desc()) << "\n";
67+
int nbDims = 4;
68+
int dimA[MIOPEN_DIM_MAX];
69+
int strideA[MIOPEN_DIM_MAX];
70+
miopenDataType_t dtype;
71+
miopenGetTensorDescriptor(d.desc(), &dtype, dimA, strideA);
72+
out << " type = " << miopenTypeToString(dtype) << "\n";
73+
out << " nbDims = " << nbDims << "\n";
74+
// Read out only nbDims of the arrays!
75+
out << " dimA = ";
76+
for (auto i : ArrayRef<int>{dimA, static_cast<size_t>(nbDims)}) {
77+
out << i << ", ";
78+
}
79+
out << "\n";
80+
out << " strideA = ";
81+
for (auto i : ArrayRef<int>{strideA, static_cast<size_t>(nbDims)}) {
82+
out << i << ", ";
83+
}
84+
out << "\n";
85+
return out;
86+
}
87+
88+
void TensorDescriptor::print() { std::cout << *this; }
89+
90+
void FilterDescriptor::set(const at::Tensor &t, int64_t pad) {
91+
auto dim = t.ndimension();
92+
if (dim > MIOPEN_DIM_MAX || pad > MIOPEN_DIM_MAX)
93+
#define _STR(X) #X
94+
#define STR(X) _STR(X)
95+
throw std::runtime_error("MIOpen supports only up to " STR(MIOPEN_DIM_MAX) " dimensions");
96+
#undef _STR
97+
#undef STR
98+
if (!t.is_contiguous()) {
99+
throw std::runtime_error("MIOpen filters (a.k.a. weights) must be contiguous");
100+
}
101+
int size[MIOPEN_DIM_MAX];
102+
int stride[MIOPEN_DIM_MAX];
103+
for (int i = 0; i < dim; ++i) {
104+
size[i] = (int) t.size(i);
105+
}
106+
for (int i = dim; i < pad; ++i) {
107+
size[i] = (int) 1;
108+
}
109+
for (int i = dim - 1; i >=0; --i) {
110+
stride[i] = (i == dim - 1) ? 1 : stride[i+1] * size[i+1];
111+
}
112+
dim = std::max(dim, pad);
113+
set(getDataType(t), (int) dim, size, stride);
114+
}
115+
116+
}}

0 commit comments

Comments
 (0)