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

Implementing cunumeric.random.BitGenerator #254

Merged
merged 96 commits into from
Jul 26, 2022

Conversation

fduguet-nv
Copy link
Contributor

No description provided.

@fduguet-nv fduguet-nv marked this pull request as ready for review April 1, 2022 19:13
.pre-commit-config.yaml Outdated Show resolved Hide resolved
cunumeric/deferred.py Outdated Show resolved Hide resolved
cunumeric/deferred.py Outdated Show resolved Hide resolved
cunumeric/deferred.py Outdated Show resolved Hide resolved
cunumeric/deferred.py Outdated Show resolved Hide resolved
src/cunumeric/random/bitgenerator_curand.inl Outdated Show resolved Hide resolved
src/cunumeric/random/bitgenerator_curand.inl Outdated Show resolved Hide resolved
src/cunumeric/random/bitgenerator_curand.inl Outdated Show resolved Hide resolved
src/cunumeric/random/bitgenerator_curand.inl Outdated Show resolved Hide resolved
src/cunumeric/random/bitgenerator_curand.inl Outdated Show resolved Hide resolved

// Match these to BitGeneratorDistribution in config.py
enum CuNumericBitGeneratorDistribution {
CUNUMERIC_BITGENDIST_INTEGERS_32 = 1,
Copy link
Contributor

Choose a reason for hiding this comment

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

are we expecting to add more distributions? if so, let's leave out enum values except the first one, otherwise you would need to renumber them every time you add a new distribution.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Yes, same point as above. Can we push that to next PR ?

)
return res

def integers(
Copy link
Contributor

Choose a reason for hiding this comment

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

any reason you added distribution methods to BitGenerator? it looks like the BitGenerator in both NumPy and CuPy is a simple opaque handle with only a couple of methods. I guess we could keep these distribution methods only in Generator.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Yes, we could. I found it easier that way: the two options I saw were:

  1. this way
  2. having Generator use internal fields from BitGenerator (self.seed, self.flags, self.generatorType), and make Generator aware of underlying implementation of BitGenerator.
    I favored the first one. Should you prefer otherwise (option 2 or another option), please advise.

def __init__(self, bit_generator):
self.bit_generator = bit_generator

def integers(
Copy link
Contributor

Choose a reason for hiding this comment

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

(I moved my comment to a more appropriate line) can this replace the legacy random.randint implementation?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Yes. That it the target. It seems that randint is used by other parts. I wanted to split the sources of potential issues: cunumeric.random errors vs other errors from tests using cunumeric.random. Though, since this is the "default generator" we could use any such as the one you proposed. Both are fine.

from cunumeric.random.bitgenerator import XORWOW, BitGenerator


class Generator:
Copy link
Contributor

Choose a reason for hiding this comment

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

any reason you left out uniform? I feel it's straightforward to add it given that you already have the implementation.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

missing test and last API link - fixed in commit 7364d28

std::vector<legate::Store>& args)
{
generator_map_t& genmap = get_generator_map();
// printtid((int)op);
Copy link
Contributor

Choose a reason for hiding this comment

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

please remove this.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

self, generatorType, seed, flags, forceCreate=False
):
self.current_random_bitgenid = self.current_random_bitgenid + 1
if forceCreate:
Copy link
Contributor

Choose a reason for hiding this comment

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

do you have any test that exercises this path? otherwise this code path can be bitrotted

Copy link
Contributor Author

Choose a reason for hiding this comment

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

test_force_build

src/cunumeric.mk Outdated
cunumeric/mapper.cc

GEN_CPU_SRC += cunumeric/cephes/chbevl.cc \
GEN_CPU_SRC += cunumeric/cephes/chbevl.cc \
Copy link
Contributor

Choose a reason for hiding this comment

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

please revert this bogus change

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Copy link
Contributor Author

@fduguet-nv fduguet-nv Jul 15, 2022

Choose a reason for hiding this comment

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

moved to random/random.mk to avoid such issues (and merge conflicts)

src/cunumeric.mk Outdated
cunumeric/cephes/i0.cc

ifeq ($(strip $(USE_OPENMP)),1)
GEN_CPU_SRC += cunumeric/ternary/where_omp.cc \
GEN_CPU_SRC += cunumeric/ternary/where_omp.cc \
Copy link
Contributor

Choose a reason for hiding this comment

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

same here. (I bet this is from the IDE you're using)

Copy link
Contributor Author

Choose a reason for hiding this comment

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

In the code, we don't use tabs; here we do. What is the tab <> #spaces equivalent I should use ?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

moved to random/random.mk to avoid such issues (and merge conflicts)

#include "generator.h"
#include "generator_create.inl"

extern "C" curandStatus_t CURANDAPI randutilCreateGeneratorHost(randutilGenerator_t* generator,
Copy link
Contributor

Choose a reason for hiding this comment

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

can we remove CURANDAPI? I don't think we're defining a cuRAND API here (same applies to all the other functions)

Copy link
Contributor Author

Choose a reason for hiding this comment

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

@magnatelee
Copy link
Contributor

left several questions and comments. otherwise looks good. (there are some refactoring changes I'd do but I'll do them in a follow-on PR, as this was sitting too long.) once you address the remaining comments and @marcinz confirms the CPU only build, we can merge this PR.

#include <curand_mrg32k3a.h>
#undef __device__

#define __CUDACC_RTC__
Copy link
Collaborator

Choose a reason for hiding this comment

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

@fduguet-nv Did you mean to #undef this after line 43? I've spent quite a lot of time trying to compile CPU package, and I finally figured out that this is the line that causes the following error in my compilation:

In file included from /home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657756345722/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/cuda/std/cstddef:32,
                 from /home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657756345722/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/cuda/std/type_traits:17,
                 from /home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657756345722/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/curand_kernel.h:72,
                 from cunumeric/random/randutil/randutil_curand.h:95,
                 from cunumeric/random/randutil/generator.h:22,
                 from cunumeric/random/randutil/generator_host_straightforward.cc:17:
/home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657756345722/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/cuda/std/detail/libcxx/include/cstddef:58:9: error: 'ptrdiff_t' has not been declared in '::'
   58 | using ::ptrdiff_t;
      |         ^~~~~~~~~
compilation terminated due to -Wfatal-errors.

Copy link
Collaborator

Choose a reason for hiding this comment

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

If I "fix" this, then I get a different error:

In file included from /home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657762144072/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/driver_types.h:61,
                 from /home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657762144072/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/builtin_types.h:59,
                 from /home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657762144072/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/cuda_runtime.h:91,
                 from /home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657762144072/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/curand.h:59,
                 from /home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657762144072/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/curand_kernel.h:99,
                 from cunumeric/random/randutil/randutil_curand.h:96,
                 from cunumeric/random/randutil/generator.h:22,
                 from cunumeric/random/randutil/generator_host_straightforward.cc:17:
/home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657762144072/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/vector_types.h:185:30: error: redefinition of 'struct uint2'
  185 | __cuda_builtin_vector_align8(uint2, unsigned int x; unsigned int y;);
      |                              ^~~~~
/home/mzalewski/.pyenv/versions/miniforge3/conda-bld/cunumeric_1657762144072/_h_env_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_placehold_p/include/vector_types.h:93:40: note: in definition of macro '__cuda_builtin_vector_align8'
   93 | struct __device_builtin__ __align__(8) tag         \
      |                                        ^~~
compilation terminated due to -Wfatal-errors.

@fduguet-nv fduguet-nv merged commit bf81d95 into nv-legate:branch-22.07 Jul 26, 2022
@bryevdv bryevdv mentioned this pull request Jul 26, 2022
jjwilke pushed a commit to jjwilke/cunumeric that referenced this pull request Jul 29, 2022
* Implementing BitGenerator - GPU part - creation

* Plugging destroy

* skip ahead for GPU BitGenerator random_raw

* Refectoring test

* Piping done for BitGenerator.random_raw

* LGT-260 -- implemented random_raw for CUDA

* Implementing BitGenerator - GPU part - creation

* Plugging destroy

* skip ahead for GPU BitGenerator random_raw

* Refectoring test

* Piping done for BitGenerator.random_raw

* LGT-260 -- implemented random_raw for CUDA

* Updating pre-commit

* Adding Multi-GPU support

* Refactoring

* refactoring

* CPU support - issues remain on multi-cpu mapper

* Attempts to support multi-dimension

* Finalizing multi-cpu implementation of BitGenerator

* Repro behavior of BitGenerator

* Adding random in tests

* Fixes from code review - part 1

* use of logger

* Enums improvement

* More on code review

* code review

* Removing task in python destructor

* Using Shape

* Lazy destroy also at destroy

* Using create_buffer for temporary buffer creation

* Fixing bugs following merge

* Adding alternate lazy init implementation

* Adding force create and force destroy

* Removing constraint in mapper

* Updating test

* Fixing Eager code branch

* Updating license information and adding integer generator

* PR review

* Removing reference to removed omp file

* Changes from PR comments

* Adding generator.random implementation

* Adding generator.lognormal implementation

* Adding generator.normal implementation

* Adding generator.poisson implementation

* Refectoring randutil to allow host-only compilation without CUDA enabled

* Moving distributions to HOST only capable

* Minor refactoring to split compilation of further distributions

* Fixing openmp test failure

* Adding random.generator.exponential

* Adding random.generator.gumbel

* Splitting distributions in several files

* Actually applying the WAR

* Adding random.generator.laplace

* Adding random.generator.logistic

* Adding random.generator.pareto

* Adding random.generator.power

* Adding random.generator.rayleigh

* Adding random.generator.standard_cauchy

* Adding random.generatortriangular

* Adding random.generator.weibul

* Adding random.generator.bytes

* PR review

* PR review

* PR review

* PR review

* fixing glitch

* More work on PR

* Fixing CURANDAPI

* Moving logger symbol from .cu to .cc

* removing destroy
jjwilke pushed a commit to jjwilke/cunumeric that referenced this pull request Jul 29, 2022
* Implementing BitGenerator - GPU part - creation

* Plugging destroy

* skip ahead for GPU BitGenerator random_raw

* Refectoring test

* Piping done for BitGenerator.random_raw

* LGT-260 -- implemented random_raw for CUDA

* Implementing BitGenerator - GPU part - creation

* Plugging destroy

* skip ahead for GPU BitGenerator random_raw

* Refectoring test

* Piping done for BitGenerator.random_raw

* LGT-260 -- implemented random_raw for CUDA

* Updating pre-commit

* Adding Multi-GPU support

* Refactoring

* refactoring

* CPU support - issues remain on multi-cpu mapper

* Attempts to support multi-dimension

* Finalizing multi-cpu implementation of BitGenerator

* Repro behavior of BitGenerator

* Adding random in tests

* Fixes from code review - part 1

* use of logger

* Enums improvement

* More on code review

* code review

* Removing task in python destructor

* Using Shape

* Lazy destroy also at destroy

* Using create_buffer for temporary buffer creation

* Fixing bugs following merge

* Adding alternate lazy init implementation

* Adding force create and force destroy

* Removing constraint in mapper

* Updating test

* Fixing Eager code branch

* Updating license information and adding integer generator

* PR review

* Removing reference to removed omp file

* Changes from PR comments

* Adding generator.random implementation

* Adding generator.lognormal implementation

* Adding generator.normal implementation

* Adding generator.poisson implementation

* Refectoring randutil to allow host-only compilation without CUDA enabled

* Moving distributions to HOST only capable

* Minor refactoring to split compilation of further distributions

* Fixing openmp test failure

* Adding random.generator.exponential

* Adding random.generator.gumbel

* Splitting distributions in several files

* Actually applying the WAR

* Adding random.generator.laplace

* Adding random.generator.logistic

* Adding random.generator.pareto

* Adding random.generator.power

* Adding random.generator.rayleigh

* Adding random.generator.standard_cauchy

* Adding random.generatortriangular

* Adding random.generator.weibul

* Adding random.generator.bytes

* PR review

* PR review

* PR review

* PR review

* fixing glitch

* More work on PR

* Fixing CURANDAPI

* Moving logger symbol from .cu to .cc

* removing destroy
ipdemes pushed a commit to ipdemes/cunumeric that referenced this pull request Aug 4, 2022
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.

3 participants