maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/polly/test/GPGPU/scalar-parameter.ll

426 lines
12 KiB
LLVM
Raw Normal View History

GPGPU: Make sure scops with more than one array work We use this opportunity to add a test case containing a scalar parameter. llvm-svn: 275547
2016-07-15 18:51:14 +08:00
; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \
; RUN: -disable-output < %s | \
; RUN: FileCheck -check-prefix=CODE %s
GPGPU: Pass subtree values correctly to the kernel llvm-svn: 277697
2016-08-04 14:55:49 +08:00
; RUN: opt %loadPolly -polly-codegen-ppcg \
; RUN: -S < %s | \
; RUN: FileCheck -check-prefix=IR %s
GPGPU: Handle scalar array references Pass the content of scalar array references to the alloca on the kernel side and do not pass them additional as normal LLVM scalar value. llvm-svn: 277699
2016-08-04 14:55:59 +08:00
; RUN: opt %loadPolly -polly-codegen-ppcg \
; RUN: -disable-output -polly-acc-dump-kernel-ir < %s | \
; RUN: FileCheck -check-prefix=KERNEL %s
[tests] Specify the dependence to NVPTX backend for Polly ACC test cases Some Polly ACC test cases fail without a working NVPTX backend. We explicitly specify this dependence in REQUIRES. Alternatively, we could have only marked polly-acc as supported in case the NVPTX backend is available, but as we might use other backends in the future, this does not seem to be the best choice. For this to work, we also need to make the 'targets_to_build' information available. Suggested-by: Michael Kruse <llvm@meinersbur.de> llvm-svn: 296853
2017-03-03 11:38:50 +08:00
; REQUIRES: pollyacc,nvptx
GPGPU: Test scalar parameters of type half/float/double/fp128/x86_fp80/ppc_fp128 We currently only test that the code structure we generate for these scalar parameters is correct and we add these types to make sure later code generation additions have sufficient test coverage. In case some of these types cannot be mapped due to missing hardware support on the GPU some of these test cases may need to be updated later on. llvm-svn: 275548
2016-07-15 19:12:29 +08:00
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
GPGPU: Detect read-only scalar arrays ... and pass these by value rather than by reference. llvm-svn: 281837
2016-09-18 03:22:18 +08:00
; KERNEL: define ptx_kernel void @kernel_0(i8* %MemRef_A, float %MemRef_b)
GPGPU: Handle scalar array references Pass the content of scalar array references to the alloca on the kernel side and do not pass them additional as normal LLVM scalar value. llvm-svn: 277699
2016-08-04 14:55:59 +08:00
GPGPU: Make sure scops with more than one array work We use this opportunity to add a test case containing a scalar parameter. llvm-svn: 275547
2016-07-15 18:51:14 +08:00
; CODE: Code
; CODE-NEXT: ====
; CODE-NEXT: # host
; CODE-NEXT: {
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(float), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(32);
GPGPU: Detect read-only scalar arrays ... and pass these by value rather than by reference. llvm-svn: 281837
2016-09-18 03:22:18 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, MemRef_b);
GPGPU: Make sure scops with more than one array work We use this opportunity to add a test case containing a scalar parameter. llvm-svn: 275547
2016-07-15 18:51:14 +08:00
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(float), cudaMemcpyDeviceToHost));
; CODE-NEXT: }
; CODE: # kernel0
; CODE-NEXT: Stmt_bb2(32 * b0 + t0);
; void foo(float A[], float b) {
; for (long i = 0; i < 1024; i++)
; A[i] += b;
; }
;
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
GPGPU: Test scalar parameters of type half/float/double/fp128/x86_fp80/ppc_fp128 We currently only test that the code structure we generate for these scalar parameters is correct and we add these types to make sure later code generation additions have sufficient test coverage. In case some of these types cannot be mapped due to missing hardware support on the GPU some of these test cases may need to be updated later on. llvm-svn: 275548
2016-07-15 19:12:29 +08:00
define void @float(float* %A, float %b) {
GPGPU: Make sure scops with more than one array work We use this opportunity to add a test case containing a scalar parameter. llvm-svn: 275547
2016-07-15 18:51:14 +08:00
bb:
br label %bb1
bb1: ; preds = %bb5, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp6, %bb5 ]
%exitcond = icmp ne i64 %i.0, 1024
br i1 %exitcond, label %bb2, label %bb7
bb2: ; preds = %bb1
%tmp = getelementptr inbounds float, float* %A, i64 %i.0
%tmp3 = load float, float* %tmp, align 4
%tmp4 = fadd float %tmp3, %b
store float %tmp4, float* %tmp, align 4
br label %bb5
bb5: ; preds = %bb2
%tmp6 = add nuw nsw i64 %i.0, 1
br label %bb1
bb7: ; preds = %bb1
ret void
}
GPGPU: Test scalar parameters of type half/float/double/fp128/x86_fp80/ppc_fp128 We currently only test that the code structure we generate for these scalar parameters is correct and we add these types to make sure later code generation additions have sufficient test coverage. In case some of these types cannot be mapped due to missing hardware support on the GPU some of these test cases may need to be updated later on. llvm-svn: 275548
2016-07-15 19:12:29 +08:00
GPGPU: Detect read-only scalar arrays ... and pass these by value rather than by reference. llvm-svn: 281837
2016-09-18 03:22:18 +08:00
; KERNEL: define ptx_kernel void @kernel_0(i8* %MemRef_A, double %MemRef_b)
GPGPU: Handle scalar array references Pass the content of scalar array references to the alloca on the kernel side and do not pass them additional as normal LLVM scalar value. llvm-svn: 277699
2016-08-04 14:55:59 +08:00
; KERNEL-NEXT: entry:
GPGPU: Detect read-only scalar arrays ... and pass these by value rather than by reference. llvm-svn: 281837
2016-09-18 03:22:18 +08:00
; KERNEL-NEXT: %b.s2a = alloca double
; KERNEL-NEXT: store double %MemRef_b, double* %b.s2a
GPGPU: Handle scalar array references Pass the content of scalar array references to the alloca on the kernel side and do not pass them additional as normal LLVM scalar value. llvm-svn: 277699
2016-08-04 14:55:59 +08:00
GPGPU: Test scalar parameters of type half/float/double/fp128/x86_fp80/ppc_fp128 We currently only test that the code structure we generate for these scalar parameters is correct and we add these types to make sure later code generation additions have sufficient test coverage. In case some of these types cannot be mapped due to missing hardware support on the GPU some of these test cases may need to be updated later on. llvm-svn: 275548
2016-07-15 19:12:29 +08:00
; CODE: Code
; CODE-NEXT: ====
; CODE-NEXT: # host
; CODE-NEXT: {
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(double), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(32);
GPGPU: Detect read-only scalar arrays ... and pass these by value rather than by reference. llvm-svn: 281837
2016-09-18 03:22:18 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, MemRef_b);
GPGPU: Test scalar parameters of type half/float/double/fp128/x86_fp80/ppc_fp128 We currently only test that the code structure we generate for these scalar parameters is correct and we add these types to make sure later code generation additions have sufficient test coverage. In case some of these types cannot be mapped due to missing hardware support on the GPU some of these test cases may need to be updated later on. llvm-svn: 275548
2016-07-15 19:12:29 +08:00
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(double), cudaMemcpyDeviceToHost));
; CODE-NEXT: }
; CODE: # kernel0
; CODE-NEXT: Stmt_bb2(32 * b0 + t0);
; void foo(double A[], double b) {
; for (long i = 0; i < 1024; i++)
; A[i] += b;
; }
;
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
define void @double(double* %A, double %b) {
bb:
br label %bb1
bb1: ; preds = %bb5, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp6, %bb5 ]
%exitcond = icmp ne i64 %i.0, 1024
br i1 %exitcond, label %bb2, label %bb7
bb2: ; preds = %bb1
%tmp = getelementptr inbounds double, double* %A, i64 %i.0
%tmp3 = load double, double* %tmp, align 4
%tmp4 = fadd double %tmp3, %b
store double %tmp4, double* %tmp, align 4
br label %bb5
bb5: ; preds = %bb2
%tmp6 = add nuw nsw i64 %i.0, 1
br label %bb1
bb7: ; preds = %bb1
ret void
}
GPGPU: Test scalar/array types i1/i3/i8/i32/i60/i64/i80/i120/i128/i3000 Arrays with integer base type are similar to arrays with floating point types, with the exception that LLVM's integer types can take some odd values. We add a selection of different values to make sure we correctly round these types when necessary. References to scalar integer types are special, as we currently do not model these types as array accesses as they are considered 'synthesizable' by Polly. As a result, we do not generate explicit data-transfers for them, but instead will need to keep track of all references to 'synthesizable' values separately. At the current stage, this is only visible by missing host-to-device data-transfer calls. In the future, we will also require special code generation strategies. llvm-svn: 275551
2016-07-15 19:33:47 +08:00
; CODE: Code
; CODE-NEXT: ====
; CODE-NEXT: # host
; CODE-NEXT: {
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(i1), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(32);
GPGPU: collect array references Initialize the list of references to a GPU array to ensure that the arrays that need to be passed to kernel calls are computed correctly. Furthermore, the very same information is also necessary to compute synchronization correctly. As the functionality to compute these references is already available, what is left for us to do is only to connect the necessary functionality to compute array reference information. llvm-svn: 275798
2016-07-18 23:44:32 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
GPGPU: Test scalar/array types i1/i3/i8/i32/i60/i64/i80/i120/i128/i3000 Arrays with integer base type are similar to arrays with floating point types, with the exception that LLVM's integer types can take some odd values. We add a selection of different values to make sure we correctly round these types when necessary. References to scalar integer types are special, as we currently do not model these types as array accesses as they are considered 'synthesizable' by Polly. As a result, we do not generate explicit data-transfers for them, but instead will need to keep track of all references to 'synthesizable' values separately. At the current stage, this is only visible by missing host-to-device data-transfer calls. In the future, we will also require special code generation strategies. llvm-svn: 275551
2016-07-15 19:33:47 +08:00
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(i1), cudaMemcpyDeviceToHost));
; CODE-NEXT: }
; CODE: # kernel0
; CODE-NEXT: Stmt_bb2(32 * b0 + t0);
; void foo(i1 A[], i1 b) {
; for (long i = 0; i < 1024; i++)
; A[i] += b;
; }
;
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
define void @i1(i1* %A, i1 %b) {
bb:
br label %bb1
bb1: ; preds = %bb5, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp6, %bb5 ]
%exitcond = icmp ne i64 %i.0, 1024
br i1 %exitcond, label %bb2, label %bb7
bb2: ; preds = %bb1
%tmp = getelementptr inbounds i1, i1* %A, i64 %i.0
%tmp3 = load i1, i1* %tmp, align 4
%tmp4 = add i1 %tmp3, %b
store i1 %tmp4, i1* %tmp, align 4
br label %bb5
bb5: ; preds = %bb2
%tmp6 = add nuw nsw i64 %i.0, 1
br label %bb1
bb7: ; preds = %bb1
ret void
}
; CODE: Code
; CODE-NEXT: ====
; CODE-NEXT: # host
; CODE-NEXT: {
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(i3), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(32);
GPGPU: collect array references Initialize the list of references to a GPU array to ensure that the arrays that need to be passed to kernel calls are computed correctly. Furthermore, the very same information is also necessary to compute synchronization correctly. As the functionality to compute these references is already available, what is left for us to do is only to connect the necessary functionality to compute array reference information. llvm-svn: 275798
2016-07-18 23:44:32 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
GPGPU: Test scalar/array types i1/i3/i8/i32/i60/i64/i80/i120/i128/i3000 Arrays with integer base type are similar to arrays with floating point types, with the exception that LLVM's integer types can take some odd values. We add a selection of different values to make sure we correctly round these types when necessary. References to scalar integer types are special, as we currently do not model these types as array accesses as they are considered 'synthesizable' by Polly. As a result, we do not generate explicit data-transfers for them, but instead will need to keep track of all references to 'synthesizable' values separately. At the current stage, this is only visible by missing host-to-device data-transfer calls. In the future, we will also require special code generation strategies. llvm-svn: 275551
2016-07-15 19:33:47 +08:00
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(i3), cudaMemcpyDeviceToHost));
; CODE-NEXT: }
; CODE: # kernel0
; CODE-NEXT: Stmt_bb2(32 * b0 + t0);
; void foo(i3 A[], i3 b) {
; for (long i = 0; i < 1024; i++)
; A[i] += b;
; }
;
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
define void @i3(i3* %A, i3 %b) {
bb:
br label %bb1
bb1: ; preds = %bb5, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp6, %bb5 ]
%exitcond = icmp ne i64 %i.0, 1024
br i1 %exitcond, label %bb2, label %bb7
bb2: ; preds = %bb1
%tmp = getelementptr inbounds i3, i3* %A, i64 %i.0
%tmp3 = load i3, i3* %tmp, align 4
%tmp4 = add i3 %tmp3, %b
store i3 %tmp4, i3* %tmp, align 4
br label %bb5
bb5: ; preds = %bb2
%tmp6 = add nuw nsw i64 %i.0, 1
br label %bb1
bb7: ; preds = %bb1
ret void
}
; CODE: Code
; CODE-NEXT: ====
; CODE-NEXT: # host
; CODE-NEXT: {
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(i8), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(32);
GPGPU: collect array references Initialize the list of references to a GPU array to ensure that the arrays that need to be passed to kernel calls are computed correctly. Furthermore, the very same information is also necessary to compute synchronization correctly. As the functionality to compute these references is already available, what is left for us to do is only to connect the necessary functionality to compute array reference information. llvm-svn: 275798
2016-07-18 23:44:32 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
GPGPU: Test scalar/array types i1/i3/i8/i32/i60/i64/i80/i120/i128/i3000 Arrays with integer base type are similar to arrays with floating point types, with the exception that LLVM's integer types can take some odd values. We add a selection of different values to make sure we correctly round these types when necessary. References to scalar integer types are special, as we currently do not model these types as array accesses as they are considered 'synthesizable' by Polly. As a result, we do not generate explicit data-transfers for them, but instead will need to keep track of all references to 'synthesizable' values separately. At the current stage, this is only visible by missing host-to-device data-transfer calls. In the future, we will also require special code generation strategies. llvm-svn: 275551
2016-07-15 19:33:47 +08:00
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(i8), cudaMemcpyDeviceToHost));
; CODE-NEXT: }
; CODE: # kernel0
; CODE-NEXT: Stmt_bb2(32 * b0 + t0);
; void foo(i8 A[], i32 b) {
; for (long i = 0; i < 1024; i++)
; A[i] += b;
; }
;
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
define void @i8(i8* %A, i8 %b) {
bb:
br label %bb1
bb1: ; preds = %bb5, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp6, %bb5 ]
%exitcond = icmp ne i64 %i.0, 1024
br i1 %exitcond, label %bb2, label %bb7
bb2: ; preds = %bb1
%tmp = getelementptr inbounds i8, i8* %A, i64 %i.0
%tmp3 = load i8, i8* %tmp, align 4
%tmp4 = add i8 %tmp3, %b
store i8 %tmp4, i8* %tmp, align 4
br label %bb5
bb5: ; preds = %bb2
%tmp6 = add nuw nsw i64 %i.0, 1
br label %bb1
bb7: ; preds = %bb1
ret void
}
GPGPU: Pass subtree values correctly to the kernel llvm-svn: 277697
2016-08-04 14:55:49 +08:00
; IR-LABEL: @i8
GPGPU: Make test cases independent of register numbering [NFC] llvm-svn: 281847
2016-09-18 14:50:28 +08:00
; IR: [[REGA:%.+]] = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef_A)
; IR-NEXT: [[REGB:%.+]] = getelementptr [2 x i8*], [2 x i8*]* %polly_launch_0_params, i64 0, i64 0
; IR-NEXT: store i8* [[REGA:%.+]], i8** %polly_launch_0_param_0
; IR-NEXT: [[REGC:%.+]] = bitcast i8** %polly_launch_0_param_0 to i8*
; IR-NEXT: store i8* [[REGC]], i8** [[REGB]]
GPGPU: Pass subtree values correctly to the kernel llvm-svn: 277697
2016-08-04 14:55:49 +08:00
; IR-NEXT: store i8 %b, i8* %polly_launch_0_param_1
GPGPU: Make test cases independent of register numbering [NFC] llvm-svn: 281847
2016-09-18 14:50:28 +08:00
; IR-NEXT: [[REGD:%.+]] = getelementptr [2 x i8*], [2 x i8*]* %polly_launch_0_params, i64 0, i64 1
; IR-NEXT: store i8* %polly_launch_0_param_1, i8** [[REGD]]
GPGPU: Pass subtree values correctly to the kernel llvm-svn: 277697
2016-08-04 14:55:49 +08:00
GPGPU: Test scalar/array types i1/i3/i8/i32/i60/i64/i80/i120/i128/i3000 Arrays with integer base type are similar to arrays with floating point types, with the exception that LLVM's integer types can take some odd values. We add a selection of different values to make sure we correctly round these types when necessary. References to scalar integer types are special, as we currently do not model these types as array accesses as they are considered 'synthesizable' by Polly. As a result, we do not generate explicit data-transfers for them, but instead will need to keep track of all references to 'synthesizable' values separately. At the current stage, this is only visible by missing host-to-device data-transfer calls. In the future, we will also require special code generation strategies. llvm-svn: 275551
2016-07-15 19:33:47 +08:00
; CODE: Code
; CODE-NEXT: ====
; CODE-NEXT: # host
; CODE-NEXT: {
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(i32), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(32);
GPGPU: collect array references Initialize the list of references to a GPU array to ensure that the arrays that need to be passed to kernel calls are computed correctly. Furthermore, the very same information is also necessary to compute synchronization correctly. As the functionality to compute these references is already available, what is left for us to do is only to connect the necessary functionality to compute array reference information. llvm-svn: 275798
2016-07-18 23:44:32 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
GPGPU: Test scalar/array types i1/i3/i8/i32/i60/i64/i80/i120/i128/i3000 Arrays with integer base type are similar to arrays with floating point types, with the exception that LLVM's integer types can take some odd values. We add a selection of different values to make sure we correctly round these types when necessary. References to scalar integer types are special, as we currently do not model these types as array accesses as they are considered 'synthesizable' by Polly. As a result, we do not generate explicit data-transfers for them, but instead will need to keep track of all references to 'synthesizable' values separately. At the current stage, this is only visible by missing host-to-device data-transfer calls. In the future, we will also require special code generation strategies. llvm-svn: 275551
2016-07-15 19:33:47 +08:00
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(i32), cudaMemcpyDeviceToHost));
; CODE-NEXT: }
; CODE: # kernel0
; CODE-NEXT: Stmt_bb2(32 * b0 + t0);
; void foo(i32 A[], i32 b) {
; for (long i = 0; i < 1024; i++)
; A[i] += b;
; }
;
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
define void @i32(i32* %A, i32 %b) {
bb:
br label %bb1
bb1: ; preds = %bb5, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp6, %bb5 ]
%exitcond = icmp ne i64 %i.0, 1024
br i1 %exitcond, label %bb2, label %bb7
bb2: ; preds = %bb1
%tmp = getelementptr inbounds i32, i32* %A, i64 %i.0
%tmp3 = load i32, i32* %tmp, align 4
%tmp4 = add i32 %tmp3, %b
store i32 %tmp4, i32* %tmp, align 4
br label %bb5
bb5: ; preds = %bb2
%tmp6 = add nuw nsw i64 %i.0, 1
br label %bb1
bb7: ; preds = %bb1
ret void
}
; CODE: Code
; CODE-NEXT: ====
; CODE-NEXT: # host
; CODE-NEXT: {
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(i60), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(32);
GPGPU: collect array references Initialize the list of references to a GPU array to ensure that the arrays that need to be passed to kernel calls are computed correctly. Furthermore, the very same information is also necessary to compute synchronization correctly. As the functionality to compute these references is already available, what is left for us to do is only to connect the necessary functionality to compute array reference information. llvm-svn: 275798
2016-07-18 23:44:32 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
GPGPU: Test scalar/array types i1/i3/i8/i32/i60/i64/i80/i120/i128/i3000 Arrays with integer base type are similar to arrays with floating point types, with the exception that LLVM's integer types can take some odd values. We add a selection of different values to make sure we correctly round these types when necessary. References to scalar integer types are special, as we currently do not model these types as array accesses as they are considered 'synthesizable' by Polly. As a result, we do not generate explicit data-transfers for them, but instead will need to keep track of all references to 'synthesizable' values separately. At the current stage, this is only visible by missing host-to-device data-transfer calls. In the future, we will also require special code generation strategies. llvm-svn: 275551
2016-07-15 19:33:47 +08:00
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(i60), cudaMemcpyDeviceToHost));
; CODE-NEXT: }
; CODE: # kernel0
; CODE-NEXT: Stmt_bb2(32 * b0 + t0);
; void foo(i60 A[], i60 b) {
; for (long i = 0; i < 1024; i++)
; A[i] += b;
; }
;
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
define void @i60(i60* %A, i60 %b) {
bb:
br label %bb1
bb1: ; preds = %bb5, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp6, %bb5 ]
%exitcond = icmp ne i64 %i.0, 1024
br i1 %exitcond, label %bb2, label %bb7
bb2: ; preds = %bb1
%tmp = getelementptr inbounds i60, i60* %A, i64 %i.0
%tmp3 = load i60, i60* %tmp, align 4
%tmp4 = add i60 %tmp3, %b
store i60 %tmp4, i60* %tmp, align 4
br label %bb5
bb5: ; preds = %bb2
%tmp6 = add nuw nsw i64 %i.0, 1
br label %bb1
bb7: ; preds = %bb1
ret void
}
; CODE: Code
; CODE-NEXT: ====
; CODE-NEXT: # host
; CODE-NEXT: {
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(i64), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(32);
GPGPU: collect array references Initialize the list of references to a GPU array to ensure that the arrays that need to be passed to kernel calls are computed correctly. Furthermore, the very same information is also necessary to compute synchronization correctly. As the functionality to compute these references is already available, what is left for us to do is only to connect the necessary functionality to compute array reference information. llvm-svn: 275798
2016-07-18 23:44:32 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
GPGPU: Test scalar/array types i1/i3/i8/i32/i60/i64/i80/i120/i128/i3000 Arrays with integer base type are similar to arrays with floating point types, with the exception that LLVM's integer types can take some odd values. We add a selection of different values to make sure we correctly round these types when necessary. References to scalar integer types are special, as we currently do not model these types as array accesses as they are considered 'synthesizable' by Polly. As a result, we do not generate explicit data-transfers for them, but instead will need to keep track of all references to 'synthesizable' values separately. At the current stage, this is only visible by missing host-to-device data-transfer calls. In the future, we will also require special code generation strategies. llvm-svn: 275551
2016-07-15 19:33:47 +08:00
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(i64), cudaMemcpyDeviceToHost));
; CODE-NEXT: }
; CODE: # kernel0
; CODE-NEXT: Stmt_bb2(32 * b0 + t0);
; void foo(i64 A[], i64 b) {
; for (long i = 0; i < 1024; i++)
; A[i] += b;
; }
;
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
define void @i64(i64* %A, i64 %b) {
bb:
br label %bb1
bb1: ; preds = %bb5, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp6, %bb5 ]
%exitcond = icmp ne i64 %i.0, 1024
br i1 %exitcond, label %bb2, label %bb7
bb2: ; preds = %bb1
%tmp = getelementptr inbounds i64, i64* %A, i64 %i.0
%tmp3 = load i64, i64* %tmp, align 4
%tmp4 = add i64 %tmp3, %b
store i64 %tmp4, i64* %tmp, align 4
br label %bb5
bb5: ; preds = %bb2
%tmp6 = add nuw nsw i64 %i.0, 1
br label %bb1
bb7: ; preds = %bb1
ret void
}