Mercurial > hg > CbC > CbC_gcc

comparison gcc/config/nvptx/nvptx.md @ 131:84e7813d76e9

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
gcc-8.2
author mir3636
date Thu, 25 Oct 2018 07:37:49 +0900
parents 04ced10e8804
children 1830386684a0
comparison
equal deleted inserted replaced
111:04ced10e8804 131:84e7813d76e9
1 ;; Machine description for NVPTX. 1 ;; Machine description for NVPTX.
2 ;; Copyright (C) 2014-2017 Free Software Foundation, Inc. 2 ;; Copyright (C) 2014-2018 Free Software Foundation, Inc.
3 ;; Contributed by Bernd Schmidt <bernds@codesourcery.com> 3 ;; Contributed by Bernd Schmidt <bernds@codesourcery.com>
4 ;; 4 ;;
5 ;; This file is part of GCC. 5 ;; This file is part of GCC.
6 ;; 6 ;;
7 ;; GCC is free software; you can redistribute it and/or modify 7 ;; GCC is free software; you can redistribute it and/or modify
53 (define_c_enum "unspecv" [ 53 (define_c_enum "unspecv" [
54 UNSPECV_LOCK 54 UNSPECV_LOCK
55 UNSPECV_CAS 55 UNSPECV_CAS
56 UNSPECV_XCHG 56 UNSPECV_XCHG
57 UNSPECV_BARSYNC 57 UNSPECV_BARSYNC
58 UNSPECV_MEMBAR
59 UNSPECV_MEMBAR_CTA
58 UNSPECV_DIM_POS 60 UNSPECV_DIM_POS
59 61
60 UNSPECV_FORK 62 UNSPECV_FORK
61 UNSPECV_FORKED 63 UNSPECV_FORKED
62 UNSPECV_JOINING 64 UNSPECV_JOINING
992 (define_insn "nop" 994 (define_insn "nop"
993 [(const_int 0)] 995 [(const_int 0)]
994 "" 996 ""
995 "") 997 "")
996 998
999 (define_insn "exit"
1000 [(const_int 1)]
1001 ""
1002 "exit;")
1003
1004 (define_insn "fake_nop"
1005 [(const_int 2)]
1006 ""
1007 "{
1008 .reg .u32 %%nop_src;
1009 .reg .u32 %%nop_dst;
1010 mov.u32 %%nop_dst, %%nop_src;
1011 }")
1012
997 (define_insn "return" 1013 (define_insn "return"
998 [(return)] 1014 [(return)]
999 "" 1015 ""
1000 { 1016 {
1001 return nvptx_output_return (); 1017 return nvptx_output_return ();
1083 }) 1099 })
1084 1100
1085 (define_insn "trap" 1101 (define_insn "trap"
1086 [(trap_if (const_int 1) (const_int 0))] 1102 [(trap_if (const_int 1) (const_int 0))]
1087 "" 1103 ""
1088 "trap;") 1104 "trap; exit;")
1089 1105
1090 (define_insn "trap_if_true" 1106 (define_insn "trap_if_true"
1091 [(trap_if (ne (match_operand:BI 0 "nvptx_register_operand" "R") 1107 [(trap_if (ne (match_operand:BI 0 "nvptx_register_operand" "R")
1092 (const_int 0)) 1108 (const_int 0))
1093 (const_int 0))] 1109 (const_int 0))]
1094 "" 1110 ""
1095 "%j0 trap;" 1111 "%j0 trap; %j0 exit;"
1096 [(set_attr "predicable" "false")]) 1112 [(set_attr "predicable" "false")])
1097 1113
1098 (define_insn "trap_if_false" 1114 (define_insn "trap_if_false"
1099 [(trap_if (eq (match_operand:BI 0 "nvptx_register_operand" "R") 1115 [(trap_if (eq (match_operand:BI 0 "nvptx_register_operand" "R")
1100 (const_int 0)) 1116 (const_int 0))
1101 (const_int 0))] 1117 (const_int 0))]
1102 "" 1118 ""
1103 "%J0 trap;" 1119 "%J0 trap; %J0 exit;"
1104 [(set_attr "predicable" "false")]) 1120 [(set_attr "predicable" "false")])
1105 1121
1106 (define_expand "ctrap<mode>4" 1122 (define_expand "ctrap<mode>4"
1107 [(trap_if (match_operator 0 "nvptx_comparison_operator" 1123 [(trap_if (match_operator 0 "nvptx_comparison_operator"
1108 [(match_operand:SDIM 1 "nvptx_register_operand") 1124 [(match_operand:SDIM 1 "nvptx_register_operand")
1422 [(set_attr "atomic" "true")]) 1438 [(set_attr "atomic" "true")])
1423 1439
1424 (define_code_iterator any_logic [and ior xor]) 1440 (define_code_iterator any_logic [and ior xor])
1425 (define_code_attr logic [(and "and") (ior "or") (xor "xor")]) 1441 (define_code_attr logic [(and "and") (ior "or") (xor "xor")])
1426 1442
1427 ;; Currently disabled until we add better subtarget support - requires sm_32.
1428 (define_insn "atomic_fetch_<logic><mode>" 1443 (define_insn "atomic_fetch_<logic><mode>"
1429 [(set (match_operand:SDIM 1 "memory_operand" "+m") 1444 [(set (match_operand:SDIM 1 "memory_operand" "+m")
1430 (unspec_volatile:SDIM 1445 (unspec_volatile:SDIM
1431 [(any_logic:SDIM (match_dup 1) 1446 [(any_logic:SDIM (match_dup 1)
1432 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri")) 1447 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))
1433 (match_operand:SI 3 "const_int_operand")] ;; model 1448 (match_operand:SI 3 "const_int_operand")] ;; model
1434 UNSPECV_LOCK)) 1449 UNSPECV_LOCK))
1435 (set (match_operand:SDIM 0 "nvptx_register_operand" "=R") 1450 (set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
1436 (match_dup 1))] 1451 (match_dup 1))]
1437 "0" 1452 "<MODE>mode == SImode || TARGET_SM35"
1438 "%.\\tatom%A1.b%T0.<logic>\\t%0, %1, %2;" 1453 "%.\\tatom%A1.b%T0.<logic>\\t%0, %1, %2;"
1439 [(set_attr "atomic" "true")]) 1454 [(set_attr "atomic" "true")])
1440 1455
1441 (define_insn "nvptx_barsync" 1456 (define_insn "nvptx_barsync"
1442 [(unspec_volatile [(match_operand:SI 0 "const_int_operand" "")] 1457 [(unspec_volatile [(match_operand:SI 0 "const_int_operand" "")]
1443 UNSPECV_BARSYNC)] 1458 UNSPECV_BARSYNC)]
1444 "" 1459 ""
1445 "\\tbar.sync\\t%0;" 1460 "\\tbar.sync\\t%0;"
1446 [(set_attr "predicable" "false")]) 1461 [(set_attr "predicable" "false")])
1447 1462
1463 (define_expand "memory_barrier"
1464 [(set (match_dup 0)
1465 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR))]
1466 ""
1467 {
1468 operands[0] = gen_rtx_MEM (BLKmode, gen_rtx_SCRATCH (Pmode));
1469 MEM_VOLATILE_P (operands[0]) = 1;
1470 })
1471
1472 ;; Ptx defines the memory barriers membar.cta, membar.gl and membar.sys
1473 ;; (corresponding to cuda functions threadfence_block, threadfence and
1474 ;; threadfence_system). For the insn memory_barrier we use membar.sys. This
1475 ;; may be overconservative, but before using membar.gl instead we'll need to
1476 ;; explain in detail why it's safe to use. For now, use membar.sys.
1477 (define_insn "*memory_barrier"
1478 [(set (match_operand:BLK 0 "" "")
1479 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR))]
1480 ""
1481 "\\tmembar.sys;"
1482 [(set_attr "predicable" "false")])
1483
1484 (define_expand "nvptx_membar_cta"
1485 [(set (match_dup 0)
1486 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR_CTA))]
1487 ""
1488 {
1489 operands[0] = gen_rtx_MEM (BLKmode, gen_rtx_SCRATCH (Pmode));
1490 MEM_VOLATILE_P (operands[0]) = 1;
1491 })
1492
1493 (define_insn "*nvptx_membar_cta"
1494 [(set (match_operand:BLK 0 "" "")
1495 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR_CTA))]
1496 ""
1497 "\\tmembar.cta;"
1498 [(set_attr "predicable" "false")])
1499
1448 (define_insn "nvptx_nounroll" 1500 (define_insn "nvptx_nounroll"
1449 [(unspec_volatile [(const_int 0)] UNSPECV_NOUNROLL)] 1501 [(unspec_volatile [(const_int 0)] UNSPECV_NOUNROLL)]
1450 "" 1502 ""
1451 "\\t.pragma \\\"nounroll\\\";" 1503 "\\t.pragma \\\"nounroll\\\";"
1452 [(set_attr "predicable" "false")]) 1504 [(set_attr "predicable" "false")])