Regenerate PTX test

Overwrites all generated PTX tests and runs `pre-commit run --all-files`
NVIDIA · Nov 25, 2024 · 59a6584 · 59a6584
1 parent 0b36a7d
commit 59a6584
Show file tree

Hide file tree

Showing 49 changed files with 42 additions and 42 deletions.
diff --git a/...xx/cuda/ptx/generated/barrier_cluster.inc → ...acxx/cuda/ptx/generated/barrier_cluster.h b/...xx/cuda/ptx/generated/barrier_cluster.inc → ...acxx/cuda/ptx/generated/barrier_cluster.h
diff --git a/...acxx/cuda/ptx/generated/cp_async_bulk.inc → ...udacxx/cuda/ptx/generated/cp_async_bulk.h b/...acxx/cuda/ptx/generated/cp_async_bulk.inc → ...udacxx/cuda/ptx/generated/cp_async_bulk.h
diff --git a/.../generated/cp_async_bulk_commit_group.inc → ...tx/generated/cp_async_bulk_commit_group.h b/.../generated/cp_async_bulk_commit_group.inc → ...tx/generated/cp_async_bulk_commit_group.h
diff --git a/...ptx/generated/cp_async_bulk_multicast.inc → ...a/ptx/generated/cp_async_bulk_multicast.h b/...ptx/generated/cp_async_bulk_multicast.inc → ...a/ptx/generated/cp_async_bulk_multicast.h
diff --git a/...da/ptx/generated/cp_async_bulk_tensor.inc → ...cuda/ptx/generated/cp_async_bulk_tensor.h b/...da/ptx/generated/cp_async_bulk_tensor.inc → ...cuda/ptx/generated/cp_async_bulk_tensor.h
diff --git a/...erated/cp_async_bulk_tensor_multicast.inc → ...enerated/cp_async_bulk_tensor_multicast.h b/...erated/cp_async_bulk_tensor_multicast.inc → ...enerated/cp_async_bulk_tensor_multicast.h
diff --git a/...tx/generated/cp_async_bulk_wait_group.inc → .../ptx/generated/cp_async_bulk_wait_group.h b/...tx/generated/cp_async_bulk_wait_group.inc → .../ptx/generated/cp_async_bulk_wait_group.h
diff --git a/...da/ptx/generated/cp_reduce_async_bulk.inc → ...cuda/ptx/generated/cp_reduce_async_bulk.h b/...da/ptx/generated/cp_reduce_async_bulk.inc → ...cuda/ptx/generated/cp_reduce_async_bulk.h
diff --git a/...x/generated/cp_reduce_async_bulk_bf16.inc → ...ptx/generated/cp_reduce_async_bulk_bf16.h b/...x/generated/cp_reduce_async_bulk_bf16.inc → ...ptx/generated/cp_reduce_async_bulk_bf16.h
@@ -1,6 +1,6 @@
 __global__ void test_cp_reduce_async_bulk_bf16(void** fn_ptr)
 {
-#  if __cccl_ptx_isa >= 800
+#if __cccl_ptx_isa >= 800
   NV_IF_TARGET(
     NV_PROVIDES_SM_90,
     (
@@ -12,9 +12,9 @@ __global__ void test_cp_reduce_async_bulk_bf16(void** fn_ptr)
                                __nv_bfloat16*,
                                const __nv_bfloat16*,
                                uint32_t)>(cuda::ptx::cp_reduce_async_bulk));));
-#  endif // __cccl_ptx_isa >= 800
+#endif // __cccl_ptx_isa >= 800
 
-#  if __cccl_ptx_isa >= 800
+#if __cccl_ptx_isa >= 800
   NV_IF_TARGET(
     NV_PROVIDES_SM_90,
     (
@@ -26,9 +26,9 @@ __global__ void test_cp_reduce_async_bulk_bf16(void** fn_ptr)
                                __nv_bfloat16*,
                                const __nv_bfloat16*,
                                uint32_t)>(cuda::ptx::cp_reduce_async_bulk));));
-#  endif // __cccl_ptx_isa >= 800
+#endif // __cccl_ptx_isa >= 800
 
-#  if __cccl_ptx_isa >= 800
+#if __cccl_ptx_isa >= 800
   NV_IF_TARGET(
     NV_PROVIDES_SM_90,
     (
@@ -40,5 +40,5 @@ __global__ void test_cp_reduce_async_bulk_bf16(void** fn_ptr)
                                __nv_bfloat16*,
                                const __nv_bfloat16*,
                                uint32_t)>(cuda::ptx::cp_reduce_async_bulk));));
