diff --git a/src/pass/LambdaAnalysis.cpp b/src/pass/LambdaAnalysis.cpp index 3fb601f0..265ee193 100644 --- a/src/pass/LambdaAnalysis.cpp +++ b/src/pass/LambdaAnalysis.cpp @@ -552,8 +552,11 @@ class LambdaPassImpl { for (const auto &JitVarInfo : It->second) { auto *SlotC = dyn_cast(JitVarInfo.Slot); auto *OffsetC = dyn_cast(JitVarInfo.Offset); - auto RCType = getRCTypeForLLVMType(JitVarInfo.Type); - if (!SlotC || !OffsetC || !RCType) + if (!SlotC || !OffsetC) + reportFatalError("Failed to emit lambda schema metadata"); + auto RCType = getRCTypeForLLVMType( + LambdaStorageType->getElementType(SlotC->getZExtValue())); + if (!RCType) reportFatalError("Failed to emit lambda schema metadata"); LLVMContext &Ctx = M.getContext(); diff --git a/tests/cpu/CMakeLists.txt b/tests/cpu/CMakeLists.txt index fd8a3009..e33f1c29 100644 --- a/tests/cpu/CMakeLists.txt +++ b/tests/cpu/CMakeLists.txt @@ -66,6 +66,7 @@ CREATE_CPU_TEST(types_api types_api.cpp) CREATE_CPU_TEST(lambda lambda.cpp) CREATE_CPU_TEST(lambda_ptrtoint_capture lambda_ptrtoint_capture.cpp) CREATE_CPU_TEST(lambda_def lambda_def.cpp) +CREATE_CPU_TEST(lambda_cast_capture lambda_cast_capture.cpp) CREATE_CPU_TEST(lambda_nested lambda_nested.cpp) CREATE_CPU_TEST(annot_nested annot_nested.cpp) CREATE_CPU_TEST(lambda_def_register_once lambda_def_register_once.cpp) diff --git a/tests/cpu/lambda_cast_capture.cpp b/tests/cpu/lambda_cast_capture.cpp new file mode 100644 index 00000000..f1f3efe7 --- /dev/null +++ b/tests/cpu/lambda_cast_capture.cpp @@ -0,0 +1,52 @@ +// clang-format off +// RUN: rm -rf "%t.$$.proteus" +// RUN: PROTEUS_CACHE_DIR="%t.$$.proteus" %build/lambda_cast_capture | %FILECHECK %s --check-prefixes=CHECK +// RUN: rm -rf "%t.$$.proteus" +// clang-format on + +#include + +#include + +template void run(F &&Func) { proteus::register_lambda(Func)(); } + +int main() { + int I = -5; + unsigned U = 0x80000000u; + float F = 1.5f; + double D = 2.7; + + auto IntToLong = + [=, X = (long)proteus::jit_variable(I)]() + __attribute__((annotate("jit"))) { printf("IntToLong %ld\n", X); }; + run(IntToLong); + + auto UnsignedToULong = [=, X = (unsigned long)proteus::jit_variable(U)]() + __attribute__((annotate("jit"))) { + printf("UnsignedToULong %lu\n", X); + }; + run(UnsignedToULong); + + auto FloatToDouble = + [=, X = (double)proteus::jit_variable(F)]() + __attribute__((annotate("jit"))) { printf("FloatToDouble %f\n", X); }; + run(FloatToDouble); + + auto IntToDouble = + [=, X = (double)proteus::jit_variable(I)]() + __attribute__((annotate("jit"))) { printf("IntToDouble %f\n", X); }; + run(IntToDouble); + + auto DoubleToInt = + [=, X = (int)proteus::jit_variable(D)]() + __attribute__((annotate("jit"))) { printf("DoubleToInt %d\n", X); }; + run(DoubleToInt); + + return 0; +} + +// CHECK: IntToLong -5 +// CHECK: UnsignedToULong 2147483648 +// CHECK: FloatToDouble 1.500000 +// CHECK: IntToDouble -5.000000 +// CHECK: DoubleToInt 2 diff --git a/tests/gpu/CMakeLists.txt b/tests/gpu/CMakeLists.txt index 5e9b321d..0e122d55 100644 --- a/tests/gpu/CMakeLists.txt +++ b/tests/gpu/CMakeLists.txt @@ -262,6 +262,7 @@ CREATE_GPU_TEST(shared_array shared_array.cpp) CREATE_GPU_TEST(enable_disable enable_disable.cpp) CREATE_GPU_TEST(lambda_multiple lambda_multiple.cpp) CREATE_GPU_TEST(lambda_def lambda_def.cpp) +CREATE_GPU_TEST(lambda_cast_capture lambda_cast_capture.cpp) CREATE_GPU_TEST(lambda_factory lambda_factory.cpp) CREATE_GPU_TEST(lambda_factory_forwarded lambda_factory_forwarded.cpp) CREATE_GPU_TEST(lambda_thread_privatizer lambda_thread_privatizer.cpp) diff --git a/tests/gpu/lambda_cast_capture.cpp b/tests/gpu/lambda_cast_capture.cpp new file mode 100644 index 00000000..61224def --- /dev/null +++ b/tests/gpu/lambda_cast_capture.cpp @@ -0,0 +1,63 @@ +// clang-format off +// RUN: rm -rf "%t.$$.proteus" +// RUN: PROTEUS_CACHE_DIR="%t.$$.proteus" %build/lambda_cast_capture.%ext | %FILECHECK %s --check-prefixes=CHECK +// RUN: rm -rf "%t.$$.proteus" +// clang-format on + +#include + +#include "gpu_common.h" +#include + +template +__global__ __attribute__((annotate("jit"))) void kernel(T LB) { + LB(); +} + +template void run(T &&LB) { + kernel<<<1, 1>>>(proteus::register_lambda(LB)); + gpuErrCheck(gpuDeviceSynchronize()); + fflush(stdout); +} + +int main() { + int I = -5; + unsigned U = 0x80000000u; + float F = 1.5f; + double D = 2.7; + + auto IntToLong = + [=, X = (long)proteus::jit_variable(I)] __device__() + __attribute__((annotate("jit"))) { printf("IntToLong %ld\n", X); }; + run(IntToLong); + + auto UnsignedToULong = + [=, X = (unsigned long)proteus::jit_variable(U)] __device__() + __attribute__((annotate("jit"))) { + printf("UnsignedToULong %lu\n", X); + }; + run(UnsignedToULong); + + auto FloatToDouble = + [=, X = (double)proteus::jit_variable(F)] __device__() + __attribute__((annotate("jit"))) { printf("FloatToDouble %f\n", X); }; + run(FloatToDouble); + + auto IntToDouble = + [=, X = (double)proteus::jit_variable(I)] __device__() + __attribute__((annotate("jit"))) { printf("IntToDouble %f\n", X); }; + run(IntToDouble); + + auto DoubleToInt = + [=, X = (int)proteus::jit_variable(D)] __device__() + __attribute__((annotate("jit"))) { printf("DoubleToInt %d\n", X); }; + run(DoubleToInt); + + return 0; +} + +// CHECK: IntToLong -5 +// CHECK: UnsignedToULong 2147483648 +// CHECK: FloatToDouble 1.500000 +// CHECK: IntToDouble -5.000000 +// CHECK: DoubleToInt 2