[LDC 1.42.0] invalidPtx error occurs in CUDA

[takinutani]

At LDC 1.42.0 (1.42.0-beta2 or later), dub test cause invalidPtx error.

PS D:\labo\dlang\dcompute_master> dub test

dcompute.driver.error.DComputeDriverException@source\dcompute\driver\error.d(207): invalidPtx
----------------
0x00007FF696623F31 in dcompute.driver.error.__lambda_L205_C5 at D:\labo\dlang\dcompute_master\source\dcompute\driver\error.d(207)
0x00007FF696623EA0 in dcompute.driver.error.__lambda_L200_C43 at D:\labo\dlang\dcompute_master\source\dcompute\driver\error.d(203)
0x00007FF696623D85 in dcompute.driver.error.checkErrors at D:\labo\dlang\dcompute_master\source\dcompute\driver\error.d(223)
0x00007FF696621BD1 in dcompute.driver.cuda.program.Program.fromFile at D:\labo\dlang\dcompute_master\source\dcompute\driver\cuda\program.d(31)
0x00007FF6966212FD in main.D main at D:\labo\dlang\dcompute_master\source\dcompute\tests\main.d(121)
0x00007FF696680691 in void rt.dmain2._d_run_main2(char[][], ulong, extern (C) int function(char[][])*).runAll()
0x00007FF696680320 in d_run_main2
0x00007FF6966805F3 in d_wrun_main
0x00007FF6966224B4 in main.wmain at D:\scoop\user\apps\ldc\1.42.0\import\core\internal\entrypoint.d(32)
0x00007FF6966A5F0C in __scrt_common_main_seh at d:\agent\_work\2\s\src\vctools\crt\vcstartup\src\startup\exe_common.inl(288)
0x00007FFBF89AE8D7 in BaseThreadInitThunk
0x00007FFBF9CEC48C in RtlUserThreadStart
Error Program exited with code 1

The direct reason is that the following function is not converted to NVPTX IR.

• source\dcompute\std\cuda\sync.d

pragma(LDC_intrinsic, "llvm.nvvm.barrier0")
void barrier0();

becomes bellow, leaving LLVM IR in PTX.

.extern .func llvm.nvvm.barrier0
()
;

.visible .func _D8dcompute3std4sync7barrierFZv()
{


	{ // callseq 6, 0
	call.uni llvm.nvvm.barrier0, ();
	} // callseq 6
	ret;

}

[takinutani]

When checking the built-in functions via gccbuiltins_nvvm.di, XX may have been dropped.

$ diff -c ldc/1.41.0/import/ldc/gccbuiltins_nvvm.di ldc/1.42.0/import/ldc/gccbuiltins_nvvm.di
*** ldc/1.41.0/import/ldc/gccbuiltins_nvvm.di   Sat Jun 07 11:54:37 2025
--- ldc/1.42.0/import/ldc/gccbuiltins_nvvm.di   Mon Mar 02 03:55:46 2026
***************
*** 43,60 ****
  pragma(LDC_intrinsic, "llvm.nvvm.add.rz.ftz.f")
      float __nvvm_add_rz_ftz_f(float, float) pure @safe;

- pragma(LDC_intrinsic, "llvm.nvvm.bar.sync")
-     void __nvvm_bar_sync(int);
-
  pragma(LDC_intrinsic, "llvm.nvvm.bar.warp.sync")
      void __nvvm_bar_warp_sync(int);

- pragma(LDC_intrinsic, "llvm.nvvm.barrier")
-     void __nvvm_bar(int, int);
-
- pragma(LDC_intrinsic, "llvm.nvvm.barrier0")
-     void __syncthreads();
-
  pragma(LDC_intrinsic, "llvm.nvvm.barrier0.and")
      int __nvvm_bar0_and(int);

--- 43,51 ----
***************
*** 64,78 ****
  pragma(LDC_intrinsic, "llvm.nvvm.barrier0.popc")
      int __nvvm_bar0_popc(int);

- pragma(LDC_intrinsic, "llvm.nvvm.barrier.n")
-     void __nvvm_bar_n(int);
-
- pragma(LDC_intrinsic, "llvm.nvvm.barrier.sync")
-     void __nvvm_barrier_sync(int);
-
- pragma(LDC_intrinsic, "llvm.nvvm.barrier.sync.cnt")
-     void __nvvm_barrier_sync_cnt(int, int);
-
  pragma(LDC_intrinsic, "llvm.nvvm.ceil.d")
      double __nvvm_ceil_d(double) pure @safe;

[takinutani]

At [llvm-project] llvm/IR/IntrinsicsNVVM.td

// The following intrinsics were once defined here, but are now auto-upgraded
// to target-generic LLVM intrinsics.

//   * llvm.nvvm.barrier0            --> llvm.nvvm.barrier.cta.sync.aligned.all(0)
//   * llvm.nvvm.barrier.n           --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
//   * llvm.nvvm.bar.sync            --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
//   * llvm.nvvm.barrier             --> llvm.nvvm.barrier.cta.sync.aligned(x, y)
//   * llvm.nvvm.barrier.sync        --> llvm.nvvm.barrier.cta.sync.all(x)
//   * llvm.nvvm.barrier.sync.cnt    --> llvm.nvvm.barrier.cta.sync(x, y)
//   * llvm.nvvm.barrier0.popc       --> llvm.nvvm.barrier.cta.red.popc.aligned.all(0, c)
//   * llvm.nvvm.barrier0.and        --> llvm.nvvm.barrier.cta.red.and.aligned.all(0, z)
//   * llvm.nvvm.barrier0.or         --> llvm.nvvm.barrier.cta.red.or.aligned.all(0, z)

It seems that llvm.nvvm.barrier0 is outdated, and LLVM automatically converts it to llvm.nvvm.barrier.cta.sync.aligned.all, so idea is that we can use llvm.nvvm.barrier.cta.sync.aligned.all instead of llvm.nvvm.barrier0.