-#  endif // __cccl_ptx_isa >= 800
+#endif // __cccl_ptx_isa >= 800
 }
diff --git a/...tx/generated/cp_reduce_async_bulk_f16.inc → .../ptx/generated/cp_reduce_async_bulk_f16.h b/...tx/generated/cp_reduce_async_bulk_f16.inc → .../ptx/generated/cp_reduce_async_bulk_f16.h
@@ -1,6 +1,6 @@
 __global__ void test_cp_reduce_async_bulk_f16(void** fn_ptr)
 {
-#  if __cccl_ptx_isa >= 800
+#if __cccl_ptx_isa >= 800
   NV_IF_TARGET(
     NV_PROVIDES_SM_90,
     (
@@ -9,9 +9,9 @@ __global__ void test_cp_reduce_async_bulk_f16(void** fn_ptr)
           static_cast<void (*)(
             cuda::ptx::space_global_t, cuda::ptx::space_shared_t, cuda::ptx::op_min_t, __half*, const __half*, uint32_t)>(
             cuda::ptx::cp_reduce_async_bulk));));
-#  endif // __cccl_ptx_isa >= 800
+#endif // __cccl_ptx_isa >= 800
 
-#  if __cccl_ptx_isa >= 800
+#if __cccl_ptx_isa >= 800
   NV_IF_TARGET(
     NV_PROVIDES_SM_90,
     (
@@ -20,9 +20,9 @@ __global__ void test_cp_reduce_async_bulk_f16(void** fn_ptr)
           static_cast<void (*)(
             cuda::ptx::space_global_t, cuda::ptx::space_shared_t, cuda::ptx::op_max_t, __half*, const __half*, uint32_t)>(
             cuda::ptx::cp_reduce_async_bulk));));
-#  endif // __cccl_ptx_isa >= 800
+#endif // __cccl_ptx_isa >= 800
 
-#  if __cccl_ptx_isa >= 800
+#if __cccl_ptx_isa >= 800
   NV_IF_TARGET(
     NV_PROVIDES_SM_90,
     (
@@ -31,5 +31,5 @@ __global__ void test_cp_reduce_async_bulk_f16(void** fn_ptr)
           static_cast<void (*)(
             cuda::ptx::space_global_t, cuda::ptx::space_shared_t, cuda::ptx::op_add_t, __half*, const __half*, uint32_t)>(
             cuda::ptx::cp_reduce_async_bulk));));
-#  endif // __cccl_ptx_isa >= 800
+#endif // __cccl_ptx_isa >= 800
 }
