[all-commits] [llvm/llvm-project] 369adb: [NVPTX] 64-bit atom.{and, or, xor, min, max} require s...

[Andrew Savonichev via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 369adba0435e22722c6291142b2ce4265ee36ca3
      https://github.com/llvm/llvm-project/commit/369adba0435e22722c6291142b2ce4265ee36ca3
  Author: Andrew Savonichev <andrew.savonichev at gmail.com>
  Date:   2022-04-14 (Thu, 14 Apr 2022)

  Changed paths:
    M llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
    M llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
    M llvm/test/CodeGen/NVPTX/atomics.ll

  Log Message:
  -----------
  [NVPTX] 64-bit atom.{and,or,xor,min,max} require sm_32 or higher

PTX ISA spec, s9.7.12.4. Parallel Synchronization and Communication
Instructions: atom

Target ISA Notes
64-bit atom.{and,or,xor,min,max} require sm_32 or higher.

Differential Revision: https://reviews.llvm.org/D123038


  Commit: 230f32696497bf788ac7f4365aecabb26f6670f1
      https://github.com/llvm/llvm-project/commit/230f32696497bf788ac7f4365aecabb26f6670f1
  Author: Andrew Savonichev <andrew.savonichev at gmail.com>
  Date:   2022-04-14 (Thu, 14 Apr 2022)

  Changed paths:
    M llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
    M llvm/test/CodeGen/NVPTX/shfl-sync-p.ll
    M llvm/test/CodeGen/NVPTX/shfl-sync.ll

  Log Message:
  -----------
  [NVPTX] shfl.sync is introduced in PTX 6.0

PTX ISA spec, s9.7.8.6. Data Movement and Conversion Instructions:
shfl.sync

PTX ISA Notes
Introduced in PTX ISA version 6.0.

Target ISA Notes
Requires sm_30 or higher.

Differential Revision: https://reviews.llvm.org/D123039


  Commit: 4cef5c397d5fae8256318e8c74a2653f5c54eeb7
      https://github.com/llvm/llvm-project/commit/4cef5c397d5fae8256318e8c74a2653f5c54eeb7
  Author: Andrew Savonichev <andrew.savonichev at gmail.com>
  Date:   2022-04-14 (Thu, 14 Apr 2022)

  Changed paths:
    M llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
    M llvm/lib/Target/NVPTX/NVPTXAsmPrinter.h
    M llvm/test/CodeGen/NVPTX/managed.ll

  Log Message:
  -----------
  [NVPTX] .attribute(.managed) is only supported for sm_30 and PTX 4.0

PTX ISA spec, s5.4.8. Variable Attribute Directive: .attribute

PTX ISA Notes
Introduced in PTX ISA version 4.0.

Target ISA Notes
.managed attribute requires sm_30 or higher.

Differential Revision: https://reviews.llvm.org/D123040


  Commit: 32949401a86aea1cfc0a04bef566c2af30c82e39
      https://github.com/llvm/llvm-project/commit/32949401a86aea1cfc0a04bef566c2af30c82e39
  Author: Andrew Savonichev <andrew.savonichev at gmail.com>
  Date:   2022-04-14 (Thu, 14 Apr 2022)

  Changed paths:
    M llvm/test/CodeGen/NVPTX/b52037.ll
    M llvm/test/CodeGen/NVPTX/barrier.ll
    M llvm/test/CodeGen/NVPTX/ld-st-addrrspace.py
    M llvm/test/CodeGen/NVPTX/match.ll
    M llvm/test/CodeGen/NVPTX/shfl-p.ll
    M llvm/test/CodeGen/NVPTX/shfl-sync-p.ll
    M llvm/test/CodeGen/NVPTX/shfl-sync.ll
    M llvm/test/CodeGen/NVPTX/shfl.ll
    M llvm/test/CodeGen/NVPTX/vote.ll

  Log Message:
  -----------
  [NVPTX] Avoid dots in global names

It seems that ptxas cannot parse them:
ptxas fatal: Parsing error near '.2': syntax error

Differential Revision: https://reviews.llvm.org/D123041


  Commit: b6183a57a10b03bdad83e4bef02990673c155011
      https://github.com/llvm/llvm-project/commit/b6183a57a10b03bdad83e4bef02990673c155011
  Author: Andrew Savonichev <andrew.savonichev at gmail.com>
  Date:   2022-04-14 (Thu, 14 Apr 2022)

  Changed paths:
    M llvm/test/CodeGen/NVPTX/barrier.ll

  Log Message:
  -----------
  [NVPTX] Fix barrier.ll LIT test

The second parameter should be a multiple of the warp size (32).

PTX ISA spec, s9.7.12.1. Parallel Synchronization and Communication
Instructions: bar, barrier

barrier.sync{.aligned}      a{, b};

Operand b specifies the number of threads participating in the
barrier. If no thread count is specified, all threads in the CTA
participate in the barrier. When specifying a thread count, the value
must be a multiple of the warp size.

Differential Revision: https://reviews.llvm.org/D123470


Compare: https://github.com/llvm/llvm-project/compare/5bf9aa38abc6...b6183a57a10b