mirror of
https://github.com/c64scene-ar/llvm-6502.git
synced 2024-08-21 18:29:46 +00:00
5a70dd1d82
Similar to gep (r230786) and load (r230794) changes.
Similar migration script can be used to update test cases, which
successfully migrated all of LLVM and Polly, but about 4 test cases
needed manually changes in Clang.
(this script will read the contents of stdin and massage it into stdout
- wrap it in the 'apply.sh' script shown in previous commits + xargs to
apply it over a large set of test cases)
import fileinput
import sys
import re
rep = re.compile(r"(getelementptr(?:\s+inbounds)?\s*\()((<\d*\s+x\s+)?([^@]*?)(|\s*addrspace\(\d+\))\s*\*(?(3)>)\s*)(?=$|%|@|null|undef|blockaddress|getelementptr|addrspacecast|bitcast|inttoptr|zeroinitializer|<|\[\[[a-zA-Z]|\{\{)", re.MULTILINE | re.DOTALL)
def conv(match):
line = match.group(1)
line += match.group(4)
line += ", "
line += match.group(2)
return line
line = sys.stdin.read()
off = 0
for match in re.finditer(rep, line):
sys.stdout.write(line[off:match.start()])
sys.stdout.write(conv(match))
off = match.end()
sys.stdout.write(line[off:])
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@232184 91177308-0d34-0410-b5e6-96231b3b80d8
92 lines
3.5 KiB
LLVM
92 lines
3.5 KiB
LLVM
; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix PTX
|
|
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix PTX
|
|
; RUN: opt < %s -S -nvptx-favor-non-generic -dce | FileCheck %s --check-prefix IR
|
|
|
|
@array = internal addrspace(3) global [10 x float] zeroinitializer, align 4
|
|
@scalar = internal addrspace(3) global float 0.000000e+00, align 4
|
|
|
|
; Verifies nvptx-favor-non-generic correctly optimizes generic address space
|
|
; usage to non-generic address space usage for the patterns we claim to handle:
|
|
; 1. load cast
|
|
; 2. store cast
|
|
; 3. load gep cast
|
|
; 4. store gep cast
|
|
; gep and cast can be an instruction or a constant expression. This function
|
|
; tries all possible combinations.
|
|
define float @ld_st_shared_f32(i32 %i, float %v) {
|
|
; IR-LABEL: @ld_st_shared_f32
|
|
; IR-NOT: addrspacecast
|
|
; PTX-LABEL: ld_st_shared_f32(
|
|
; load cast
|
|
%1 = load float, float* addrspacecast (float addrspace(3)* @scalar to float*), align 4
|
|
; PTX: ld.shared.f32 %f{{[0-9]+}}, [scalar];
|
|
; store cast
|
|
store float %v, float* addrspacecast (float addrspace(3)* @scalar to float*), align 4
|
|
; PTX: st.shared.f32 [scalar], %f{{[0-9]+}};
|
|
; use syncthreads to disable optimizations across components
|
|
call void @llvm.cuda.syncthreads()
|
|
; PTX: bar.sync 0;
|
|
|
|
; cast; load
|
|
%2 = addrspacecast float addrspace(3)* @scalar to float*
|
|
%3 = load float, float* %2, align 4
|
|
; PTX: ld.shared.f32 %f{{[0-9]+}}, [scalar];
|
|
; cast; store
|
|
store float %v, float* %2, align 4
|
|
; PTX: st.shared.f32 [scalar], %f{{[0-9]+}};
|
|
call void @llvm.cuda.syncthreads()
|
|
; PTX: bar.sync 0;
|
|
|
|
; load gep cast
|
|
%4 = load float, float* getelementptr inbounds ([10 x float], [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5), align 4
|
|
; PTX: ld.shared.f32 %f{{[0-9]+}}, [array+20];
|
|
; store gep cast
|
|
store float %v, float* getelementptr inbounds ([10 x float], [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5), align 4
|
|
; PTX: st.shared.f32 [array+20], %f{{[0-9]+}};
|
|
call void @llvm.cuda.syncthreads()
|
|
; PTX: bar.sync 0;
|
|
|
|
; gep cast; load
|
|
%5 = getelementptr inbounds [10 x float], [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5
|
|
%6 = load float, float* %5, align 4
|
|
; PTX: ld.shared.f32 %f{{[0-9]+}}, [array+20];
|
|
; gep cast; store
|
|
store float %v, float* %5, align 4
|
|
; PTX: st.shared.f32 [array+20], %f{{[0-9]+}};
|
|
call void @llvm.cuda.syncthreads()
|
|
; PTX: bar.sync 0;
|
|
|
|
; cast; gep; load
|
|
%7 = addrspacecast [10 x float] addrspace(3)* @array to [10 x float]*
|
|
%8 = getelementptr inbounds [10 x float], [10 x float]* %7, i32 0, i32 %i
|
|
%9 = load float, float* %8, align 4
|
|
; PTX: ld.shared.f32 %f{{[0-9]+}}, [%{{(r|rl|rd)[0-9]+}}];
|
|
; cast; gep; store
|
|
store float %v, float* %8, align 4
|
|
; PTX: st.shared.f32 [%{{(r|rl|rd)[0-9]+}}], %f{{[0-9]+}};
|
|
call void @llvm.cuda.syncthreads()
|
|
; PTX: bar.sync 0;
|
|
|
|
%sum2 = fadd float %1, %3
|
|
%sum3 = fadd float %sum2, %4
|
|
%sum4 = fadd float %sum3, %6
|
|
%sum5 = fadd float %sum4, %9
|
|
ret float %sum5
|
|
}
|
|
|
|
; When hoisting an addrspacecast between different pointer types, replace the
|
|
; addrspacecast with a bitcast.
|
|
define i32 @ld_int_from_float() {
|
|
; IR-LABEL: @ld_int_from_float
|
|
; IR: load i32, i32 addrspace(3)* bitcast (float addrspace(3)* @scalar to i32 addrspace(3)*)
|
|
; PTX-LABEL: ld_int_from_float(
|
|
; PTX: ld.shared.u{{(32|64)}}
|
|
%1 = load i32, i32* addrspacecast(float addrspace(3)* @scalar to i32*), align 4
|
|
ret i32 %1
|
|
}
|
|
|
|
declare void @llvm.cuda.syncthreads() #3
|
|
|
|
attributes #3 = { noduplicate nounwind }
|
|
|