diff --git a/...generated/cp_reduce_async_bulk_tensor.inc → ...x/generated/cp_reduce_async_bulk_tensor.h b/...generated/cp_reduce_async_bulk_tensor.inc → ...x/generated/cp_reduce_async_bulk_tensor.h
diff --git a/...t/libcudacxx/cuda/ptx/generated/fence.inc → ...est/libcudacxx/cuda/ptx/generated/fence.h b/...t/libcudacxx/cuda/ptx/generated/fence.inc → ...est/libcudacxx/cuda/ptx/generated/fence.h
diff --git a/...uda/ptx/generated/fence_mbarrier_init.inc → .../cuda/ptx/generated/fence_mbarrier_init.h b/...uda/ptx/generated/fence_mbarrier_init.inc → .../cuda/ptx/generated/fence_mbarrier_init.h
diff --git a/.../cuda/ptx/generated/fence_proxy_alias.inc → ...xx/cuda/ptx/generated/fence_proxy_alias.h b/.../cuda/ptx/generated/fence_proxy_alias.inc → ...xx/cuda/ptx/generated/fence_proxy_alias.h
diff --git a/.../cuda/ptx/generated/fence_proxy_async.inc → ...xx/cuda/ptx/generated/fence_proxy_async.h b/.../cuda/ptx/generated/fence_proxy_async.inc → ...xx/cuda/ptx/generated/fence_proxy_async.h
diff --git a/...nerated/fence_proxy_tensormap_generic.inc → ...generated/fence_proxy_tensormap_generic.h b/...nerated/fence_proxy_tensormap_generic.inc → ...generated/fence_proxy_tensormap_generic.h
diff --git a/...ibcudacxx/cuda/ptx/generated/get_sreg.inc → .../libcudacxx/cuda/ptx/generated/get_sreg.h b/...ibcudacxx/cuda/ptx/generated/get_sreg.inc → .../libcudacxx/cuda/ptx/generated/get_sreg.h
diff --git a/...cudacxx/cuda/ptx/generated/getctarank.inc → ...ibcudacxx/cuda/ptx/generated/getctarank.h b/...cudacxx/cuda/ptx/generated/getctarank.inc → ...ibcudacxx/cuda/ptx/generated/getctarank.h
diff --git a/...xx/cuda/ptx/generated/mbarrier_arrive.inc → ...acxx/cuda/ptx/generated/mbarrier_arrive.h b/...xx/cuda/ptx/generated/mbarrier_arrive.inc → ...acxx/cuda/ptx/generated/mbarrier_arrive.h
diff --git a/...x/generated/mbarrier_arrive_expect_tx.inc → ...ptx/generated/mbarrier_arrive_expect_tx.h b/...x/generated/mbarrier_arrive_expect_tx.inc → ...ptx/generated/mbarrier_arrive_expect_tx.h
diff --git a/...generated/mbarrier_arrive_no_complete.inc → ...x/generated/mbarrier_arrive_no_complete.h b/...generated/mbarrier_arrive_no_complete.inc → ...x/generated/mbarrier_arrive_no_complete.h
diff --git a/...acxx/cuda/ptx/generated/mbarrier_init.inc → ...udacxx/cuda/ptx/generated/mbarrier_init.h b/...acxx/cuda/ptx/generated/mbarrier_init.inc → ...udacxx/cuda/ptx/generated/mbarrier_init.h
diff --git a/.../cuda/ptx/generated/mbarrier_try_wait.inc → ...xx/cuda/ptx/generated/mbarrier_try_wait.h b/.../cuda/ptx/generated/mbarrier_try_wait.inc → ...xx/cuda/ptx/generated/mbarrier_try_wait.h
diff --git a/...tx/generated/mbarrier_try_wait_parity.inc → .../ptx/generated/mbarrier_try_wait_parity.h b/...tx/generated/mbarrier_try_wait_parity.inc → .../ptx/generated/mbarrier_try_wait_parity.h
diff --git a/...acxx/cuda/ptx/generated/mbarrier_wait.inc → ...udacxx/cuda/ptx/generated/mbarrier_wait.h b/...acxx/cuda/ptx/generated/mbarrier_wait.inc → ...udacxx/cuda/ptx/generated/mbarrier_wait.h
diff --git a/...da/ptx/generated/mbarrier_wait_parity.inc → ...cuda/ptx/generated/mbarrier_wait_parity.h b/...da/ptx/generated/mbarrier_wait_parity.inc → ...cuda/ptx/generated/mbarrier_wait_parity.h
diff --git a/...bcudacxx/cuda/ptx/generated/red_async.inc → ...libcudacxx/cuda/ptx/generated/red_async.h b/...bcudacxx/cuda/ptx/generated/red_async.inc → ...libcudacxx/cuda/ptx/generated/red_async.h
diff --git a/...ibcudacxx/cuda/ptx/generated/st_async.inc → .../libcudacxx/cuda/ptx/generated/st_async.h b/...ibcudacxx/cuda/ptx/generated/st_async.inc → .../libcudacxx/cuda/ptx/generated/st_async.h
diff --git a/...ptx/generated/tensormap_cp_fenceproxy.inc → ...a/ptx/generated/tensormap_cp_fenceproxy.h b/...ptx/generated/tensormap_cp_fenceproxy.inc → ...a/ptx/generated/tensormap_cp_fenceproxy.h
diff --git a/.../cuda/ptx/generated/tensormap_replace.inc → ...xx/cuda/ptx/generated/tensormap_replace.h b/.../cuda/ptx/generated/tensormap_replace.inc → ...xx/cuda/ptx/generated/tensormap_replace.h
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.barrier.cluster.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.barrier.cluster.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/barrier_cluster.inc"
+#include "generated/barrier_cluster.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.commit_group.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.commit_group.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/cp_async_bulk_commit_group.inc"
+#include "generated/cp_async_bulk_commit_group.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/cp_async_bulk.inc"
+#include "generated/cp_async_bulk.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.multicast.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.multicast.compile.pass.cpp
@@ -33,7 +33,7 @@
  *
  */
 
-#include "generated/cp_async_bulk_multicast.inc"
+#include "generated/cp_async_bulk_multicast.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.tensor.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.tensor.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/cp_async_bulk_tensor.inc"
+#include "generated/cp_async_bulk_tensor.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.tensor.multicast.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.tensor.multicast.compile.pass.cpp
@@ -33,7 +33,7 @@
  *
  */
 
