experiments/STT-tensorflow
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[tfdbg] Fix DebugNumericSummaryV2OP kernel mode=CURT_HEALTH breakage on Windows GPU builds

Browse Source 
PiperOrigin-RevId: 286585970
Change-Id: I2db43f37def35fb0f58a15f09f1d9c026b27a536
This commit is contained in:
Shanqing Cai 2019-12-20 08:50:57 -08:00 committed by TensorFlower Gardener
parent 09a10c139a
commit f079d31f77
1 changed files with 8 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
tensorflow/core/kernels/debug_ops_gpu.cu.cc
View File

@ -42,9 +42,9 @@ __global__ void CurtHealthKernel(const Tin* __restrict__ data, int size,
const int32 total_thread_count = gridDim.x * blockDim.x; const int32 total_thread_count = gridDim.x * blockDim.x;
int32 offset = thread_id; int32 offset = thread_id;
while (offset < size) { while (offset < size) {
if (isinf(data[offset]) || isnan(data[offset])) { if (Eigen::numext::isinf(data[offset]) ||
Eigen::numext::isnan(data[offset])) {
output[0] = 1.0; output[0] = 1.0;
} }
offset += total_thread_count; offset += total_thread_count;
@ -63,14 +63,14 @@ __global__ void ConciseHealthKernel(const Tin* __restrict__ data, int size,
Tout accum[3] = {0.0, 0.0, 0.0}; Tout accum[3] = {0.0, 0.0, 0.0};
while (offset < size) { while (offset < size) {
if (isinf(data[offset])) { if (Eigen::numext::isinf(data[offset])) {
if (data[offset] < static_cast<Tin>(0.f)) { if (data[offset] < static_cast<Tin>(0.f)) {
++accum[0]; ++accum[0];
} else { } else {
++accum[1]; ++accum[1];
} }
} }
if (isnan(data[offset])) { if (Eigen::numext::isnan(data[offset])) {
++accum[2]; ++accum[2];
} }
offset += total_thread_count; offset += total_thread_count;
@ -94,13 +94,13 @@ __global__ void FullHealthKernel(const Tin* __restrict__ data, int size,
Tout accum[6] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0}; Tout accum[6] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0};
while (offset < size) { while (offset < size) {
if (isinf(data[offset])) { if (Eigen::numext::isinf(data[offset])) {
if (data[offset] < static_cast<Tin>(0.f)) { if (data[offset] < static_cast<Tin>(0.f)) {
++accum[0]; ++accum[0];
} else { } else {
++accum[1]; ++accum[1];
} }
} else if (isnan(data[offset])) { } else if (Eigen::numext::isnan(data[offset])) {
++accum[2]; ++accum[2];
} else { } else {
if (data[offset] < static_cast<Tin>(0.f)) { if (data[offset] < static_cast<Tin>(0.f)) {
@ -136,14 +136,14 @@ __global__ void ReduceInfNanThreeSlotsKernel(const Tin* __restrict__ data,
int32 offset = thread_id; int32 offset = thread_id;
while (offset < size) { while (offset < size) {
if (isinf(data[offset])) { if (Eigen::numext::isinf(data[offset])) {
if (data[offset] < static_cast<Tin>(0.f)) { if (data[offset] < static_cast<Tin>(0.f)) {
output[0] = -std::numeric_limits<Tout>::infinity(); output[0] = -std::numeric_limits<Tout>::infinity();
} else { } else {
output[1] = std::numeric_limits<Tout>::infinity(); output[1] = std::numeric_limits<Tout>::infinity();
} }
} }
if (isnan(data[offset])) { if (Eigen::numext::isnan(data[offset])) {
output[2] = std::numeric_limits<Tout>::quiet_NaN(); output[2] = std::numeric_limits<Tout>::quiet_NaN();
} }
offset += total_thread_count; offset += total_thread_count;