130 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py2// REQUIRES: amdgpu-registered-target3// REQUIRES: x86-registered-target4 5// RUN: %clang_cc1 "-aux-triple" "x86_64-unknown-linux-gnu" "-triple" "spirv64-amd-amdhsa" \6// RUN: -fcuda-is-device "-aux-target-cpu" "x86-64" -emit-llvm -o - %s | FileCheck %s7 8#include "Inputs/cuda.h"9 10// CHECK-LABEL: @_Z8test_argPDF16bDF16b(11// CHECK-NEXT: entry:12// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(4), align 813// CHECK-NEXT: [[IN_ADDR:%.*]] = alloca bfloat, align 214// CHECK-NEXT: [[BF16:%.*]] = alloca bfloat, align 215// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr [[OUT_ADDR]] to ptr addrspace(4)16// CHECK-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr [[IN_ADDR]] to ptr addrspace(4)17// CHECK-NEXT: [[BF16_ASCAST:%.*]] = addrspacecast ptr [[BF16]] to ptr addrspace(4)18// CHECK-NEXT: store ptr addrspace(4) [[OUT:%.*]], ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 819// CHECK-NEXT: store bfloat [[IN:%.*]], ptr addrspace(4) [[IN_ADDR_ASCAST]], align 220// CHECK-NEXT: [[TMP0:%.*]] = load bfloat, ptr addrspace(4) [[IN_ADDR_ASCAST]], align 221// CHECK-NEXT: store bfloat [[TMP0]], ptr addrspace(4) [[BF16_ASCAST]], align 222// CHECK-NEXT: [[TMP1:%.*]] = load bfloat, ptr addrspace(4) [[BF16_ASCAST]], align 223// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 824// CHECK-NEXT: store bfloat [[TMP1]], ptr addrspace(4) [[TMP2]], align 225// CHECK-NEXT: ret void26//27__device__ void test_arg(__bf16 *out, __bf16 in) {28 __bf16 bf16 = in;29 *out = bf16;30}31 32// CHECK-LABEL: @_Z9test_loadPDF16bS_(33// CHECK-NEXT: entry:34// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(4), align 835// CHECK-NEXT: [[IN_ADDR:%.*]] = alloca ptr addrspace(4), align 836// CHECK-NEXT: [[BF16:%.*]] = alloca bfloat, align 237// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr [[OUT_ADDR]] to ptr addrspace(4)38// CHECK-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr [[IN_ADDR]] to ptr addrspace(4)39// CHECK-NEXT: [[BF16_ASCAST:%.*]] = addrspacecast ptr [[BF16]] to ptr addrspace(4)40// CHECK-NEXT: store ptr addrspace(4) [[OUT:%.*]], ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 841// CHECK-NEXT: store ptr addrspace(4) [[IN:%.*]], ptr addrspace(4) [[IN_ADDR_ASCAST]], align 842// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[IN_ADDR_ASCAST]], align 843// CHECK-NEXT: [[TMP1:%.*]] = load bfloat, ptr addrspace(4) [[TMP0]], align 244// CHECK-NEXT: store bfloat [[TMP1]], ptr addrspace(4) [[BF16_ASCAST]], align 245// CHECK-NEXT: [[TMP2:%.*]] = load bfloat, ptr addrspace(4) [[BF16_ASCAST]], align 246// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 847// CHECK-NEXT: store bfloat [[TMP2]], ptr addrspace(4) [[TMP3]], align 248// CHECK-NEXT: ret void49//50__device__ void test_load(__bf16 *out, __bf16 *in) {51 __bf16 bf16 = *in;52 *out = bf16;53}54 55// CHECK-LABEL: @_Z8test_retDF16b(56// CHECK-NEXT: entry:57// CHECK-NEXT: [[RETVAL:%.*]] = alloca bfloat, align 258// CHECK-NEXT: [[IN_ADDR:%.*]] = alloca bfloat, align 259// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr [[RETVAL]] to ptr addrspace(4)60// CHECK-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr [[IN_ADDR]] to ptr addrspace(4)61// CHECK-NEXT: store bfloat [[IN:%.*]], ptr addrspace(4) [[IN_ADDR_ASCAST]], align 262// CHECK-NEXT: [[TMP0:%.*]] = load bfloat, ptr addrspace(4) [[IN_ADDR_ASCAST]], align 263// CHECK-NEXT: ret bfloat [[TMP0]]64//65__device__ __bf16 test_ret( __bf16 in) {66 return in;67}68 69// CHECK-LABEL: @_Z9test_callDF16b(70// CHECK-NEXT: entry:71// CHECK-NEXT: [[RETVAL:%.*]] = alloca bfloat, align 272// CHECK-NEXT: [[IN_ADDR:%.*]] = alloca bfloat, align 273// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr [[RETVAL]] to ptr addrspace(4)74// CHECK-NEXT: [[IN_ADDR_ASCAST:%.*]] = addrspacecast ptr [[IN_ADDR]] to ptr addrspace(4)75// CHECK-NEXT: store bfloat [[IN:%.*]], ptr addrspace(4) [[IN_ADDR_ASCAST]], align 276// CHECK-NEXT: [[TMP0:%.*]] = load bfloat, ptr addrspace(4) [[IN_ADDR_ASCAST]], align 277// CHECK-NEXT: [[CALL:%.*]] = call contract spir_func noundef addrspace(4) bfloat @_Z8test_retDF16b(bfloat noundef [[TMP0]]) #[[ATTR1:[0-9]+]]78// CHECK-NEXT: ret bfloat [[CALL]]79//80__device__ __bf16 test_call( __bf16 in) {81 return test_ret(in);82}83 84 85// CHECK-LABEL: @_Z15test_vec_assignv(86// CHECK-NEXT: entry:87// CHECK-NEXT: [[VEC2_A:%.*]] = alloca <2 x bfloat>, align 488// CHECK-NEXT: [[VEC2_B:%.*]] = alloca <2 x bfloat>, align 489// CHECK-NEXT: [[VEC4_A:%.*]] = alloca <4 x bfloat>, align 890// CHECK-NEXT: [[VEC4_B:%.*]] = alloca <4 x bfloat>, align 891// CHECK-NEXT: [[VEC8_A:%.*]] = alloca <8 x bfloat>, align 1692// CHECK-NEXT: [[VEC8_B:%.*]] = alloca <8 x bfloat>, align 1693// CHECK-NEXT: [[VEC16_A:%.*]] = alloca <16 x bfloat>, align 3294// CHECK-NEXT: [[VEC16_B:%.*]] = alloca <16 x bfloat>, align 3295// CHECK-NEXT: [[VEC2_A_ASCAST:%.*]] = addrspacecast ptr [[VEC2_A]] to ptr addrspace(4)96// CHECK-NEXT: [[VEC2_B_ASCAST:%.*]] = addrspacecast ptr [[VEC2_B]] to ptr addrspace(4)97// CHECK-NEXT: [[VEC4_A_ASCAST:%.*]] = addrspacecast ptr [[VEC4_A]] to ptr addrspace(4)98// CHECK-NEXT: [[VEC4_B_ASCAST:%.*]] = addrspacecast ptr [[VEC4_B]] to ptr addrspace(4)99// CHECK-NEXT: [[VEC8_A_ASCAST:%.*]] = addrspacecast ptr [[VEC8_A]] to ptr addrspace(4)100// CHECK-NEXT: [[VEC8_B_ASCAST:%.*]] = addrspacecast ptr [[VEC8_B]] to ptr addrspace(4)101// CHECK-NEXT: [[VEC16_A_ASCAST:%.*]] = addrspacecast ptr [[VEC16_A]] to ptr addrspace(4)102// CHECK-NEXT: [[VEC16_B_ASCAST:%.*]] = addrspacecast ptr [[VEC16_B]] to ptr addrspace(4)103// CHECK-NEXT: [[TMP0:%.*]] = load <2 x bfloat>, ptr addrspace(4) [[VEC2_B_ASCAST]], align 4104// CHECK-NEXT: store <2 x bfloat> [[TMP0]], ptr addrspace(4) [[VEC2_A_ASCAST]], align 4105// CHECK-NEXT: [[TMP1:%.*]] = load <4 x bfloat>, ptr addrspace(4) [[VEC4_B_ASCAST]], align 8106// CHECK-NEXT: store <4 x bfloat> [[TMP1]], ptr addrspace(4) [[VEC4_A_ASCAST]], align 8107// CHECK-NEXT: [[TMP2:%.*]] = load <8 x bfloat>, ptr addrspace(4) [[VEC8_B_ASCAST]], align 16108// CHECK-NEXT: store <8 x bfloat> [[TMP2]], ptr addrspace(4) [[VEC8_A_ASCAST]], align 16109// CHECK-NEXT: [[TMP3:%.*]] = load <16 x bfloat>, ptr addrspace(4) [[VEC16_B_ASCAST]], align 32110// CHECK-NEXT: store <16 x bfloat> [[TMP3]], ptr addrspace(4) [[VEC16_A_ASCAST]], align 32111// CHECK-NEXT: ret void112//113__device__ void test_vec_assign() {114 typedef __attribute__((ext_vector_type(2))) __bf16 bf16_x2;115 bf16_x2 vec2_a, vec2_b;116 vec2_a = vec2_b;117 118 typedef __attribute__((ext_vector_type(4))) __bf16 bf16_x4;119 bf16_x4 vec4_a, vec4_b;120 vec4_a = vec4_b;121 122 typedef __attribute__((ext_vector_type(8))) __bf16 bf16_x8;123 bf16_x8 vec8_a, vec8_b;124 vec8_a = vec8_b;125 126 typedef __attribute__((ext_vector_type(16))) __bf16 bf16_x16;127 bf16_x16 vec16_a, vec16_b;128 vec16_a = vec16_b;129}130