-#include "generated/cp_async_bulk_tensor_multicast.inc"
+#include "generated/cp_async_bulk_tensor_multicast.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.wait_group.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.async.bulk.wait_group.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/cp_async_bulk_wait_group.inc"
+#include "generated/cp_async_bulk_wait_group.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.reduce.async.bulk.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.reduce.async.bulk.compile.pass.cpp
@@ -31,14 +31,14 @@
  *
  */
 
-#include "generated/cp_reduce_async_bulk.inc"
+#include "generated/cp_reduce_async_bulk.h"
 
 #ifdef _LIBCUDACXX_HAS_NVF16
-#  include "generated/cp_reduce_async_bulk_f16.inc"
+#  include "generated/cp_reduce_async_bulk_f16.h"
 #endif // _LIBCUDACXX_HAS_NVF16
 
 #ifdef _LIBCUDACXX_HAS_NVBF16
-#  include "generated/cp_reduce_async_bulk_bf16.inc"
+#  include "generated/cp_reduce_async_bulk_bf16.h"
 #endif // _LIBCUDACXX_HAS_NVBF16
 
 int main(int, char**)

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.reduce.async.bulk.tensor.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.cp.reduce.async.bulk.tensor.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/cp_reduce_async_bulk_tensor.inc"
+#include "generated/cp_reduce_async_bulk_tensor.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.fence.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.fence.compile.pass.cpp
@@ -31,11 +31,11 @@
  *
  */
 
-#include "generated/fence.inc"
-#include "generated/fence_mbarrier_init.inc"
-#include "generated/fence_proxy_alias.inc"
-#include "generated/fence_proxy_async.inc"
-#include "generated/fence_proxy_tensormap_generic.inc"
+#include "generated/fence.h"
+#include "generated/fence_mbarrier_init.h"
+#include "generated/fence_proxy_alias.h"
+#include "generated/fence_proxy_async.h"
+#include "generated/fence_proxy_tensormap_generic.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.get_sreg.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.get_sreg.compile.pass.cpp
@@ -32,7 +32,7 @@
  *
  */
 
-#include "generated/get_sreg.inc"
+#include "generated/get_sreg.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.getctarank.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.getctarank.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/getctarank.inc"
+#include "generated/getctarank.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.arrive.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.arrive.compile.pass.cpp
@@ -31,9 +31,9 @@
  *
  */
 
-#include "generated/mbarrier_arrive.inc"
-#include "generated/mbarrier_arrive_expect_tx.inc"
-#include "generated/mbarrier_arrive_no_complete.inc"
+#include "generated/mbarrier_arrive.h"
+#include "generated/mbarrier_arrive_expect_tx.h"
+#include "generated/mbarrier_arrive_no_complete.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.init.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.init.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/mbarrier_init.inc"
+#include "generated/mbarrier_init.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.wait.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.mbarrier.wait.compile.pass.cpp
@@ -31,10 +31,10 @@
  *
  */
 
-#include "generated/mbarrier_try_wait.inc"
-#include "generated/mbarrier_try_wait_parity.inc"
-#include "generated/mbarrier_wait.inc"
-#include "generated/mbarrier_wait_parity.inc"
+#include "generated/mbarrier_try_wait.h"
+#include "generated/mbarrier_try_wait_parity.h"
+#include "generated/mbarrier_wait.h"
+#include "generated/mbarrier_wait_parity.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.red.async.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.red.async.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/red_async.inc"
+#include "generated/red_async.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.st.async.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.st.async.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/st_async.inc"
+#include "generated/st_async.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.tensormap.cp_fenceproxy.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.tensormap.cp_fenceproxy.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/tensormap_cp_fenceproxy.inc"
+#include "generated/tensormap_cp_fenceproxy.h"
 
 int main(int, char**)
 {

diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.tensormap.replace.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.tensormap.replace.compile.pass.cpp
@@ -31,7 +31,7 @@
  *
  */
 
-#include "generated/tensormap_replace.inc"
+#include "generated/tensormap_replace.h"
 
 int main(int, char**)
 {
-Original file line number
+Diff line change
@@ Expand Up / @@ -31,7 +31,7 @@ @@
      *
      */
-    #include "generated/barrier_cluster.inc"
+    #include "generated/barrier_cluster.h"
     int main(int, char**)
     {
@@ Expand Down @@