Skip to content

Update OpenCL to 3.0. Add checks around tests that use doubles. #3338

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
umar456 merged 17 commits into from
Jan 21, 2023
Merged

Update OpenCL to 3.0. Add checks around tests that use doubles. #3338

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
17 commits
Select commit Hold shift + click to select a range
c85e600
Add type checks for tests and different types
umar456 Nov 28, 2022
ef2e6ec
Fix fftconvolve so that floats are used for complex float values
umar456 Nov 29, 2022
cf6ebf4
Add ifdef check around powll and powul functions in jit.cl
umar456 Nov 29, 2022
0fb43d6
Update cl2hpp tag and disable building cl2hpp if found on system
umar456 Nov 29, 2022
b19a070
Update the minimum required OpenCL version to 3.0
umar456 Nov 30, 2022
b89ea4c
Fix some errors due to pow in CUDA code with fast math
umar456 Nov 30, 2022
ea34879
Add OpenCL version def to af/opencl.h. Remove FindOpenCL from test
umar456 Dec 1, 2022
d01a25e
Allow OpenCL C device version checks on older platforms
umar456 Dec 2, 2022
27ab2fe
Fix Version formatting function
umar456 Jan 11, 2023
32862dd
Update convolve tests tolerances for floating point types
umar456 Jan 11, 2023
c8fa0e6
Add support for building older OpenCL versions.
umar456 Jan 11, 2023
db4ad6f
Fix meanvar tests to avoid segfaults for unsupported types
umar456 Jan 20, 2023
ec594ca
Fix warnings related to Version class
umar456 Jan 20, 2023
c57e2d8
Update vcpkg baseline to update OpenCL version
umar456 Jan 20, 2023
53cc5cc
Add group flags around LAPACKE libraries to avoid missing symbol errs
umar456 Jan 20, 2023
54b30e6
Fix extern half include directories command in cmake
umar456 Jan 20, 2023
1d36f1f
Fix error due to an extra brace during the namespace refactor
umar456 Jan 20, 2023
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
Prev Previous commit
Next Next commit
Fix some errors due to pow in CUDA code with fast math
  • Loading branch information
@umar456
umar456 committed Jan 20, 2023
commit b89ea4c82205e2361b157c952ed423dcfc2e7dd8
35 changes: 18 additions & 17 deletions src/backend/cuda/kernel/anisotropic_diffusion.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,8 @@ __forceinline__ __device__ int index(const int x, const int y, const int dim0,
return clamp(x, 0, dim0 - 1) * stride0 + clamp(y, 0, dim1 - 1) * stride1;
}

__device__ float quadratic(const float value) { return 1.0 / (1.0 + value); }
__device__
float quadratic(const float value) { return 1.0f / (1.0f + value); }

