[clang][OpenMP][SPIR-V] Fix addrspace of global constants #134399

sarnex · 2025-04-04T15:22:27Z

SPIR-V has strict address space rules, constant globals cannot be in the default address space.

The OMPIRBuilder change was required for lit tests to pass, we were missing an addrspacecast.

Signed-off-by: Sarnie, Nick <[email protected]>

llvmbot · 2025-04-04T17:46:06Z

@llvm/pr-subscribers-hlsl
@llvm/pr-subscribers-clang-codegen

@llvm/pr-subscribers-clang

Author: Nick Sarnie (sarnex)

Changes

SPIR-V has strict address space rules, globals cannot be in the default address space. Normal globals should be in addrspace 1 (which is what we get from opencl_global in the SPIR-V address space map) and 2 for global constants (opencl_constant in the SPIR-V address space map)

This is similar to what was done for HIPSPV.

The OMPIRBuilder change was required for lit tests to pass, we were missing an addrspacecast.

Full diff: https://github.com/llvm/llvm-project/pull/134399.diff

3 Files Affected:

(modified) clang/lib/CodeGen/CodeGenModule.cpp (+9)
(added) clang/test/OpenMP/spirv_target_addrspace.c (+20)
(modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+6)

diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 8f9cf965af2b9..cc6d726445cbb 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -5384,6 +5384,11 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
     LangAS AS;
     if (OpenMPRuntime->hasAllocateAttributeForGlobalVar(D, AS))
       return AS;
+    if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
+      // SPIR-V globals should map to CrossWorkGroup instead of default
+      // AS, as generic/no address space is invalid. This is similar
+      // to what is done for HIPSPV.
+      return LangAS::opencl_global;
   }
   return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);
 }
@@ -5402,6 +5407,10 @@ LangAS CodeGenModule::GetGlobalConstantAddressSpace() const {
     // UniformConstant storage class is not viable as pointers to it may not be
     // casted to Generic pointers which are used to model HIP's "flat" pointers.
     return LangAS::cuda_device;
+  if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
+    // OpenMP SPIR-V global constants should map to UniformConstant, different
+    // from the HIPSPV case above.
+    return LangAS::opencl_constant;
   if (auto AS = getTarget().getConstantAddressSpace())
     return *AS;
   return LangAS::Default;
diff --git a/clang/test/OpenMP/spirv_target_addrspace.c b/clang/test/OpenMP/spirv_target_addrspace.c
new file mode 100644
index 0000000000000..8430a30efe0c8
--- /dev/null
+++ b/clang/test/OpenMP/spirv_target_addrspace.c
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64 -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -O0 -fopenmp -fopenmp-targets=spirv64 -fopenmp-is-target-device -triple spirv64 -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - | FileCheck %s
+
+extern int printf(char[]);
+
+#pragma omp declare target
+// CHECK: @global = addrspace(1) global i32 0, align 4
+// CHECK: @.str = private unnamed_addr addrspace(2) constant [4 x i8] c"foo\00", align 1
+int global = 0;
+#pragma omp end declare target
+int main() {
+  // CHECK: = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @__omp_offloading_{{.*}}_kernel_environment to ptr), ptr %{{.*}})
+#pragma omp target
+  {
+    for(int i = 0; i < 5; i++)
+      global++;
+    printf("foo");
+  }
+  return global;
+}
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 68b1fa42934ad..998702c1af3cd 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -6295,6 +6295,12 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetInit(
           : ConstantExpr::getAddrSpaceCast(KernelEnvironmentGV,
                                            KernelEnvironmentPtr);
   Value *KernelLaunchEnvironment = DebugKernelWrapper->getArg(0);
+  Type *KernelLaunchEnvParamTy = Fn->getFunctionType()->getParamType(1);
+  KernelLaunchEnvironment =
+      KernelLaunchEnvironment->getType() == KernelLaunchEnvParamTy
+          ? KernelLaunchEnvironment
+          : Builder.CreateAddrSpaceCast(KernelLaunchEnvironment,
+                                        KernelLaunchEnvParamTy);
   CallInst *ThreadKind =
       Builder.CreateCall(Fn, {KernelEnvironment, KernelLaunchEnvironment});

llvmbot · 2025-04-04T17:46:06Z

@llvm/pr-subscribers-flang-openmp

Author: Nick Sarnie (sarnex)

Changes

SPIR-V has strict address space rules, globals cannot be in the default address space. Normal globals should be in addrspace 1 (which is what we get from opencl_global in the SPIR-V address space map) and 2 for global constants (opencl_constant in the SPIR-V address space map)

This is similar to what was done for HIPSPV.

The OMPIRBuilder change was required for lit tests to pass, we were missing an addrspacecast.

Full diff: https://github.com/llvm/llvm-project/pull/134399.diff

3 Files Affected:

(modified) clang/lib/CodeGen/CodeGenModule.cpp (+9)
(added) clang/test/OpenMP/spirv_target_addrspace.c (+20)
(modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+6)

diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 8f9cf965af2b9..cc6d726445cbb 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -5384,6 +5384,11 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
     LangAS AS;
     if (OpenMPRuntime->hasAllocateAttributeForGlobalVar(D, AS))
       return AS;
+    if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
+      // SPIR-V globals should map to CrossWorkGroup instead of default
+      // AS, as generic/no address space is invalid. This is similar
+      // to what is done for HIPSPV.
+      return LangAS::opencl_global;
   }
   return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);
 }
