Skip to content

Commit e400d17

Browse files
t4c1t4c1
committed
removed redundant include and added tests for volatile intrinsics
1 parent be3b553 commit e400d17

File tree

2 files changed

+96
-2
lines changed

2 files changed

+96
-2
lines changed

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -54,8 +54,6 @@
5454
#include "llvm/Support/X86TargetParser.h"
5555
#include <sstream>
5656

57-
#include <iostream>
58-
5957
using namespace clang;
6058
using namespace CodeGen;
6159
using namespace llvm;

‎llvm/test/CodeGen/NVPTX/atomics-with-semantics.ll

Lines changed: 96 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2558,6 +2558,78 @@ entry:
25582558
; CHECK: st.release.cta.shared.f64
25592559
tail call void @llvm.nvvm.st.shared.f.release.cta.p3f64.f64(double addrspace(3)* %dfp3, double %df);
25602560

2561+
; CHECK: ld.volatile.s32
2562+
%tmpldst144 = tail call i32 @llvm.nvvm.ld.gen.i.volatile.i32.p0(i32* %ip);
2563+
2564+
; CHECK: ld.volatile.global.s32
2565+
%tmpldst145 = tail call i32 @llvm.nvvm.ld.global.i.volatile.i32.p1(i32 addrspace(1)* %ip1);
2566+
2567+
; CHECK: ld.volatile.shared.s32
2568+
%tmpldst146 = tail call i32 @llvm.nvvm.ld.shared.i.volatile.i32.p3(i32 addrspace(3)* %ip3);
2569+
2570+
; CHECK: ld.volatile.s64
2571+
%tmpldst147 = tail call i64 @llvm.nvvm.ld.gen.i.volatile.i64.p0(i64* %llp);
2572+
2573+
; CHECK: ld.volatile.global.s64
2574+
%tmpldst148 = tail call i64 @llvm.nvvm.ld.global.i.volatile.i64.p1(i64 addrspace(1)* %llp1);
2575+
2576+
; CHECK: ld.volatile.shared.s64
2577+
%tmpldst149 = tail call i64 @llvm.nvvm.ld.shared.i.volatile.i64.p3(i64 addrspace(3)* %llp3);
2578+
2579+
; CHECK: ld.volatile.f32
2580+
%tmpldst150 = tail call float @llvm.nvvm.ld.gen.f.volatile.f32.p0(float* %fp);
2581+
2582+
; CHECK: ld.volatile.global.f32
2583+
%tmpldst151 = tail call float @llvm.nvvm.ld.global.f.volatile.f32.p1(float addrspace(1)* %fp1);
2584+
2585+
; CHECK: ld.volatile.shared.f32
2586+
%tmpldst152 = tail call float @llvm.nvvm.ld.shared.f.volatile.f32.p3(float addrspace(3)* %fp3);
2587+
2588+
; CHECK: ld.volatile.f64
2589+
%tmpldst153 = tail call double @llvm.nvvm.ld.gen.f.volatile.f64.p0(double* %dfp);
2590+
2591+
; CHECK: ld.volatile.global.f64
2592+
%tmpldst154 = tail call double @llvm.nvvm.ld.global.f.volatile.f64.p1(double addrspace(1)* %dfp1);
2593+
2594+
; CHECK: ld.volatile.shared.f64
2595+
%tmpldst155 = tail call double @llvm.nvvm.ld.shared.f.volatile.f64.p3(double addrspace(3)* %dfp3);
2596+
2597+
; CHECK: st.volatile.s32
2598+
tail call void @llvm.nvvm.st.gen.i.volatile.p0i32.i32(i32* %ip, i32 %i);
2599+
2600+
; CHECK: st.volatile.global.s32
2601+
tail call void @llvm.nvvm.st.global.i.volatile.p1i32.i32(i32 addrspace(1)* %ip1, i32 %i);
2602+
2603+
; CHECK: st.volatile.shared.s32
2604+
tail call void @llvm.nvvm.st.shared.i.volatile.p3i32.i32(i32 addrspace(3)* %ip3, i32 %i);
2605+
2606+
; CHECK: st.volatile.s64
2607+
tail call void @llvm.nvvm.st.gen.i.volatile.p0i64.i64(i64* %llp, i64 %ll);
2608+
2609+
; CHECK: st.volatile.global.s64
2610+
tail call void @llvm.nvvm.st.global.i.volatile.p1i64.i64(i64 addrspace(1)* %llp1, i64 %ll);
2611+
2612+
; CHECK: st.volatile.shared.s64
2613+
tail call void @llvm.nvvm.st.shared.i.volatile.p3i64.i64(i64 addrspace(3)* %llp3, i64 %ll);
2614+
2615+
; CHECK: st.volatile.f32
2616+
tail call void @llvm.nvvm.st.gen.f.volatile.p0f32.f32(float* %fp, float %f);
2617+
2618+
; CHECK: st.volatile.global.f32
2619+
tail call void @llvm.nvvm.st.global.f.volatile.p1f32.f32(float addrspace(1)* %fp1, float %f);
2620+
2621+
; CHECK: st.volatile.shared.f32
2622+
tail call void @llvm.nvvm.st.shared.f.volatile.p3f32.f32(float addrspace(3)* %fp3, float %f);
2623+
2624+
; CHECK: st.volatile.f64
2625+
tail call void @llvm.nvvm.st.gen.f.volatile.p0f64.f64(double* %dfp, double %df);
2626+
2627+
; CHECK: st.volatile.global.f64
2628+
tail call void @llvm.nvvm.st.global.f.volatile.p1f64.f64(double addrspace(1)* %dfp1, double %df);
2629+
2630+
; CHECK: st.volatile.shared.f64
2631+
tail call void @llvm.nvvm.st.shared.f.volatile.p3f64.f64(double addrspace(3)* %dfp3, double %df);
2632+
25612633
; CHECK: ret
25622634
ret void
25632635
}
@@ -2720,6 +2792,30 @@ declare void @llvm.nvvm.st.shared.f.cta.p3f64.f64(double addrspace(3)* nocapture
27202792
declare void @llvm.nvvm.st.gen.f.release.cta.p0f64.f64(double* nocapture, double) #1
27212793
declare void @llvm.nvvm.st.global.f.release.cta.p1f64.f64(double addrspace(1)* nocapture, double) #1
27222794
declare void @llvm.nvvm.st.shared.f.release.cta.p3f64.f64(double addrspace(3)* nocapture, double) #1
2795+
declare i32 @llvm.nvvm.ld.gen.i.volatile.i32.p0(i32* nocapture) #1
2796+
declare i32 @llvm.nvvm.ld.global.i.volatile.i32.p1(i32 addrspace(1)* nocapture) #1
2797+
declare i32 @llvm.nvvm.ld.shared.i.volatile.i32.p3(i32 addrspace(3)* nocapture) #1
2798+
declare i64 @llvm.nvvm.ld.gen.i.volatile.i64.p0(i64* nocapture) #1
2799+
declare i64 @llvm.nvvm.ld.global.i.volatile.i64.p1(i64 addrspace(1)* nocapture) #1
2800+
declare i64 @llvm.nvvm.ld.shared.i.volatile.i64.p3(i64 addrspace(3)* nocapture) #1
2801+
declare float @llvm.nvvm.ld.gen.f.volatile.f32.p0(float* nocapture) #1
2802+
declare float @llvm.nvvm.ld.global.f.volatile.f32.p1(float addrspace(1)* nocapture) #1
2803+
declare float @llvm.nvvm.ld.shared.f.volatile.f32.p3(float addrspace(3)* nocapture) #1
2804+
declare double @llvm.nvvm.ld.gen.f.volatile.f64.p0(double* nocapture) #1
2805+
declare double @llvm.nvvm.ld.global.f.volatile.f64.p1(double addrspace(1)* nocapture) #1
2806+
declare double @llvm.nvvm.ld.shared.f.volatile.f64.p3(double addrspace(3)* nocapture) #1
2807+
declare void @llvm.nvvm.st.gen.i.volatile.p0i32.i32(i32* nocapture, i32) #1
2808+
declare void @llvm.nvvm.st.global.i.volatile.p1i32.i32(i32 addrspace(1)* nocapture, i32) #1
2809+
declare void @llvm.nvvm.st.shared.i.volatile.p3i32.i32(i32 addrspace(3)* nocapture, i32) #1
2810+
declare void @llvm.nvvm.st.gen.i.volatile.p0i64.i64(i64* nocapture, i64) #1
2811+
declare void @llvm.nvvm.st.global.i.volatile.p1i64.i64(i64 addrspace(1)* nocapture, i64) #1
2812+
declare void @llvm.nvvm.st.shared.i.volatile.p3i64.i64(i64 addrspace(3)* nocapture, i64) #1
2813+
declare void @llvm.nvvm.st.gen.f.volatile.p0f32.f32(float* nocapture, float) #1
2814+
declare void @llvm.nvvm.st.global.f.volatile.p1f32.f32(float addrspace(1)* nocapture, float) #1
2815+
declare void @llvm.nvvm.st.shared.f.volatile.p3f32.f32(float addrspace(3)* nocapture, float) #1
2816+
declare void @llvm.nvvm.st.gen.f.volatile.p0f64.f64(double* nocapture, double) #1
2817+
declare void @llvm.nvvm.st.global.f.volatile.p1f64.f64(double addrspace(1)* nocapture, double) #1
2818+
declare void @llvm.nvvm.st.shared.f.volatile.p3f64.f64(double addrspace(3)* nocapture, double) #1
27232819

27242820
declare i32 @llvm.nvvm.atomic.add.gen.i.acquire.i32.p0i32(i32* nocapture, i32) #1
27252821
declare i32 @llvm.nvvm.atomic.add.global.i.acquire.i32.p1i32(i32 addrspace(1)* nocapture, i32) #1

0 commit comments

Comments
 (0)