================
@@ -0,0 +1,85 @@
+#include "cuda.h"
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip
-fhip-new-launch-api \
+// RUN: -fclangir -I%S/../CodeGenCUDA/Inputs/ -emit-cir %s -o - \
+// RUN: | FileCheck --check-prefix=CIR %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip
-fhip-new-launch-api \
+// RUN: -fclangir -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=LLVM %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip
-fhip-new-launch-api \
+// RUN: -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=LLVM %s
+
+// The host-side kernel handle takes the linkage of the kernel's device stub
and
+// the visibility of the kernel's declaration. Kernels with internal linkage
+// must get an internal handle in every TU that uses them as an external one
+// causes multiple-definition link errors.
+
+template <class T> static __global__ void static_tmpl(T *p) {}
+template <class T> __global__ void tmpl(T *p) {}
+static __global__ void static_kernel(int *p) {}
+namespace {
+__global__ void anon_kernel(int *p) {}
+}
+__global__ void ext_kernel(int *p) {}
+__attribute__((visibility("hidden"))) __global__ void hidden_kernel(int *p) {}
+
+// Explicit instantiation definition: weak_odr, in a comdat.
+template <class T> __global__ void inst(T *p) {}
+template __global__ void inst<float>(float *p);
+
+// Explicit specialization of a template that is only declared: strong
+// external, no comdat.
+template <class T> __global__ void decl_only(T *p);
+template <> __global__ void decl_only<int>(int *p) {}
+
+// Referenced before its definition, so the handle is created before the
+// stub's linkage is known.
+static __global__ void fwd_kernel(int *p);
+
+// Defined in another TU, so the stub is never emitted here; the visibility
+// must come from the declaration.
+__attribute__((visibility("hidden"))) __global__ void hidden_ext_kernel(int
*p);
+
+// Internal linkage wins over the explicit visibility.
+static __attribute__((visibility("hidden"))) __global__ void
+static_hidden_kernel(int *p) {}
+
+void launch(int *p) {
+ static_tmpl<int><<<1, 1>>>(p);
+ tmpl<int><<<1, 1>>>(p);
+ static_kernel<<<1, 1>>>(p);
+ anon_kernel<<<1, 1>>>(p);
+ ext_kernel<<<1, 1>>>(p);
+ hidden_kernel<<<1, 1>>>(p);
+ decl_only<int><<<1, 1>>>(p);
+ fwd_kernel<<<1, 1>>>(p);
+ hidden_ext_kernel<<<1, 1>>>(p);
+ static_hidden_kernel<<<1, 1>>>(p);
+}
+
+static __global__ void fwd_kernel(int *p) {}
+
+// CIR-DAG: cir.global constant external @_Z10ext_kernelPi =
#cir.global_view<@_Z25__device_stub__ext_kernelPi>
+// CIR-DAG: cir.global hidden constant external @_Z13hidden_kernelPi =
#cir.global_view<@_Z28__device_stub__hidden_kernelPi>
+// CIR-DAG: cir.global "private" constant internal dso_local
@_ZL11static_tmplIiEvPT_ =
#cir.global_view<@_ZL26__device_stub__static_tmplIiEvPT_>
+// CIR-DAG: cir.global constant linkonce_odr comdat @_Z4tmplIiEvPT_ =
#cir.global_view<@_Z19__device_stub__tmplIiEvPT_>
+// CIR-DAG: cir.global "private" constant internal dso_local
@_ZL13static_kernelPi = #cir.global_view<@_ZL28__device_stub__static_kernelPi>
+// CIR-DAG: cir.global "private" constant internal dso_local
@_ZN12_GLOBAL__N_111anon_kernelEPi =
#cir.global_view<@_ZN12_GLOBAL__N_126__device_stub__anon_kernelEPi>
+// CIR-DAG: cir.global "private" constant internal dso_local
@_ZL10fwd_kernelPi = #cir.global_view<@_ZL25__device_stub__fwd_kernelPi>
+// CIR-DAG: cir.global constant weak_odr comdat @_Z4instIfEvPT_ =
#cir.global_view<@_Z19__device_stub__instIfEvPT_>
+// CIR-DAG: cir.global constant external @_Z9decl_onlyIiEvPT_ =
#cir.global_view<@_Z24__device_stub__decl_onlyIiEvPT_>
+// CIR-DAG: cir.global "private" hidden constant external
@_Z17hidden_ext_kernelPi :
+// CIR-DAG: cir.global "private" constant internal dso_local
@_ZL20static_hidden_kernelPi =
#cir.global_view<@_ZL35__device_stub__static_hidden_kernelPi>
----------------
Fznamznon wrote:
dso_local and comdat exist on CIR level but not in LLVM, is that expected?
https://github.com/llvm/llvm-project/pull/227267
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
