submission 68267

  #!POPCORN leaderboard vectoradd_v2
+ #!POPCORN gpus A100 H100 B200 L4
   
+ try:
+     import cuda.parallel.experimental.algorithms as algorithms
+ except:
+     import os
+     import subprocess
+  
+     if not os.path.exists("cccl"):
+         subprocess.check_call(
+             ["git", "clone", "https://github.com/NaderAlAwar/cccl.git"]
+         )
+  
+     subprocess.check_call(["git", "checkout", "gpu-mode-submissions-a100"], cwd="cccl")
+     subprocess.check_call(
+         ["git", "checkout", "49b7297dbe3abddbc25f937b132b8e6e16202100"], cwd="cccl"
+     )
+  
+     env = os.environ.copy()
+     env["CC"] = "gcc"
+     env["CXX"] = "g++"
+     env["CMAKE_ARGS"] = "-DCMAKE_CXX_STANDARD=20"
+  
+     subprocess.check_call(
+         ["pip", "install", "../cuda_cccl"], cwd="cccl/python/cuda_parallel", env=env
+     )
+     subprocess.check_call(
+         ["pip", "install", "."], cwd="cccl/python/cuda_parallel", env=env
+     )
+  
+ import functools
+ import cuda.parallel.experimental.algorithms as algorithms
  import torch
- import triton
- import triton.language as tl
+  
  from task import input_t, output_t
   
- # Autotune configs: same as before (wide sweep)
- AUTOTUNE_CONFIGS = [
-     triton.Config({"BLOCK_SIZE": 128, "NUM_WARPS": 1, "NUM_STAGES": 2}),
-     triton.Config({"BLOCK_SIZE": 256, "NUM_WARPS": 1, "NUM_STAGES": 2}),
-     triton.Config({"BLOCK_SIZE": 512, "NUM_WARPS": 2, "NUM_STAGES": 2}),
-     triton.Config({"BLOCK_SIZE": 1024, "NUM_WARPS": 4, "NUM_STAGES": 2}),
-     triton.Config({"BLOCK_SIZE": 2048, "NUM_WARPS": 4, "NUM_STAGES": 3}),
-     triton.Config({"BLOCK_SIZE": 4096, "NUM_WARPS": 8, "NUM_STAGES": 3}),
-     triton.Config({"BLOCK_SIZE": 8192, "NUM_WARPS": 8, "NUM_STAGES": 4}),
-     triton.Config({"BLOCK_SIZE": 16384, "NUM_WARPS": 8, "NUM_STAGES": 4}),
- ]
+ d_in1 = torch.tensor([1], dtype=torch.float16).cuda()
+ d_in2 = torch.tensor([1], dtype=torch.float16).cuda()
   
   
- @triton.autotune(configs=AUTOTUNE_CONFIGS, key=["size"])
- @triton.jit
- def add_kernel(
-     A_ptr,
-     B_ptr,
-     C_ptr,
-     size,
-     BLOCK_SIZE: tl.constexpr,
-     NUM_WARPS: tl.constexpr,
-     NUM_STAGES: tl.constexpr,
- ):
-     pid = tl.program_id(0)
-     offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
-     mask = offsets < size
-     a = tl.load(A_ptr + offsets, mask=mask)
-     b = tl.load(B_ptr + offsets, mask=mask)
-     tl.store(C_ptr + offsets, a + b, mask=mask)
+ def op(a, b):
+     return a + b
   
   
- def custom_kernel(data: input_t, deterministic: bool = False) -> output_t:
-     """
-     Wrapper:
-       - ensures contiguity,
-       - chooses BLOCK_SIZE from the autotune candidates (so grid matches compile-time BLOCK_SIZE),
-       - constructs grid with that BLOCK_SIZE,
-       - launches autotuned kernel.
-     If deterministic=True, use a fixed config (no autotune variability) for consistent benchmarking.
-     """
-     A, B, C = data
-     A = A.contiguous()
-     B = B.contiguous()
-     C = C.contiguous()
-     assert A.numel() == B.numel() == C.numel(), "inputs must have same numel"
+ transform = algorithms.binary_transform(d_in1, d_in2, d_in2, op)
   
-     size = A.numel()
   
-     # Grid computed using the same BLOCK_SIZE we pass to the kernel — keeps coverage exact.
-     grid = lambda meta: (triton.cdiv(size, meta["BLOCK_SIZE"]),)
+ @functools.cache
+ def initialize(shape):
+     return torch.empty(shape, dtype=torch.float16).cuda(), shape[0] * shape[1]
   
-     # Prefer letting the autotuner pick NUM_WARPS/NUM_STAGES for this BLOCK_SIZE.
-     # By passing BLOCK_SIZE here we ensure the grid is consistent with the compile-time BLOCK_SIZE.
-     add_kernel[grid](A, B, C, size)
   
-     return C
+ def custom_kernel(data: input_t) -> output_t:
+     d_in1, d_in2, _ = data
+     d_out, numel = initialize(d_in1.shape)
+     transform(d_in1, d_in2, d_out, numel)
+     return d_out