def test_parameters_kernel_regex(profile, format):
data = profile("""
__global__ void addVectors(int* p) {
int a = 5;
*p = a;
}
__global__ void subtractVectors(int* p) {
int a = 5;
*p = a;
}
__global__ void generalKernel(int* p) {
int a = 5;
*p = a;
}
int main() {
int* dptr;
cudaMalloc(&dptr, sizeof(int));
addVectors<<<1, 1>>>(dptr);
subtractVectors<<<1, 1>>>(dptr);
generalKernel<<<1, 1>>>(dptr);
cudaFree(dptr);
return 0;
}""", format=format, kernel_regex=".*a.*ors.*", with_metadata=True)
mappings = data["mappings"]
assert metadata_file("addVectors") in mappings
assert metadata_file("subtractVectors") in mappings
assert metadata_file("generalKernel") not in mappings
assert len(mappings[kernel_file("addVectors", format=format)]["warps"]) == 1
assert len(mappings[kernel_file("subtractVectors", format=format)]["warps"]) == 1
assert kernel_file("generalKernel", format=format) not in mappings
def test_metadata_type_and_name(profile):
data = profile("""
__global__ void kernel(int* p) {
int x = *p;
*p = 5;
}
int main() {
int* dptr;
cudaMalloc(&dptr, sizeof(int));
kernel<<<1, 1>>>(dptr);
cudaFree(dptr);
return 0;
}""")
metadata = data[metadata_file("kernel")]
assert metadata["type"] == "metadata"
assert metadata["kernel"] == "kernel"
def test_metadata_type_index(profile, format):
data = profile("""
__global__ void kernel(int* p) {
int x = *p;
}
int main() {
int* dptr;
cudaMalloc(&dptr, sizeof(int));
kernel<<<1, 1>>>(dptr);
cudaFree(dptr);
return 0;
}""",
format=format)
types = data[metadata_file("kernel")]["typeMap"]
assert len(types) > 0
warp = data[kernel_file("kernel", format=format)]["warps"][0]
assert warp["typeIndex"] == types.index("i32")
def test_metadata_name_index_shared_buffer(profile, format):
data = profile("""
#include <cstdio>
__global__ void kernel() {
__shared__ float arr[10];
arr[threadIdx.x] = threadIdx.x;
}
int main() {
kernel<<<1, 1>>>();
return 0;
}
""",
format=format)
names = data[metadata_file("kernel")]["nameMap"]
assert len(names) > 0
allocations = data[kernel_file("kernel", format=format)]["allocations"][0]
assert allocations["nameIndex"] == names.index("arr")
def test_metadata_type_index_shared_variable(profile, format):
data = profile("""
#include <cstdio>
__global__ void kernel() {
__shared__ float arr;
arr = threadIdx.x;
}
int main() {
kernel<<<1, 1>>>();
return 0;
}
""",
format=format)
types = data[metadata_file("kernel")]["typeMap"]
assert len(types) > 0
warp = data[kernel_file("kernel", format=format)]["warps"][0]
assert warp["typeIndex"] == types.index("float")
def test_metadata_function_debug_info(profile):
body = """// comment
__global__ void kernel(int *p)
{
p[threadIdx.x] = threadIdx.x;
}
int main() {
int* dptr;
cudaMalloc(&dptr, sizeof(int));
kernel<<<1, 1>>>(dptr);
cudaFree(dptr);
return 0;
}
"""
data = profile(body)
metadata = data[metadata_file("kernel")]
assert metadata["source"]["file"].endswith(source_file())
assert metadata["source"]["line"] == 3
assert metadata["source"]["content"].endswith(body)
def test_metadata_debug_location(profile):
data = profile("""
__global__ void kernel(int* p) {
int x = *p;
*p = 5;
}
int main() {
int* dptr;
cudaMalloc(&dptr, sizeof(int));
kernel<<<1, 1>>>(dptr);
cudaFree(dptr);
return 0;
}""",
with_metadata=True)
records = data["mappings"][metadata_file("kernel")]["locations"]
assert len(records) == 2
check_debug_record(data, records[0], "load", 3)
check_debug_record(data, records[1], "store", 4)
def test_metadata_empty_debug(profile):
data = profile("__global__ void kernel() {}", with_main=True)
assert len(data) == 1
assert metadata_file("kernel") not in data