From c8b5d37350a09fc57307697e6428872f61c0f220 Mon Sep 17 00:00:00 2001 From: Eli Bendersky Date: Tue, 15 Apr 2014 16:57:53 +0000 Subject: [PATCH] Add test case for r206302 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@206303 91177308-0d34-0410-b5e6-96231b3b80d8 --- test/CodeGenCUDA/launch-bounds.cu | 30 ++++++++++++++++++++++++++++++ 1 file changed, 30 insertions(+) create mode 100644 test/CodeGenCUDA/launch-bounds.cu diff --git a/test/CodeGenCUDA/launch-bounds.cu b/test/CodeGenCUDA/launch-bounds.cu new file mode 100644 index 0000000000..3f88aeb7e7 --- /dev/null +++ b/test/CodeGenCUDA/launch-bounds.cu @@ -0,0 +1,30 @@ +// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s + +#include "../SemaCUDA/cuda.h" + +#define MAX_THREADS_PER_BLOCK 256 +#define MIN_BLOCKS_PER_MP 2 + +// Test both max threads per block and Min cta per sm. +extern "C" { +__global__ void +__launch_bounds__( MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP ) +Kernel1() +{ +} +} + +// CHECK: !{{[0-9]+}} = metadata !{void ()* @Kernel1, metadata !"maxntidx", i32 256} +// CHECK: !{{[0-9]+}} = metadata !{void ()* @Kernel1, metadata !"minctasm", i32 2} + +// Test only max threads per block. Min cta per sm defaults to 0, and +// CodeGen doesn't output a zero value for minctasm. +extern "C" { +__global__ void +__launch_bounds__( MAX_THREADS_PER_BLOCK ) +Kernel2() +{ +} +} + +// CHECK: !{{[0-9]+}} = metadata !{void ()* @Kernel2, metadata !"maxntidx", i32 256} -- 2.40.0