eyalroz / cuda-kat

CUDA kernel author's tools
BSD 3-Clause "New" or "Revised" License
104 stars 8 forks source link

Array.fill with constexpr variable #54

Open codecircuit opened 4 years ago

codecircuit commented 4 years ago

This is a general problem of nvcc I would say:

#include <array>
#include <kat/containers/array.hpp>

constexpr int duzzle = -7;

__global__
void kernel() {
    kat::array<int, 7> arr;
    arr.fill(duzzle); // fails to compile
}

int main() {

    std::array<int, 7> arr;
    arr.fill(duzzle);
}

The problem here is the signature of fill(const value_type&). E.g. if we omit the reference it works fine or if we define the constexpr variable in the device function itself. A real funny workaround:

__global__
void kernel() {
    kat::array<int, 7> arr;
        constexpr int duzzle_ = duzzle;
    arr.fill(duzzle_);
}
eyalroz commented 4 years ago

This reminds me I don't have any unit tests for the array class. Unlike cuda-api-wrappers (where I should have done otherwise), here' I'm trying for full test coverage of most facilities in the library.

eyalroz commented 4 years ago

But to the point: I'm not quite sure what to do here. Just opened this on StackOverflow.

eyalroz commented 4 years ago

A less-verbose workaround is arr.fill(int{duzzle}). Not what I'd like to have of course.

eyalroz commented 4 years ago

@codecircuit : The array class now has unit tests, and consequently, I've fixed some bugs with it, which have now landed on the master branch. FYI.