Back to Cutlass

CUTLASS: semaphore.h Source File

docs/semaphore_8h_source.html

4.4.27.9 KB
Original Source

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

semaphore.h

Go to the documentation of this file.

1 /***************************************************************************************************

2 * Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.

3 *

4 * Redistribution and use in source and binary forms, with or without modification, are permitted

5 * provided that the following conditions are met:

6 * * Redistributions of source code must retain the above copyright notice, this list of

7 * conditions and the following disclaimer.

8 * * Redistributions in binary form must reproduce the above copyright notice, this list of

9 * conditions and the following disclaimer in the documentation and/or other materials

10 * provided with the distribution.

11 * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used

12 * to endorse or promote products derived from this software without specific prior written

13 * permission.

14 *

15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR

16 * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND

17 * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE

18 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,

19 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;

20 * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,

21 * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE

22 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

23 *

24 **************************************************************************************************/

29 #pragma once

30

31 #include "cutlass/cutlass.h"

32

33 #include "cutlass/aligned_buffer.h"

34 #include "cutlass/array.h"

35

36 #include "cutlass/numeric_types.h"

37 #include "cutlass/matrix_shape.h"

38

39 #include "cutlass/gemm/gemm.h"

40

42

43 namespace cutlass {

44

46

48 class Semaphore {

49 public:

50

51int *lock;

52bool wait_thread;

53int state;

54

55 public:

56

58CUTLASS_HOST_DEVICE

59Semaphore(int *lock_, int thread_id):

60 lock(lock_),

61 wait_thread(thread_id < 0 || thread_id == 0),

62 state(-1) {

63

64 }

65

67 CUTLASS_DEVICE

68void fetch() {

69

70asm volatile ("ld.global.cg.s32 %0, [%1];\n" : "=r"(state) : "l"(lock));

71 }

72

74 CUTLASS_DEVICE

75int get_state() const {

76return state;

77 }

78

80 CUTLASS_DEVICE

81void wait(int status = 0) {

82

83if (wait_thread) {

84while (state != status) {

85

86fetch();

87

88 __syncwarp(0x01);

89

90 };

91 }

92

93 __syncthreads();

94 }

95

97 CUTLASS_DEVICE

98void release(int status = 0) {

99 __syncthreads();

100

101if (wait_thread) {

102

103asm volatile ("st.global.cg.s32 [%0], %1;\n" : : "l"(lock), "r"(status));

104 }

105 }

106 };

107

109

110 } // namespace cutlass

111

cutlass

Definition: aligned_buffer.h:35

gemm.h

Defines common types used for all GEMM-like operators.

cutlass::Semaphore::Semaphore

CUTLASS_HOST_DEVICE Semaphore(int *lock_, int thread_id)

Implements a semaphore to wait for a flag to reach a given value.

Definition: semaphore.h:59

cutlass::Semaphore::fetch

CUTLASS_DEVICE void fetch()

Permit fetching the synchronization mechanism early.

Definition: semaphore.h:68

array.h

Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...

matrix_shape.h

Defines a Shape template for matrix tiles.

cutlass::Semaphore::get_state

CUTLASS_DEVICE int get_state() const

Gets the internal state.

Definition: semaphore.h:75

cutlass::Semaphore::wait_thread

bool wait_thread

Definition: semaphore.h:52

aligned_buffer.h

AlignedBuffer is a container for trivially copyable elements suitable for use in unions and shared me...

CUTLASS_HOST_DEVICE

#define CUTLASS_HOST_DEVICE

Definition: cutlass.h:89

numeric_types.h

Top-level include for all CUTLASS numeric types.

cutlass::Semaphore

CTA-wide semaphore for inter-CTA synchronization.

Definition: semaphore.h:48

cutlass::Semaphore::release

CUTLASS_DEVICE void release(int status=0)

Updates the lock with the given result.

Definition: semaphore.h:98

cutlass::Semaphore::wait

CUTLASS_DEVICE void wait(int status=0)

Waits until the semaphore is equal to the given value.

Definition: semaphore.h:81

cutlass.h

Basic include for CUTLASS.

cutlass::Semaphore::state

int state

Definition: semaphore.h:53

cutlass::Semaphore::lock

int * lock

Definition: semaphore.h:51


Generated by 1.8.11