asc_stcg 함수 개요

Ascend AI 프로세서를 위한 CANN (Compute Architecture for Neural Networks) 개발 킷에서는 고성능 커스텀 오퍼레이터 개발을 위해 특화된 언어 확장과 라이브러리를 제공합니다. 이 중 asc_stcg는 전역 메모리(Global Memory)에 데이터를 기록하면서 L2 캐시에만 캐싱하고, 데이터 캐시(Data Cache)에는 저장하지 않는 특수한 스토어 연산을 수행하는 인라인 함수입니다. 이를 통해 특정 메모리 계층 제어가 필요한 고도화된 최적화 시나리오에서 유용하게 사용됩니다.

지원 제품

제품군	지원 여부
Ascend 950PR / Ascend 950DT	지원됨
Atlas A3 훈련/추론 제품군	미지원
Atlas A2 훈련/추론 제품군	미지원
Atlas 200I/500 A2 추론 제품	미지원
Atlas 추론 제품 (AI Core)	미지원
Atlas 추론 제품 (Vector Core)	미지원
Atlas 훈련 제품군	미지원

기능 설명

asc_stcg는 주어진 값 val을 지정된 전역 메모리 주소 address에 기록하며, 해당 데이터는 L2 캐시에 캐싱되지만 L1 데이터 캐시에는 적재되지 않습니다. 이는 반복 액세스가 예상되지 않거나 캐시 오염을 방지해야 하는 대량 스트리밍 쓰기 작업에 적합합니다.

함수 시그니처

다양한 기본 및 벡터 타입에 대해 오버로드된 형태로 제공됩니다:

inline void asc_stcg(long int* address, long int val);
inline void asc_stcg(unsigned long int* address, unsigned long int val);
inline void asc_stcg(long long int* address, long long int val);
inline void asc_stcg(unsigned long long int* address, unsigned long long int val);

// 2/4 요소 벡터 타입 지원
inline void asc_stcg(long2* address, long2 val);
inline void asc_stcg(ulong2* address, ulong2 val);
inline void asc_stcg(long4* address, long4 val);
inline void asc_stcg(ulong4* address, ulong4 val);

// 실수형 및 반정밀도 타입
inline void asc_stcg(float* address, float val);
inline void asc_stcg(float2* address, float2 val);
inline void asc_stcg(float4* address, float4 val);

// Half-precision (FP16)
inline void asc_stcg(half* address, half val);
inline void asc_stcg(half2* address, half2 val);

// BFloat16 지원
inline void asc_stcg(bfloat16_t* address, bfloat16_t val);
inline void asc_stcg(bfloat16x2_t* address, bfloat16x2_t val);

// 정수형 벡터
inline void asc_stcg(char2* address, char2 val);
inline void asc_stcg(short4* address, short4 val);
inline void asc_stcg(int2* address, int2 val);
// 기타 유사 형식 생략
    

매개변수 설명

매개변수	방향	설명
address	입력	전역 메모리 내 대상 주소 포인터
val	입력	저장할 소스 값 또는 벡터

반환값

해당 함수는 반환값이 없습니다 (void).

헤더 포함 요구사항

사용하는 데이터 타입에 따라 다음과 같은 헤더 파일을 포함해야 합니다:

• 일반 정수/실수 타입: #include "simt_api/device_functions.h"
• FP16 타입(half, half2): #include "simt_api/asc_fp16.h"
• BFloat16 타입(bfloat16_t, bfloat16x2_t): #include "simt_api/asc_bf16.h"

사용 예제

SIMT 커널 예제:

__global__ __launch_bounds__(1024)
void kernel_stcg_write(float* input, float* output) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    asc_stcg(&output[tid], input[tid] * 2.0f);
}
    

SIMD와 SIMT 혼합 모델 예제:
글로벌 메모리 접근 시 명시적으로 __gm__ 주소 공간 한정자를 사용해야 합니다.

__simt_vf__ __launch_bounds__(1024)
inline void vector_kernel_stcg(__gm__ float* src, __gm__ float* dst) {
    int idx = get_thread_id();
    float temp = src[idx] + 1.0f;
    asc_stcg(&dst[idx], temp);
}