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.

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

Fix jitter operator illegal memory access #1914

Merged
merged 1 commit into from
Apr 28, 2020
Merged

Conversation

JanuszL
Copy link
Contributor

@JanuszL JanuszL commented Apr 28, 2020

  • JitterAugment accesses random generator on the GPU beyond allowed range
  • added a fix to make sure that access to random generator is inside allowed range

Signed-off-by: Janusz Lisiecki jlisiecki@nvidia.com

Why we need this PR?

Pick one, remove the rest

  • It fixes a bug of jitter operator illegal memory access

What happened in this PR?

Fill relevant points, put NA otherwise. Replace anything inside []

  • What solution was applied:
    added a fix to make sure that access to random generator is inside allowed range
  • Affected modules and functionalities:
    jitter
  • Key points relevant for the review:
    NA
  • Validation and testing:
    CI
  • Documentation (including examples):
    NA

In response to #966

JIRA TASK: [NA]

@dali-automaton
Copy link
Collaborator

CI MESSAGE: [1285453]: BUILD STARTED

@@ -38,7 +39,8 @@ Randomizer<GPUBackend>::Randomizer(int seed, size_t len) {
len_ = len;
cudaGetDevice(&device_);
states_ = GPUBackend::New(sizeof(curandState) * len, true);
initializeStates<<<128, 256>>>(len_, seed, reinterpret_cast<curandState*>(states_));
initializeStates<<<std::ceil(len_/static_cast<float>(block_size_)), block_size_>>>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
initializeStates<<<std::ceil(len_/static_cast<float>(block_size_)), block_size_>>>
initializeStates<<<div_ceil(len_, block_size_), block_size_>>>

using floating point arithmetic may lead to incorrect results when len_ > 256*2^23

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

done

@dali-automaton
Copy link
Collaborator

CI MESSAGE: [1285453]: BUILD FAILED

- JitterAugment accesses random generator on the GPU beyond allowed range
- added a fix to make sure that access to random generator is inside allowed
  range

Signed-off-by: Janusz Lisiecki <jlisiecki@nvidia.com>
@dali-automaton
Copy link
Collaborator

CI MESSAGE: [1285483]: BUILD STARTED

@dali-automaton
Copy link
Collaborator

CI MESSAGE: [1285483]: BUILD PASSED

@JanuszL JanuszL merged commit 0975a42 into NVIDIA:master Apr 28, 2020
@JanuszL JanuszL deleted the fix_jitter branch May 19, 2020 23:03
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

4 participants