[BUG] right_inverse with 1d layout results "#2700-D: attempt to access expired storage" on constexpr return value.

[cloudhan]

Describe the bug

Steps/Code to reproduce bug

#include "cute/tensor.hpp"
using namespace cute;
__global__ void kernel() {
  constexpr auto weird = right_inverse(make_layout(_2{}, _1{}));
  print(weird);
}

int main() {
  kernel<<<1, 1>>>();
  cudaDeviceSynchronize();
}

compile the code cause

error: expression must have a constant value
include/cute/layout.hpp(1143): note #2700-D: attempt to access expired storage
include/cute/algorithm/tuple_algorithms.hpp(114): note #2693-D: called from:
include/cute/algorithm/tuple_algorithms.hpp(248): note #2693-D: called from:
include/cute/layout.hpp(1143): note #2693-D: called from:

remove the constexpr however, the program compiles smoothly and prints _2:_1

Expected behavior
It compiles.

Environment details (please complete the following information):
reproduced with cuda 11.8 and 12.1

Additional context

a75b4ac

[hwu36]

@thakkarV

[thakkarV]

@ccecka @mhoemmen

[ccecka]

Weird one, but I think I know where this is coming from. I'll include a fix in the next CuTe review

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[github-actions]

This issue has been labeled inactive-90d due to no recent activity in the past 90 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed.

[mhoemmen]

I do get a lot of warnings with MSVC about CuTe functions returning temporary storage. I suspect MSVC doesn't follow the reference forwarding rules of other compilers

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[mhoemmen]

Commenting so the friendly bot doesn't auto-close this issue.

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[cloudhan]

comment to avoid auto-close

[mhoemmen]

@ccecka has this been fixed?

[ccecka]

Doesn't appear to be, no.

[github-actions]

This issue has been labeled inactive-90d due to no recent activity in the past 90 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed.

[leimao]

@ccecka Still no fix? I got a problem probably relevant to this when I tried to make make_tiled_copy as a constexpr function.

root@5bfb1ac7e7b2:/mnt# gcc --version
gcc (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0
Copyright (C) 2021 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.

root@5bfb1ac7e7b2:/mnt# nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Thu_Mar_28_02:18:24_PDT_2024
Cuda compilation tools, release 12.4, V12.4.131
Build cuda_12.4.r12.4/compiler.34097967_0

[thakkarV]

question, are you using the build flags generated by the CUTLASS CMake?

[leimao]

question, are you using the build flags generated by the CUTLASS CMake?

Thanks @thakkarV. I am new to this. Could you please point me to those build flags?

I was simply adding constexprs to the functions in this file and I was using my own CMake script for building. I might have missed something critical.

[leimao]

More specifically, was there any problem to make this make_tiled_copy function constexpr?

[thakkarV]

we dont support custom cmake configs. please use the one that is in our repo.

[leimao]

we dont support custom cmake configs. please use the one that is in our repo.

But... Eventually the user will write their own scripts for building the program, right? Could you please point me to the key CMake files where those build flags were set?

[thakkarV]

No, that's not the case. if you are importing CUTLASS through your own project, then you can https://stackoverflow.com/questions/56561591/link-one-project-to-another-cmake