@@ -5402,6 +5407,10 @@ LangAS CodeGenModule::GetGlobalConstantAddressSpace() const {
     // UniformConstant storage class is not viable as pointers to it may not be
     // casted to Generic pointers which are used to model HIP's "flat" pointers.
     return LangAS::cuda_device;
+  if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
+    // OpenMP SPIR-V global constants should map to UniformConstant, different
+    // from the HIPSPV case above.
+    return LangAS::opencl_constant;
   if (auto AS = getTarget().getConstantAddressSpace())
     return *AS;
   return LangAS::Default;
diff --git a/clang/test/OpenMP/spirv_target_addrspace.c b/clang/test/OpenMP/spirv_target_addrspace.c
new file mode 100644
index 0000000000000..8430a30efe0c8
--- /dev/null
+++ b/clang/test/OpenMP/spirv_target_addrspace.c
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64 -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -O0 -fopenmp -fopenmp-targets=spirv64 -fopenmp-is-target-device -triple spirv64 -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - | FileCheck %s
+
+extern int printf(char[]);
+
+#pragma omp declare target
+// CHECK: @global = addrspace(1) global i32 0, align 4
+// CHECK: @.str = private unnamed_addr addrspace(2) constant [4 x i8] c"foo\00", align 1
+int global = 0;
+#pragma omp end declare target
+int main() {
+  // CHECK: = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @__omp_offloading_{{.*}}_kernel_environment to ptr), ptr %{{.*}})
+#pragma omp target
+  {
+    for(int i = 0; i < 5; i++)
+      global++;
+    printf("foo");
+  }
+  return global;
+}
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 68b1fa42934ad..998702c1af3cd 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -6295,6 +6295,12 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetInit(
           : ConstantExpr::getAddrSpaceCast(KernelEnvironmentGV,
                                            KernelEnvironmentPtr);
   Value *KernelLaunchEnvironment = DebugKernelWrapper->getArg(0);
+  Type *KernelLaunchEnvParamTy = Fn->getFunctionType()->getParamType(1);
+  KernelLaunchEnvironment =
+      KernelLaunchEnvironment->getType() == KernelLaunchEnvParamTy
+          ? KernelLaunchEnvironment
+          : Builder.CreateAddrSpaceCast(KernelLaunchEnvironment,
+                                        KernelLaunchEnvParamTy);
   CallInst *ThreadKind =
       Builder.CreateCall(Fn, {KernelEnvironment, KernelLaunchEnvironment});

jhuber6 · 2025-04-04T17:48:17Z

clang/lib/CodeGen/CodeGenModule.cpp

+      // SPIR-V globals should map to CrossWorkGroup instead of default
+      // AS, as generic/no address space is invalid. This is similar
+      // to what is done for HIPSPV.
+      return LangAS::opencl_global;
  }
  return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);


Wouldn't this logic fit in here?

jhuber6 · 2025-04-04T17:52:46Z

clang/lib/CodeGen/CodeGenModule.cpp

@@ -5384,6 +5384,11 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
    LangAS AS;
    if (OpenMPRuntime->hasAllocateAttributeForGlobalVar(D, AS))
      return AS;
