Skip to content

pycuda do not support triton.autotune when len(configs) > 1 #465

Description

@l1351868270

Describe the bug
I am using openai triton to write some operator. I want to use pycuda array instead of pytorch tensor, so i will not dependent on the pytorch(for some platform pytorch is not support). but it do not support triton.autotune when len(configs) > 1. The source code is here:
https://github.com/l1351868270/implicit_gemm.triton/blob/main/vector-add.py

To Reproduce
Steps to reproduce the behavior:

  1. git clone https://github.com/l1351868270/implicit_gemm.triton.git
  2. cd implicit_gemm.triton
  3. python vector-add.py
  4. when i use https://github.com/l1351868270/implicit_gemm.triton/blob/main/vector-add.py#L18-L21 it is ok.
  5. but when i use https://github.com/l1351868270/implicit_gemm.triton/blob/ab5211fab7ecc49665b7440443e6e9b87b3cca89/vector-add.py#L14-L17 , it will occur an error
  6. Traceback (most recent call last):
    File "/root/implicit_gemm.triton/vector-add.py", line 97, in
    output_triton = add(tl_a, tl_b, tl_c, size)
    ^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/root/implicit_gemm.triton/vector-add.py", line 80, in add
    add_kernel[grid](x, y, output, n_elements)
    File "/root/miniconda3/envs/cu121/lib/python3.11/site-packages/triton/runtime/jit.py", line 347, in
    return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)
    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/root/miniconda3/envs/cu121/lib/python3.11/site-packages/triton/runtime/autotuner.py", line 156, in run
    timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}
    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/root/miniconda3/envs/cu121/lib/python3.11/site-packages/triton/runtime/autotuner.py", line 156, in
    timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}
    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/root/miniconda3/envs/cu121/lib/python3.11/site-packages/triton/runtime/autotuner.py", line 133, in _bench
    return do_bench(kernel_call, warmup=self.num_warmups, rep=self.num_reps, quantiles=(0.5, 0.2, 0.8))
    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/root/miniconda3/envs/cu121/lib/python3.11/site-packages/triton/testing.py", line 120, in do_bench
    fn()
    File "/root/miniconda3/envs/cu121/lib/python3.11/site-packages/triton/runtime/autotuner.py", line 114, in kernel_call
    self.fn.run(
    File "/root/miniconda3/envs/cu121/lib/python3.11/site-packages/triton/runtime/jit.py", line 693, in run
    kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,
    File "/root/miniconda3/envs/cu121/lib/python3.11/site-packages/triton/backends/nvidia/driver.py", line 366, in call
    self.launch(*args, **kwargs)
    ValueError: Pointer argument (at 0) cannot be accessed from Triton (cpu tensor?)

7.the triton code is here https://github.com/triton-lang/triton/blob/185299e204ea6e163cc717da137526939002659b/third_party/nvidia/backend/driver.py#L273-L278

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions