• [x] Simple cmake support, maybe this can be done better or expanded on.
• [x] Tested with the tests folder, works really well.
• [x] Removed std::binary_function, apparently, you don't need it: https://stackoverflow.com/a/22387805/5729690 Definitely the cause of problem with C++17.
• [x] Enable all SMs.
• [x] We now add ModernGPU version to cmake as well, this increments the minor version to 2.13.
• [x] Addressing the use of __attribute__((alignment))
• [x] Added support for demos and tutorials, tested and compiles.
• [x] Added CI using github actions (workflows for both Windows and Ubuntu).
• [x] Not doing this; ~~Switch to using C++ alignas?~~

NOTE: This PR removes the old Makefile.

Hi,

I am having trouble getting reduce_by_key to compile. I think there are some compiler issues. My system is VS2015 with Cuda 8.

The following "attribute((aligned))" in tuple.hxx

makes the compiler choke.

The second problem is the restrict keyword usage in meta.hxx:

It seems, that the __restrict__keyword ist not allowed for type definitions. It throws this compiler error:

https://msdn.microsoft.com/en-us/library/097sy9kt.aspx

Removing problem 1 and all types that include restrict works. However, I am not sure how this influences the functionality of the library.

Thanks for the great work!

using cta_launch with NT not being a power of two results in wrong tid values passed to the lambda. simple test case:

static const int NT = 96
mgpu::cta_launch<NT>([=] MGPU_DEVICE (const int tid, const int block) {
      printf("thread %d %d %d\n", threadIdx.x, tid, threadIdx.x & (NT - 1));
}, 1, ctx);

For NT = 32, 64 and 128 this works fine. However setting NT to any multiple of 32 should be valid, right?

I am trying to run test on mergesort and segreduce, only to find the folling error

NVIDIA GeForce RTX 2080 : 1710.000 Mhz   (Ordinal 0)
46 SMs enabled. Compute Capability sm_75
FreeMem:   7862MB   TotalMem:   7981MB   64-bit pointers.
Mem Clock: 7000.000 Mhz x 256 bits   (448.0 GB/s)
ECC Disabled


Floating point exception (core dumped)

How could I fix the bug?

Hello, I am testing both moderngpu's mergesort and cub's radixsort for an implementation that I am currently working on.

I have benchmarked both and came to the conclusion that (at least for my data) the radix sort implementation from the CUB library is slightly faster due to the fact that it does not need several mallocs while sorting. On the other hand, moderngpu's mergesort appears to be performing mallocs every iteration, see below:

The box highlight in red is composed of several cudaMalloc and cudaFree. I was wondering if it is possible to get rid of these allocs, since this would make moderngpu faster than CUB in my application. For instance, in CUB I can simply pass the pointers with pre-allocated memory to the kernel launch.

Is there anyway to do this in moderngpu? Thank you for your time, Esteban

----------Edit: more information: I timed the time of one my apps run and found that the sys time is higher than that of CUB (I assume due to the cudaMallocs, because as it is, mgpu takes approx 10 to 30% less time.) Run with CUB: real 3m26.482s user 3m10.728s sys 0m15.584s

Run with Mgpu: real 3m34.427s user 3m11.548s sys 0m22.884s

Variable is just used to store and is not being used for anything else, this will cause warnings and errors while compiling with options. Just to be safer side, it would be better without the variable.

hi, thanks for the nice library. I'm using transform_lbs in a scenario where every work item needs to perform quite some computation and I'm unsure how to properly set the launch_params. Just reducing the values from the default launch_params_t<128, 11, 8> to something like launch_params_t<128, 3, 3> already increased the performance quite a bit. However, I'm unsure about how to set vt and vt0. should it be odd numbers? always vt > vt0? is the correctness affected if these numbers are not properly set? thanks for your help.

anyone having (recent) experience on converting *.sln files to *.cmake? successfully used one perl script in the past, now trying https://gist.github.com/sixman9/1781094

make: *** No rule to make target demo.o', needed bydemo'. Stop.

[email protected]:~/Development/quantcrunch/demo$ make --version GNU Make 3.81 Copyright (C) 2006 Free Software Foundation, Inc. This is free software; see the source for copying conditions. There is NO warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.

This program built for x86_64-pc-linux-gnu

Hi,

I have CUDA 5.5 installed and VS2012.

I tried to open the project but I am getting following error -

C:\Alenka\moderngpu\demo\demo.vcxproj(32,5): The imported project "C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\V110\BuildCustomizations\CUDA 5.0.props" was not found. Confirm that the path in the declaration is correct, and that the file exists on disk.

I am getting the same error for all the projects.

Thanks, Sawan

When I run make I get the following error:

demo/cities.cu(265): internal error: assertion failed: gen_paren_or_brace_dynamic_init: bad kind (cp_gen_be.c, line 21252 in gen_paren_or_brace_dynamic_init)

1 catastrophic error detected in the compilation of "demo/cities.cu". Compilation aborted. Aborted (core dumped)