template<af_flux_function FluxEnum>
__device__ float gradientUpdate(const float mct, const float C, const float S,
Expand All @@ -39,13 +40,13 @@ __device__ float gradientUpdate(const float mct, const float C, const float S,
db = C - W;

if (FluxEnum == AF_FLUX_EXPONENTIAL) {
cx = expf((df * df + 0.25f * powf(dy + 0.5f * (SE - NE), 2)) * mct);
cxd = expf((db * db + 0.25f * powf(dy + 0.5f * (SW - NW), 2)) * mct);
cx = expf((df * df + 0.25f * afpowf(dy + 0.5f * (SE - NE), 2)) * mct);
cxd = expf((db * db + 0.25f * afpowf(dy + 0.5f * (SW - NW), 2)) * mct);
} else {
cx =
quadratic((df * df + 0.25f * powf(dy + 0.5f * (SE - NE), 2)) * mct);
quadratic((df * df + 0.25f * afpowf(dy + 0.5f * (SE - NE), 2)) * mct);
cxd =
quadratic((db * db + 0.25f * powf(dy + 0.5f * (SW - NW), 2)) * mct);
quadratic((db * db + 0.25f * afpowf(dy + 0.5f * (SW - NW), 2)) * mct);
}
delta += (cx * df - cxd * db);

Expand All @@ -54,13 +55,13 @@ __device__ float gradientUpdate(const float mct, const float C, const float S,
db = C - N;

if (FluxEnum == AF_FLUX_EXPONENTIAL) {
cx = expf((df * df + 0.25f * powf(dx + 0.5f * (SE - SW), 2)) * mct);
cxd = expf((db * db + 0.25f * powf(dx + 0.5f * (NE - NW), 2)) * mct);
cx = expf((df * df + 0.25f * afpowf(dx + 0.5f * (SE - SW), 2)) * mct);
cxd = expf((db * db + 0.25f * afpowf(dx + 0.5f * (NE - NW), 2)) * mct);
} else {
cx =
quadratic((df * df + 0.25f * powf(dx + 0.5f * (SE - SW), 2)) * mct);
quadratic((df * df + 0.25f * afpowf(dx + 0.5f * (SE - SW), 2)) * mct);
cxd =
quadratic((db * db + 0.25f * powf(dx + 0.5f * (NE - NW), 2)) * mct);
quadratic((db * db + 0.25f * afpowf(dx + 0.5f * (NE - NW), 2)) * mct);
}
delta += (cx * df - cxd * db);

Expand All @@ -87,8 +88,8 @@ __device__ float curvatureUpdate(const float mct, const float C, const float S,
df0 = df;
db0 = db;

gmsqf = (df * df + 0.25f * powf(dy + 0.5f * (SE - NE), 2));
gmsqb = (db * db + 0.25f * powf(dy + 0.5f * (SW - NW), 2));
gmsqf = (df * df + 0.25f * afpowf(dy + 0.5f * (SE - NE), 2));
gmsqb = (db * db + 0.25f * afpowf(dy + 0.5f * (SW - NW), 2));

gmf = sqrtf(1.0e-10 + gmsqf);
gmb = sqrtf(1.0e-10 + gmsqb);
Expand All @@ -102,8 +103,8 @@ __device__ float curvatureUpdate(const float mct, const float C, const float S,
df = S - C;
db = C - N;

gmsqf = (df * df + 0.25f * powf(dx + 0.5f * (SE - SW), 2));
gmsqb = (db * db + 0.25f * powf(dx + 0.5f * (NE - NW), 2));
gmsqf = (df * df + 0.25f * afpowf(dx + 0.5f * (SE - SW), 2));
gmsqb = (db * db + 0.25f * afpowf(dx + 0.5f * (NE - NW), 2));
gmf = sqrtf(1.0e-10 + gmsqf);
gmb = sqrtf(1.0e-10 + gmsqb);

Expand All @@ -114,14 +115,14 @@ __device__ float curvatureUpdate(const float mct, const float C, const float S,

if (delta > 0) {
prop_grad +=
(powf(fminf(db0, 0.0f), 2.0f) + powf(fmaxf(df0, 0.0f), 2.0f));
(afpowf(fminf(db0, 0.0f), 2.0f) + afpowf(fmaxf(df0, 0.0f), 2.0f));
prop_grad +=
(powf(fminf(db, 0.0f), 2.0f) + powf(fmaxf(df, 0.0f), 2.0f));
(afpowf(fminf(db, 0.0f), 2.0f) + afpowf(fmaxf(df, 0.0f), 2.0f));
} else {
prop_grad +=
(powf(fmaxf(db0, 0.0f), 2.0f) + powf(fminf(df0, 0.0f), 2.0f));
(afpowf(fmaxf(db0, 0.0f), 2.0f) + afpowf(fminf(df0, 0.0f), 2.0f));
prop_grad +=
(powf(fmaxf(db, 0.0f), 2.0f) + powf(fminf(df, 0.0f), 2.0f));
(afpowf(fmaxf(db, 0.0f), 2.0f) + afpowf(fminf(df, 0.0f), 2.0f));
}

return sqrtf(prop_grad) * delta;
Expand Down
2 changes: 1 addition & 1 deletion src/backend/cuda/kernel/jit.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,7 +65,7 @@ typedef cuDoubleComplex cdouble;
pow(static_cast<double>(lhs), static_cast<double>(rhs)));
#else
#define __pow(lhs, rhs) \
__float2int_rn(pow(__int2float_rn((int)lhs), __int2float_rn((int)rhs)))
__float2int_rn(powf(__int2float_rn((int)lhs), __int2float_rn((int)rhs)))
#endif
#define __powll(lhs, rhs) \
__double2ll_rn(pow(__ll2double_rn(lhs), __ll2double_rn(rhs)))
Expand Down
2 changes: 1 addition & 1 deletion src/backend/cuda/kernel/susan.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,7 +73,7 @@ __global__ void susan(T* out, const T* in, const unsigned idim0,
if (i * i + j * j < rSqrd) {
float c = m_0;
float m = shrdMem[b * shrdLen + a];
float exp_pow = powf((m - c) / t, 6.0f);
float exp_pow = afpowf((m - c) / t, 6.0f);
float cM = expf(-exp_pow);
nM += cM;
}
Expand Down
13 changes: 13 additions & 0 deletions src/backend/cuda/math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -392,6 +392,19 @@ template<typename T>
constexpr const __DH__ T clamp(const T value, const T lo, const T hi) {
return clamp(value, lo, hi, [](auto lhs, auto rhs) { return lhs < rhs; });
}

#ifdef AF_WITH_FAST_MATH
/// The pow function with fast math is constantly wrong with fast math
/// so this function converts the operation to double when fast-math
/// is used
__device__ inline double afpowf(double x, double y) { return pow(x, y); }
#else
/// The pow function with fast math is constantly wrong with fast math
/// so this function converts the operation to double when fast-math
/// is used
__device__ inline float afpowf(float x, float y) { return powf(x, y); }
#endif

} // namespace cuda
} // namespace arrayfire

Expand Down