This is an archive of the discontinued LLVM Phabricator instance.

[OPENMP][NVPTX]General formatting/code improvement, NFC.
ClosedPublic

Authored by ABataev on Jan 3 2019, 1:18 PM.

Diff Detail

Repository
rL LLVM

Event Timeline

ABataev created this revision.Jan 3 2019, 1:18 PM
grokos added inline comments.Jan 4 2019, 8:24 AM
libomptarget/deviceRTLs/nvptx/src/loop.cu
420 ↗(On Diff #180129)

What about the Pascal issue when inlining this function?

libomptarget/deviceRTLs/nvptx/src/reduction.cu
197 ↗(On Diff #180129)

Can the rest of __BALLOT_SYNC instances in this file be replaced with __ACTIVEMASK()? If so, then we can remove the definition of __BALLOT_SYNC altogether from omptarget-nvptx.h.

ABataev marked 2 inline comments as done.Jan 4 2019, 9:38 AM
ABataev added inline comments.
libomptarget/deviceRTLs/nvptx/src/loop.cu
420 ↗(On Diff #180129)

Everything works with the inlined functions. Seems to me, the problem was caused again with the __syncthreads() problem. Just need to remove the comment.

libomptarget/deviceRTLs/nvptx/src/reduction.cu
197 ↗(On Diff #180129)

Will try

ABataev updated this revision to Diff 180268.Jan 4 2019, 10:13 AM

Removed __BALLOT_SYNC completely.

grokos accepted this revision.Jan 4 2019, 12:01 PM

Looks good. Thanks!

This revision is now accepted and ready to land.Jan 4 2019, 12:01 PM
This revision was automatically updated to reflect the committed changes.