I am not sure how to fix this. Any suggestions?

The "Reduce and Scan" chapter tries to link to "NVIDIA's Mark Harris on Optimizing Parallel Reduction in CUDA" but uses

`https://moderngpu.github.io/http|//developer.download.nvidia.com/assets/cuda/files/reduction.pdf`

instead of

`http://developer.download.nvidia.com/assets/cuda/files/reduction.pdf`

cd bf-knn nvcc -arch=sm_21 -I ../moderngpu/include/ bf_knn_device.cu bf_knn_host.cu demo.cc -o bf-knn-demo i am facing a problem when compiling this code in CMD. help me

ERROR: bf_knn_device.cu<44> : fatal error C1083: can not open include file 'moderngpu,cuh' No such a file or directory kindly help me how to include this file

I am trying to optimize the runtime of many calls to mgpu::lbs_segreduce with the same segment configuration. If I understand correctly, there is a way to use caching iterators. Could you maybe give an example how to use them? There is no documentation on that topic yet.

moderngpu git tag: v2.12_june_8_2016 CentOS 7, Linux version 3.10.0-327.22.2.el7.x86_64 gcc (GCC) 4.9.2 20150212 (Red Hat 4.9.2-6) Cuda compilation tools, release 7.5, V7.5.17

Calling segmented_sort() on a trivial data set (e.g., 1 segment of 24 values) results in cudaErrorInvalidConfiguration (error 9). Determined the cause to be located at kernel_segsort.hxx:69:

op_counters = fill<int2>(int2(), num_passes, context);

Per kernel_segsort.hxx:43-45, num_passes is set to 0 if num_ctas is equal to 1. Allocating an extra element in op_counters fixes the issue by allowing fill() to succeed:

op_counters = fill<int2>(int2(), num_passes + 1, context);

Minimum working example:

using uint = unsigned int;

#include "moderngpu/context.hxx"
#include "moderngpu/kernel_segsort.hxx"

#include <exception>
#include <sstream>

int main(int, char**) {
    using namespace mgpu;
    using namespace std;

    standard_context_t context;

    uint nValues = 24, nSegments = 1;
    uint *dKeys, *dVals, *dSegs;

    cudaMalloc(&dKeys, sizeof(*dKeys) * nValues);
    cudaMalloc(&dVals, sizeof(*dVals) * nValues);
    cudaMalloc(&dSegs, sizeof(*dSegs) * (nSegments + 1));

    uint hKeys[] = {
        98, 63, 82, 50, 40, 36, 44, 36, 49, 17, 21, 71,
        82, 67, 54, 74, 60, 15, 85, 58, 11, 42, 22, 97
    };

    uint hVals[] = {
        34, 11, 98, 41, 83, 34, 49, 40, 61, 14, 10, 82,
        19, 27, 80, 66, 54, 99, 79, 90, 70, 42, 16, 78
    };

    uint hSegs[] = {0, nValues};

    cudaMemcpy(dKeys, &hKeys[0], sizeof(*dKeys) * nValues, cudaMemcpyHostToDevice);
    cudaMemcpy(dVals, &hVals[0], sizeof(*dVals) * nValues, cudaMemcpyHostToDevice);
    cudaMemcpy(dSegs, &hSegs[0], sizeof(*dSegs) * (nSegments + 1), cudaMemcpyHostToDevice);

    segmented_sort(
        dKeys, dVals, (int)nValues,
        dSegs, (int)nSegments,
        less_t<uint>(), context);

    cudaDeviceSynchronize();
    cudaError_t error = cudaGetLastError();

    if (error != cudaSuccess) {
        ostringstream msg;
        msg << "cudaGetLastError(): " << cudaGetErrorString(error) << endl;
        throw runtime_error(msg.str());
    }

    cudaFree(dKeys);
    cudaFree(dVals);
    cudaFree(dSegs);
    return 0;
}

I'm currently in the process of replacing thrust with moderngpu because of https://github.com/Oblomov/titanxstall/pull/2

My problem now is to choose an adequate launch_params_t depending on the number of elements to be sorted, as that number can be changed by the user.

Thrust chooses this automatically in a way, I did not find out how, but it works quite good.

Any recommendations?

Hi! Would it be possible to complete support for mgpu::tie? As is, it's missing support for the assignment of mgpu::tuple<args_t...> to
mgpu::tuple<args_t&...>.

Example

#include <moderngpu/tuple.hxx>

MGPU_DEVICE
mgpu::tuple<int, float>
example_tuple() {
  return mgpu::make_tuple(1, 2.f);
}


MGPU_DEVICE
mgpu::tuple<int, float>
example_using_tie() {
  int a;
  float b;
  mgpu::tie(a,b) = example_tuple();
  return mgpu::make_tuple(a,b);
}

yields

main.cu(15): error: no operator "=" matches these operands
            operand types are: mgpu::tuple<int &, float &> = mgpu::tuple<int, float>

Thanks!