+    if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())


Is this really OpenMP specific? Sounds like a target info thing to me.

Yes I am a bit confused as to why this is necessary, DataLayout already encodes that global is AS1. If you're seeing globals end up in generic (I am excluding llvm.used and llvm.compiler.used here, since they are special and should be in generic/0) it might just be a case where CodeGen has a subtle bug. Could you please say a bit more as to what is motivating this change? Thank you!

Thanks for the feedback guys, the SPIR-V address space stuff is a total nightmare so I'll take any feedback I can get.

Here's the problem I'm trying to solve. For the code in the test I have:

extern int printf(char[]); #pragma omp declare target int global = 0; #pragma omp end declare target int main() { #pragma omp target { for(int i = 0; i < 5; i++) global++; printf("foo"); } return global; }

Currently we get this IR

@global = global i32 0, align 4 @.str = private unnamed_addr constant [4 x i8] c"foo\00", align 1

Clearly the address space of both is wrong, addrspace(0) is not valid in SPIR-V for globals.

I think doing it in the target itself is much better, let me update the PR doing that, thanks.

This feels like a spot where we are missing something in Clang - the string should've at least been AS1; some time ago I had a pop at fixing a bunch of places in CodeGen where we just used 0 / unqual rather than getting the GlobalVar AS or the Constant AS, but issues clearly remain - I think we should try to address this in Clang. Are you seeing the above with spirv64-unknown-unknown?

Yep, I see the problem with the string even with pure spirv64-unknown-unknown. Repro:

extern int printf(const char*); int main() { printf("foo"); return 0; }

clang++ -cc1 -triple spirv64-unknown-unknown -emit-llvm test.cpp -o -

; ModuleID = 'test.cpp' source_filename = "test.cpp" target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" target triple = "spirv64-unknown-unknown" @.str = private unnamed_addr constant [4 x i8] c"foo\00", align 1 ; Function Attrs: mustprogress noinline norecurse nounwind optnone define noundef i32 @main() #0 { entry: %retval = alloca i32, align 4 store i32 0, ptr %retval, align 4 %call = call spir_func noundef i32 @_Z6printfPKc(ptr noundef @.str) ret i32 0 } declare spir_func noundef i32 @_Z6printfPKc(ptr noundef) #1 attributes #0 = { mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } attributes #1 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" } !llvm.module.flags = !{!0} !llvm.ident = !{!1} !0 = !{i32 1, !"wchar_size", i32 4} !1 = !{!"clang version 21.0.0git (https://github.com/llvm/llvm-project.git 06bfbba877c26630b6c5b0ffef7f6623aa2e9ee8)"}

Here's where we get the addrspace from:

unsigned AddrSpace = CGM.getContext().getTargetAddressSpace( CGM.GetGlobalConstantAddressSpace());

In GetGlobalConstantAddressSpace, we do

if (auto AS = getTarget().getConstantAddressSpace()) return *AS;

and since there's no override for SPIR-V we just get the default:

virtual std::optional<LangAS> getConstantAddressSpace() const { return LangAS::Default; }

If you see something wrong in this callstack let me know, I'm happy to fix it!

This seems ok. I would suggest that we consider, in the override, returning opencl_constant (2) only for OCL, and otherwise returning the global var AS (1), to prevent crashing into the invalid cast problem.

Sure, let me try that. Probably AS1 will be fine for my use case.

BTW codegen already does the right thing because of this code in GetGlobalConstantAddressSpace, and CodeGenOpenCL/str_literals.cl already locks it down, so the OpenCL part of my change is basically NFC.

AlexVlx · 2025-04-04T18:36:26Z

clang/lib/CodeGen/CodeGenModule.cpp

+  if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
+    // OpenMP SPIR-V global constants should map to UniformConstant, different
+    // from the HIPSPV case above.
+    return LangAS::opencl_constant;


I believe that whilst this makes sense you might run into obnoxious issues where valid source ends up generating a SPIR-V invalid cast to/from generic/constant, which will fail in the translator / fail SPIR-V validation (HIP code runs into this). I think we need to relax this restriction at least in the translator / BE, or potentially extend SPIR-V itself in this direction.

Yeah I hit something similar when trying to use addrspace(4) for generic. For constants it seems to be working okay for now, hopefully it's okay with you if I try this (but moved into the target) for now and then extend the solution if a problem comes up.

