diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 2fb758bc75d6..6e3ae96b3afc 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -9681,7 +9681,8 @@ public: /// Check if the specified variable is captured by 'target' directive. /// \param Level Relative level of nested OpenMP construct for that the check /// is performed. - bool isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level) const; + bool isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level, + unsigned CaptureLevel) const; ExprResult PerformOpenMPImplicitIntegerConversion(SourceLocation OpLoc, Expr *Op); diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index ea4b93ee6a5a..67b68ffb5eda 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -16280,8 +16280,10 @@ bool Sema::tryCaptureVariable( captureVariablyModifiedType(Context, QTy, OuterRSI); } } - bool IsTargetCap = !IsOpenMPPrivateDecl && - isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel); + bool IsTargetCap = + !IsOpenMPPrivateDecl && + isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel, + RSI->OpenMPCaptureLevel); // When we detect target captures we are looking from inside the // target region, therefore we need to propagate the capture from the // enclosing region. Therefore, the capture is not initially nested. diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 3fce0e27e9b3..72afe63749ac 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -31,6 +31,7 @@ #include "clang/Sema/SemaInternal.h" #include "llvm/ADT/IndexedMap.h" #include "llvm/ADT/PointerEmbeddedInt.h" +#include "llvm/ADT/STLExtras.h" #include "llvm/Frontend/OpenMP/OMPConstants.h" using namespace clang; using namespace llvm::omp; @@ -2010,7 +2011,23 @@ VarDecl *Sema::isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo, // if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) return nullptr; - return VD; + CapturedRegionScopeInfo *CSI = nullptr; + for (FunctionScopeInfo *FSI : llvm::drop_begin( + llvm::reverse(FunctionScopes), + CheckScopeInfo ? (FunctionScopes.size() - (StopAt + 1)) : 0)) { + if (!isa(FSI)) + return nullptr; + if (auto *RSI = dyn_cast(FSI)) + if (RSI->CapRegionKind == CR_OpenMP) { + CSI = RSI; + break; + } + } + SmallVector Regions; + getOpenMPCaptureRegions(Regions, + DSAStack->getDirective(CSI->OpenMPLevel)); + if (Regions[CSI->OpenMPCaptureLevel] != OMPD_task) + return VD; } } @@ -2151,15 +2168,18 @@ void Sema::setOpenMPCaptureKind(FieldDecl *FD, const ValueDecl *D, FD->addAttr(OMPCaptureKindAttr::CreateImplicit(Context, OMPC)); } -bool Sema::isOpenMPTargetCapturedDecl(const ValueDecl *D, - unsigned Level) const { +bool Sema::isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level, + unsigned CaptureLevel) const { assert(LangOpts.OpenMP && "OpenMP is not allowed"); // Return true if the current level is no longer enclosed in a target region. + SmallVector Regions; + getOpenMPCaptureRegions(Regions, DSAStack->getDirective(Level)); const auto *VD = dyn_cast(D); return VD && !VD->hasLocalStorage() && DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective, - Level); + Level) && + Regions[CaptureLevel] != OMPD_task; } void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; } diff --git a/clang/test/OpenMP/target_depend_codegen.cpp b/clang/test/OpenMP/target_depend_codegen.cpp index e2810f946b04..e5473dad8055 100644 --- a/clang/test/OpenMP/target_depend_codegen.cpp +++ b/clang/test/OpenMP/target_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], diff --git a/clang/test/OpenMP/target_messages.cpp b/clang/test/OpenMP/target_messages.cpp index 5f22fbcaf606..ad04e9306cb0 100644 --- a/clang/test/OpenMP/target_messages.cpp +++ b/clang/test/OpenMP/target_messages.cpp @@ -112,4 +112,12 @@ int main(int argc, char **argv) { return 0; } + +template struct a { static bool b; }; +template ::b> void e(c) { // expected-note {{candidate template ignored: substitution failure [with c = int]: non-type template argument is not a constant expression}} +#pragma omp target + { + int d ; e(d); // expected-error {{no matching function for call to 'e'}} + } +} #endif diff --git a/clang/test/OpenMP/target_parallel_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_depend_codegen.cpp index dceb585b14c4..de7e37efc484 100644 --- a/clang/test/OpenMP/target_parallel_depend_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], diff --git a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp index 6b2325592c6a..73a5cb1a8beb 100644 --- a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], diff --git a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp index 74ff678316d7..6b9950428dba 100644 --- a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], diff --git a/clang/test/OpenMP/target_simd_depend_codegen.cpp b/clang/test/OpenMP/target_simd_depend_codegen.cpp index 0fb75b0b7fc4..0c28c43dc734 100644 --- a/clang/test/OpenMP/target_simd_depend_codegen.cpp +++ b/clang/test/OpenMP/target_simd_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], diff --git a/clang/test/OpenMP/target_teams_depend_codegen.cpp b/clang/test/OpenMP/target_teams_depend_codegen.cpp index 85fc5e297ce8..4cc402c89344 100644 --- a/clang/test/OpenMP/target_teams_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], diff --git a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp index 37b80b097224..c8694761bf7b 100644 --- a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp index b136e7b75e4e..df1ba2772afc 100644 --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp index 16c73e7406d2..69cdd99e3f46 100644 --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], diff --git a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp index 4fffbc50a5eb..06fb3daa3240 100644 --- a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp @@ -209,7 +209,7 @@ int foo(int n) { // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], @@ -223,7 +223,7 @@ int foo(int n) { // CHECK: call void (i8*, ...) % // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % -// CHECK: [[BP1_I32:%.+]] = load i32, i32* % +// CHECK: [[BP1_I32:%.+]] = load i32, i32* @ // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],