Skip to content
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

Jump to bottom

Patch for DoublePrec_ACC build using SYCL #1049

Merged
merged 1 commit into from
Dec 15, 2023
Merged

Patch for DoublePrec_ACC build using SYCL #1049

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
18 changes: 18 additions & 0 deletions src/acc/sycl/sycl_kernels/diff2_gpu.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -146,17 +146,27 @@ void sycl_kernel_diff2_coarse(
{
// translatePixel(x, y, z, tx, ty, tz, s_real[pix+init_pixel%block_sz], s_imag[pix+init_pixel%block_sz], real, imag);
XFLOAT val {x*tx + y*ty + z*tz};
#ifdef ACC_DOUBLE_PRECISION
XFLOAT s {sycl::sin(val)};
XFLOAT c {sycl::cos(val)};
#else
XFLOAT s {sycl::native::sin(val)};
XFLOAT c {sycl::native::cos(val)};
#endif
real = c * s_real[pix + init_pixel%block_sz] - s * s_imag[pix + init_pixel%block_sz];
imag = c * s_imag[pix + init_pixel%block_sz] + s * s_real[pix + init_pixel%block_sz];
}
else
{
// translatePixel(x, y, tx, ty, s_real[pix+init_pixel%block_sz], s_imag[pix+init_pixel%block_sz], real, imag);
XFLOAT val {x*tx + y*ty};
#ifdef ACC_DOUBLE_PRECISION
XFLOAT s {sycl::sin(val)};
XFLOAT c {sycl::cos(val)};
#else
XFLOAT s {sycl::native::sin(val)};
XFLOAT c {sycl::native::cos(val)};
#endif
real = c*s_real[pix + init_pixel%block_sz] - s*s_imag[pix + init_pixel%block_sz];
imag = c*s_imag[pix + init_pixel%block_sz] + s*s_real[pix + init_pixel%block_sz];
}
Expand Down Expand Up @@ -396,7 +406,11 @@ void sycl_kernel_diff2_CC_coarse(
}

if (tid == 0)
#ifdef ACC_DOUBLE_PRECISION
g_diff2[iorient*trans_num + itrans] += -s_weight[0] / sycl::sqrt(s_norm[0]);
#else
g_diff2[iorient*trans_num + itrans] += -s_weight[0] / sycl::native::sqrt(s_norm[0]);
#endif
}

template<bool REF3D, bool DATA3D, int block_sz>
Expand Down Expand Up @@ -517,7 +531,11 @@ void sycl_kernel_diff2_CC_fine(
}

if (tid < trans_num)
#ifdef ACC_DOUBLE_PRECISION
g_diff2s[d_job_idx[bid] + tid] += -s[tid * block_sz] / sycl::sqrt(s_cc[tid * block_sz]);
#else
g_diff2s[d_job_idx[bid] + tid] += -s[tid * block_sz] / sycl::native::sqrt(s_cc[tid * block_sz]);
#endif
}
}

