search

NVIDIA / cutlass

CUDA Templates for Linear Algebra Subroutines
Other
5k stars 853 forks source link

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

Open cloudhan opened 7 months ago

cloudhan commented 7 months ago

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

a75b4ac483166189a45290783cb0a18af5ff0ea5

hwu36 commented 6 months ago

@thakkarV

thakkarV commented 6 months ago

@ccecka @mhoemmen

ccecka commented 6 months ago

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

github-actions[bot] commented 5 months ago

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[bot] commented 2 months ago

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 commented 2 months ago

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[bot] commented 1 month ago

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 commented 1 month ago

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

github-actions[bot] commented 2 weeks ago

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.