Signed-off-by: Sarnie, Nick <[email protected]>

jhuber6 · 2025-04-08T14:50:10Z

clang/lib/Basic/Targets/SPIR.h

@@ -37,8 +37,8 @@ static const unsigned SPIRDefIsPrivMap[] = {
    0, // cuda_device
    0, // cuda_constant
    0, // cuda_shared
-    // SYCL address space values for this map are dummy
-    0, // sycl_global
+    // Most SYCL address space values for this map are dummy


Anyone know why this is the case?

Probably because for SYCL it doesn't use that map, it uses this one because of this logic.

This is a hack for OpenCL, and is not meant to work otherwise, it should only be used for OCL (looks like an unfortunate import from some brainrot we put in AMDGPU). We've made an effort to fix this recently, so perhaps we could look at moving SPIR-V away from it as well, see #112442 and its children.

@AlexVlx Sorry do you mean the map with default AS as 0 should be OCL only and the one with default AS as 4 should be used otherwise even if temporary? If so I can update this PR to do that instead and deal with the fallout sooner rather than later.

Yes, not only should it be OCL only, but it probably should only be OCL with no generic AS support only (please check out the PR I linked where we switched that over for AMDGPU). PrivateAsDefault was/is a bad hack. Note that you might want to fork the AS map switch into a different PR, as the fallout might end up a bit of a slog. E.g., this is still stuck in limbo #113930, and it will bite.

Got it, so are you okay with this PR if I remove the part changing the default AS 0 map (and just have the constant global addrspace change)?

Oh, yes, apologies for the segue, this LGTM in general, thanks.

thanks for the review, should have a PR for the map fix shortly

AlexVlx

LGTM, thanks!

Signed-off-by: Sarnie, Nick <[email protected]>

SPIR-V has strict address space rules, constant globals cannot be in the default address space. The OMPIRBuilder change was required for lit tests to pass, we were missing an addrspacecast. --------- Signed-off-by: Sarnie, Nick <[email protected]>

alexfh · 2025-04-15T23:56:44Z

We've started seeing errors similar to this one after this commit:

Can't translate llvm instruction:
 Global variable can not have Function storage class. Consider setting a proper address space.
 Original LLVM value:
@.str.2 = private unnamed_addr constant [19 x i8] c"%s at %s:%d in %s\0A\00", align 1

Is this the expected effect of the change?

sarnex · 2025-04-16T14:08:02Z

No, the string should have a non-zero address space. Do you have a reproduction?

alexfh · 2025-04-16T23:20:39Z

This is being looked into. If this turns out to be a problem in LLVM rather than in the SPIRV-LLVM-Translator, I'll let you know here.

SPIR-V has strict address space rules, constant globals cannot be in the default address space. The OMPIRBuilder change was required for lit tests to pass, we were missing an addrspacecast. --------- Signed-off-by: Sarnie, Nick <[email protected]>

sarnex · 2025-04-18T14:42:48Z

@alexfh FYI I expect the problem is in LLVM and not the translator, but it's not that this change is totally wrong it's just a missed case as the goal of this commit was to add the address space, not remove it, which seems to be what happened based on that error.

ShangwuYao · 2025-04-20T23:22:45Z

This test reproduces the issue above:

// RUN: %clang_cc1 -fcuda-is-device -triple spirv32 -o - -emit-llvm -x cuda %s  | FileCheck %s
// RUN: %clang_cc1 -fcuda-is-device -triple spirv64 -o - -emit-llvm -x cuda %s  | FileCheck %s

// CHECK: @.str = private unnamed_addr addrspace(4) constant [13 x i8] c"Hello World\0A\00", align 1 

extern "C" __attribute__((device)) int printf(const char* format, ...);

__attribute__((global)) void printf_kernel() {
  printf("Hello World\n");
}

Could you also add the test case as test/CodeGenCUDASPIRV/printf.cu or something? Thanks!!

sarnex · 2025-04-21T14:22:55Z

Investigating now, thanks!

sarnex · 2025-04-21T18:00:45Z

This will be fixed by my existing open PR #135251

I will add that test case and add you guys as reviewers, thanks and sorry for the trouble.

This fixes a CUDA SPIR-V regression introduced in #134399. --------- Signed-off-by: Sarnie, Nick <[email protected]>

…e 1 (#136753) This fixes a CUDA SPIR-V regression introduced in llvm/llvm-project#134399. --------- Signed-off-by: Sarnie, Nick <[email protected]>

…36753) This fixes a CUDA SPIR-V regression introduced in llvm#134399. --------- Signed-off-by: Sarnie, Nick <[email protected]>

[clang][OpenMP][SPIR-V] Fix addrspace of globals and global constants

3812b13

Signed-off-by: Sarnie, Nick <[email protected]>

sarnex force-pushed the addrspace branch from 1eb8258 to 3812b13 Compare April 4, 2025 15:46

sarnex marked this pull request as ready for review April 4, 2025 17:45

llvmbot added clang Clang issues not falling into any other category clang:codegen IR generation bugs: mangling, exceptions, etc. flang:openmp clang:openmp OpenMP related changes to Clang labels Apr 4, 2025

sarnex requested review from svenvh, MrSidims, AlexVlx and jhuber6 April 4, 2025 17:46

jhuber6 reviewed Apr 4, 2025

View reviewed changes

AlexVlx reviewed Apr 4, 2025

View reviewed changes

do it in the target

93ff204

Signed-off-by: Sarnie, Nick <[email protected]>

llvmbot added clang:frontend Language frontend issues, e.g. anything involving "Sema" HLSL HLSL Language Support labels Apr 7, 2025

github-project-automation bot added this to HLSL Support Apr 7, 2025

sarnex added 2 commits April 7, 2025 11:31

whitespace

06bfbba

Signed-off-by: Sarnie, Nick <[email protected]>

opencl_constant only for ocl

c4941ea

Signed-off-by: Sarnie, Nick <[email protected]>

sarnex requested review from AlexVlx and jhuber6 April 8, 2025 14:26

jhuber6 reviewed Apr 8, 2025

View reviewed changes

AlexVlx approved these changes Apr 8, 2025

View reviewed changes

revert global changes

c9a24d0

Signed-off-by: Sarnie, Nick <[email protected]>

sarnex changed the title ~~[clang][OpenMP][SPIR-V] Fix addrspace of globals and global constants~~ [clang][OpenMP][SPIR-V] Fix addrspace of global constants Apr 8, 2025

sarnex merged commit 68ee56d into llvm:main Apr 9, 2025
11 checks passed

sarnex mentioned this pull request Apr 10, 2025

[clang][OpenMP][SPIR-V] Fix AS of globals and set the default AS to 4 #135251

Draft

sarnex mentioned this pull request Apr 22, 2025

[clang][SPIR-V] Addrspace of opencl_global should always be 1 #136753

Merged

sarnex added a commit that referenced this pull request Apr 24, 2025

[clang][SPIR-V] Addrspace of opencl_global should always be 1 (#136753)

52a9649

This fixes a CUDA SPIR-V regression introduced in #134399. --------- Signed-off-by: Sarnie, Nick <[email protected]>

damyanp moved this to Closed in HLSL Support Apr 25, 2025

[clang][OpenMP][SPIR-V] Fix addrspace of global constants #134399

[clang][OpenMP][SPIR-V] Fix addrspace of global constants #134399

Uh oh!

Conversation

sarnex commented Apr 4, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Apr 4, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Apr 4, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

sarnex Apr 7, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

sarnex Apr 8, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

sarnex Apr 8, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

AlexVlx left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

alexfh commented Apr 15, 2025

Uh oh!

sarnex commented Apr 16, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

alexfh commented Apr 16, 2025

Uh oh!

sarnex commented Apr 18, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

ShangwuYao commented Apr 20, 2025

Uh oh!

sarnex commented Apr 21, 2025

Uh oh!

sarnex commented Apr 21, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

sarnex commented Apr 4, 2025 •

edited

Loading

llvmbot commented Apr 4, 2025 •

edited

Loading

sarnex Apr 7, 2025 •

edited

Loading

sarnex Apr 8, 2025 •

edited

Loading

sarnex Apr 8, 2025 •

edited

Loading

sarnex commented Apr 16, 2025 •

edited

Loading

sarnex commented Apr 18, 2025 •

edited

Loading

sarnex commented Apr 21, 2025 •

edited

Loading