Expand Down
47 changes: 47 additions & 0 deletions src/acc/sycl/sycl_kernels/helper_gpu.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -345,7 +345,11 @@ static void sycl_kernel_exponentiate(
assert(count <= std::numeric_limits<int>::max());
sycl::queue *Q = dynamic_cast<devSYCL*>(devAcc)->getQueue();
#ifdef USE_MORE_ONEDPL
#ifdef ACC_DOUBLE_PRECISION
dpl::transform(dynamic_cast<devSYCL*>(devAcc)->getDevicePolicy(), g_array, g_array+count, g_array, [=](T v) { v+=add; v = (v<KERNEL_EXP_VALUE) ? static_cast<T>(0) : sycl::exp(v); return v;});
#else
dpl::transform(dynamic_cast<devSYCL*>(devAcc)->getDevicePolicy(), g_array, g_array+count, g_array, [=](T v) { v+=add; v = (v<KERNEL_EXP_VALUE) ? static_cast<T>(0) : sycl::native::exp(v); return v;});
#endif
#else
Q->submit([&](sycl::handler &cgh)
{
Expand All @@ -359,7 +363,11 @@ static void sycl_kernel_exponentiate(
if (a < KERNEL_EXP_VALUE)
g_array[idx] = static_cast<T>(0);
else
#ifdef ACC_DOUBLE_PRECISION
g_array[idx] = sycl::exp(a);
#else
g_array[idx] = sycl::native::exp(a);
#endif
}
});
}).wait_and_throw();
Expand Down Expand Up @@ -852,7 +860,11 @@ inline void powerClass(int gridSize,

inline std::pair<XFLOAT, XFLOAT> sycl_sincos(XFLOAT val)
{
#ifdef ACC_DOUBLE_PRECISION
return std::make_pair(sycl::sin(val), sycl::cos(val));
#else
return std::make_pair(sycl::native::sin(val), sycl::native::cos(val));
#endif
}

inline void translatePixel(
Expand All @@ -866,8 +878,13 @@ inline void translatePixel(
XFLOAT &tImag)
{
XFLOAT v = x * tx + y * ty;
#ifdef ACC_DOUBLE_PRECISION
XFLOAT s = sycl::sin(v);
XFLOAT c = sycl::cos(v);
#else
XFLOAT s = sycl::native::sin(v);
XFLOAT c = sycl::native::cos(v);
#endif

tReal = c * real - s * imag;
tImag = c * imag + s * real;
Expand All @@ -886,8 +903,13 @@ inline void translatePixel(
XFLOAT &tImag)
{
XFLOAT v = x * tx + y * ty + z * tz;
#ifdef ACC_DOUBLE_PRECISION
XFLOAT s = sycl::sin(v);
XFLOAT c = sycl::cos(v);
#else
XFLOAT s = sycl::native::sin(v);
XFLOAT c = sycl::native::cos(v);
#endif

tReal = c * real - s * imag;
tImag = c * imag + s * real;
Expand Down Expand Up @@ -916,15 +938,25 @@ inline void computeSincosLookupTable2D(unsigned long trans_num,
for(int x=0; x<xSize; x++) {
unsigned long index = i * xSize + x;
XFLOAT v = x * tx;
#ifdef ACC_DOUBLE_PRECISION
sin_x[index] = sycl::sin(v);
cos_x[index] = sycl::cos(v);
#else
sin_x[index] = sycl::native::sin(v);
cos_x[index] = sycl::native::cos(v);
#endif
}

for(int y=0; y<ySize; y++) {
unsigned long index = i * ySize + y;
XFLOAT v = y * ty;
#ifdef ACC_DOUBLE_PRECISION
sin_y[index] = sycl::sin(v);
cos_y[index] = sycl::cos(v);
#else
sin_y[index] = sycl::native::sin(v);
cos_y[index] = sycl::native::cos(v);
#endif
}
}
}
Expand All @@ -951,22 +983,37 @@ inline void computeSincosLookupTable3D(unsigned long trans_num,
for(int x=0; x<xSize; x++) {
unsigned long index = i * xSize + x;
XFLOAT v = x * tx;
#ifdef ACC_DOUBLE_PRECISION
sin_x[index] = sycl::sin(v);
cos_x[index] = sycl::cos(v);
#else
sin_x[index] = sycl::native::sin(v);
cos_x[index] = sycl::native::cos(v);
#endif
}

for(int y=0; y<ySize; y++) {
unsigned long index = i * ySize + y;
XFLOAT v = y * ty;
#ifdef ACC_DOUBLE_PRECISION
sin_y[index] = sycl::sin(v);
cos_y[index] = sycl::cos(v);
#else
sin_y[index] = sycl::native::sin(v);
cos_y[index] = sycl::native::cos(v);
#endif
}

for(int z=0; z<zSize; z++) {
unsigned long index = i * zSize + z;
XFLOAT v = z * tz;
#ifdef ACC_DOUBLE_PRECISION
sin_z[index] = sycl::sin(v);
cos_z[index] = sycl::cos(v);
#else
sin_z[index] = sycl::native::sin(v);
cos_z[index] = sycl::native::cos(v);
#endif
}
}
}
Expand Down
Loading