Skip to content

[WIP] Enable disabled functions for ROCm #252

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 10 commits into from
Oct 9, 2018
Merged

[WIP] Enable disabled functions for ROCm #252

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
b6e525d
rocRAND brings Poisson RNG - don't disable functions
iotamudelta Oct 4, 2018
9a09103
Merge branch 'master' into enable_disabled
iotamudelta Oct 8, 2018
11ada85
Should also only need philox which rocRAND now has.
iotamudelta Oct 8, 2018
0924383
19:27:19 /var/lib/jenkins/workspace/aten/src/ATen/native/cuda/Distrib…
iotamudelta Oct 8, 2018
6407c91
At the very least the WARP_SIZE is wrong for ROCm.
iotamudelta Oct 8, 2018
3578393
Do not disable renorm kernel
iotamudelta Oct 8, 2018
eee4e0f
Put the ifdef directly in to ensure better maintainability.
iotamudelta Oct 9, 2018
5088c20
scalar_cast does not exist any longer.
iotamudelta Oct 9, 2018
ce8162c
Implement the one ifdef directly, CUDA_VERSION check is not in the file
iotamudelta Oct 9, 2018
0805856
Merge branch 'master' into enable_disabled
iotamudelta Oct 9, 2018
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion aten/src/THC/THCNumerics.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -209,7 +209,7 @@ struct THCNumerics<at::Half> {
static inline __host__ __device__ at::Half round(at::Half a) { return ::round(a); }

static inline __host__ __device__ at::Half frac(at::Half a) {
#ifdef __CUDA_ARCH__
#if defined(__CUDA_ARCH__) || defined(__HIP_PLATFORM_HCC__)
return a - ::trunc(a);
#else // __CUDA_ARCH__
return a - ::floor(a);
Expand Down
4 changes: 4 additions & 0 deletions aten/src/THCUNN/LookupTable.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,11 @@
#include "THCTensorSort.cuh"
#include "../THC/THCTensorMathReduce.cuh"

#ifdef __HIP_PLATFORM_HCC__
const int WARP_SIZE = 64;
#else
const int WARP_SIZE = 32;
#endif

template
<typename Dtype,
Expand Down
27 changes: 0 additions & 27 deletions tools/amd_build/disabled_features.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,31 +32,12 @@
"_OPENMP": "_OPENMP_STUB"
}
},
{
"path": "aten/src/ATen/Context.cpp",
"s_constants": {
"#ifdef USE_SSE3": "#if defined(USE_SSE3) && !defined(__HIP_DEVICE_COMPILE__)"
}
},
{
"path": "aten/src/ATen/native/Distributions.h",
"s_constants": {
"scalar_cast": "static_cast"
}
},
{
"path": "aten/src/ATen/native/cuda/Distributions.cu",
"s_constants": {
"#include <nvfunctional>": ""
}
},
{
"path": "aten/src/THC/THCNumerics.cuh",
"s_constants": {
"#ifdef __CUDA_ARCH__": "#if defined(__CUDA_ARCH__) || defined(__HIP_PLATFORM_HCC__)",
"#if CUDA_VERSION < 9000": "#if CUDA_VERSION < 9000 && !defined(__HIP_PLATFORM_HCC__)"
}
},
{
"path": "aten/src/ATen/native/cuda/RoiPooling.cu",
"s_constants": {
Expand Down Expand Up @@ -144,8 +125,6 @@
{
"path": "aten/src/ATen/native/cuda/Distributions.cu",
"functions": [
"_s_poisson_cuda",
"poisson_cuda_kernel",
"gamma_cuda_kernel",
"gamma_grad_cuda_kernel",
]
Expand All @@ -164,12 +143,6 @@
"THNN_(LookupTable_renorm)"
]
},
{
"path": "aten/src/THCUNN/LookupTable.cu",
"functions": [
"calculate_norms_and_renorm"
]
},
{
"path": "aten/src/THC/generic/THCTensor.cu",
"functions": [
Expand Down